phyronnaz / hashdag Goto Github PK

When running with a clean compile I get this error:

Starting...
ERROR /data/transfer/HashDAG/src/dag_tracer.cu:135: cuda error "no kernel image is available for execution on the device"

This because my hardware only supports CUDA compute capability 5.0.
Adding the following lines to CMakeLists.txt and clean compiling resolves the issue. This will make sure that various compute capacity images are generated:

if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
  set(CMAKE_CUDA_ARCHITECTURES 35 50 61 72)
endif()

On a side note: Setting the -gencode arch=compute_50,code=sm_50 flags to the correct version is not removing the error.

Hello,

I have been working with the HashDAG for a little while now and am still working to get a feel for things but I was wondering if you had any thoughts on these ideas:

1. Animation of a group of voxels?
2. Including external object and animating them?
3. Enhancing world size? I have some simple functions for this which also allows for SVO to lattice coordinates and visa versa.
4. Dynamically reducing/enlarging voxel sizes? (perhaps for the millimeter or micrometer to move towards more photorealistic scenes)

There are some additional ideas that I have as well but these cover basic animation and voxel resolution.

One other really interesting idea that I am playing around with is "Symmetry-aware Sparse Voxel DAGs (SSVDAGs) for compression-domain tracing of high-resolution geometric scenes" (https://jcgt.org/published/0006/02/01/) which could integrated into the HashDAG as well as having multiple light-sources for ray tracing/casting effects.

Looking forward to hearing back from you on these ideas.

Hi there,

I'm running into some compilation issues on Windows (Visual Studio 2019):

1>hash_table.cpp
1>C:\Users\mathi\Documents\GitHub\HashDAG\src\dags\hash_dag\hash_table.cpp(22,5): error C2653: 'Memory': is not a class or namespace name
1>C:\Users\mathi\Documents\GitHub\HashDAG\src\dags\hash_dag\hash_table.cpp(22,20): error C2062: type 'int' unexpected
1>C:\Users\mathi\Documents\GitHub\HashDAG\src\dags\hash_dag\hash_table.cpp(22,38): error C2653: 'EMemoryType': is not a class or namespace name
1>C:\Users\mathi\Documents\GitHub\HashDAG\src\dags\hash_dag\hash_table.cpp(24,23): error C2653: 'Memory': is not a class or namespace name
1>C:\Users\mathi\Documents\GitHub\HashDAG\src\dags\hash_dag\hash_table.cpp(24,38): error C2275: 'uint32': illegal use of this type as an expression
1>C:\Users\mathi\Documents\GitHub\HashDAG\src\typedefs.h(668): message : see declaration of 'uint32'
...

The memory file memory.h is included at the top of hash_table.cpp but it is not picking up the Memory class declared inside.
The cause of the issue is that the compiler is that Microsoft also ships a memory.h file in their UCRT library:
https://docs.microsoft.com/en-us/cpp/c-runtime-library/reference/memset-wmemset?view=msvc-160

This can be verified by using the go-to-definition tool in Visual Studio.
By default it will jump to the projects memory.h which unfortunately is not what the compiler does (IntelliSense is using a different back-end that is not based on MSVC).
When I change the name of memory.h to something else, using go-to-definition will indeed take me to memory.h in the UCRT library.

My suggested change is to rename memory.h to something that is less likely to clash with system header files.

Best regards,
Mathijs

There is a bug in the function find_interior_node_in_bucket() in hash_table.h that prevents the hash table from finding some nodes.
When iterating through the memory pages that make up a hash table bucket the loop stops one element too early:

uint32 pageEndIndex = std::min(bucketSize, pindex + C_pageSize);
if (pindex + nodeSize >= pageEndIndex)
    return 0xFFFFFFFF;

// ...

pageEndIndex -= nodeSize;

// ...

for (uint32 index = pindex; index < pageEndIndex;) {
    ...
}

When the page contains only the node we're looking for then pageEndIndex == pindex + nodeSize. However the bounds check if statement will automatically return not found due. This can be easily fixed by changing the comparison operator:

if (pindex + nodeSize > pageEndIndex)
    return 0xFFFFFFFF;

The second issue is the loop itself. At this point the size of the node being searched for has been subtracted from pageEndIndex which now points to the final index in the page that may contain the node without going out of bounds. Hence the comparison operator in the if statement is also wrong and should be greater than or equal:

if (pindex + nodeSize >= pageEndIndex)
    return 0xFFFFFFFF;

Due to a bug in Memory::realloc_impl DynamicArray cannot grow/shrink if they point to GPU memory (either GPU_Malloc or GPU_Managed).
Grow/shrinking calls the Memory::realloc_impl function which checks whether the allocation contains CPU memory or GPU accessible memory.
In case of the latter cuda_memcpy_impl is called with cudaMemcpyDefault as cudaMemcpyKind.

if (is_gpu_type(oldAlloc.type)) {
        ptr = malloc_impl(newSize, oldAlloc.name, oldAlloc.type);
        if (ptr)
            cuda_memcpy_impl(static_cast<uint8*>(ptr), static_cast<uint8*>(oldPtr), oldAlloc.size, cudaMemcpyDefault);
        free_impl(oldPtr);
    } else {
        ...
    }

However, cuda_memcpy_impl does not support cudaMemcpyDefault and silently fails, causing the DynamicArray to contain junk values (whatever malloc returned):

void Memory::cuda_memcpy_impl(uint8* dst, const uint8* src, uint64 size, cudaMemcpyKind memcpyKind)
{
    const auto BlockCopy = [&]() {
        const double Start = Utils::seconds();
        CUDA_CHECKED_CALL cudaMemcpy(dst, src, size, memcpyKind);
        const double End = Utils::seconds();

        return size / double(1u << 30) / (End - Start);
    };

    if (memcpyKind == cudaMemcpyDeviceToDevice) {
        PROFILE_SCOPEF("Memcpy HtH %fMB", size / double(1u << 20));
        [[maybe_unused]] const double Bandwidth = BlockCopy();
        ZONE_METADATA("%fGB/s", Bandwidth);
    } else if (memcpyKind == cudaMemcpyDeviceToHost) {
        PROFILE_SCOPEF("Memcpy DtH %fMB", size / double(1u << 20));
        [[maybe_unused]] const double Bandwidth = BlockCopy();
        ZONE_METADATA("%fGB/s", Bandwidth);
    } else if (memcpyKind == cudaMemcpyHostToDevice) {
        PROFILE_SCOPEF("Memcpy HtD %fMB", size / double(1u << 20));
        [[maybe_unused]] const double Bandwidth = BlockCopy();
        ZONE_METADATA("%fGB/s", Bandwidth);
    }
}

This could be fixed by adding an extra else if statement to cuda_memcpy_impl to handle the cudaMemcpyDefault case.

Hello,

I am highly interested in the HashDAG approach and was able to run the Windows demo that used Epic Citadel and I was wondering what other types of scenes could be used as I am still trying to get a feel for things?

In particular, I am interested in what type of files could be loaded and what limitations there might be on this.

Any guidance would be greatly appreciated.

Thanks