Skip L1 cache when loading nodes.

Node access is very irregular as each thread descends the BVH tree.
Each node is only 16 bytes, so the 128 byte cache line size in the L1
cache means that a lot of useless data is often fetched.  Using some
embedded PTX, we can force the L1 cache to be skipped, going directly
to L2.  The L2 cache line is 32 bytes long, which means that both
children in a binary tree will be cached at the same time.

This improves the speed on the default generated binary trees, but
does not help an optimized tree yet.

1 files changed, 10 insertions, 1 deletions

diff --git a/chroma/cuda/geometry.h b/chroma/cuda/geometry.h
index 86ced6c..577da72 100644
--- a/chroma/cuda/geometry.h
+++ b/chroma/cuda/geometry.h
@@ -13,10 +13,19 @@ to_float3(const uint3 &a)
   return make_float3(a.x, a.y, a.z);
 }
 
+__device__ uint4 read_skip_l1(uint4 *ptr)
+{
+  uint4 val;
+  asm(" ld.cg.v4.u32 {%0, %1, %2, %3}, [%4];"
+      : "=r"(val.x), "=r"(val.y), "=r"(val.z), "=r"(val.w) 
+      : "l"(ptr) : "memory");
+  return val;
+}
+
 __device__ Node
 get_node(Geometry *geometry, const unsigned int &i)
 {
-    uint4 node = geometry->nodes[i];
+    uint4 node = read_skip_l1(geometry->nodes + i);
     Node node_struct;
 
     if (node.x == 0) {