I’m currently working on voxel project in university. My work is based on a highly modified version of the gvdb-library. One of the changes is, that I want to use only one Bit per voxel instead of char/float etc…
An important part is removing the voxel at places where they collide with the position of points. This is easy if each voxel is of size one byte or larger: the position of the points is read in a kernel and zero is written to the location where each of the intersecting voxel is stored.
However, if one voxel is only the size of a bit - one dimension eg. In x direction extends by 8. If we now launch kernel with one thread per point and if two points lie within one char(8voxel), two threads would each set a voxel/bit to zero and write it to global memory. We’d now have a race condition and the whole world collapses.
So the question is: What is the best way to do this?
Before I present my current thoughts, here are a few facts: The number of points that are interacting with the voxelstructure is in the order of 10^4 to 10^6. The points do often cover a surface and are separated by approximately one voxel size apart from each other.
My current approach would be to sort the points via cub::radixsort in x,y and z and bin all points that have the same value with respect to x/8 to a byte-long bitmask, such that I can deal with all Voxel within a byte at once.
I’d be really glad, if you guys have suggestions, ideas or feedback. If you need any more information, just ask away, I just wanted to keep it short and on point for now.
All the best from munich,