From 4f3e0b7709bb64ffc24f2a003509d5f480848239 Mon Sep 17 00:00:00 2001
From: Anthony LaTorre <tlatorre9@gmail.com>
Date: Tue, 9 Aug 2011 16:39:40 -0400
Subject: 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.

---
 src/tools.cu | 17 +++++++++++++++++
 1 file changed, 17 insertions(+)
 create mode 100644 src/tools.cu

(limited to 'src/tools.cu')

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);
+}
+
+}
-- 
cgit