switch to indexing child nodes by start and stop indices instead of start and length; this reduces a bit of arithmetic when traversing the bounding volume hierarchy and makes the code in Geometry.build() more concise. add an untested cuda kernel to interleave the bits in three uint64 arrays.

1 files changed, 17 insertions, 0 deletions

diff --git a/src/tools.cu b/src/tools.cu
new file mode 100644
index 0000000..3d3fed7
--- /dev/null
+++ b/src/tools.cu
@@ -0,0 +1,17 @@
+//--*-c-*-
+
+extern "C"
+{
+
+__global__ void interleave(int nthreads, unsigned long long *x, unsigned long long *y, unsigned long long *z, int bits, unsigned long long *dest)
+{
+	int id = blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (id >= nthreads)
+		return;
+
+	for (int i=0; i < bits; i++)
+		dest[id] |= (x[id] & 1 << i) << (2*i) | (y[id] & 1 << i) << (2*i+1) | (z[id] & 1 << i) << (2*i+2);
+}
+
+}