diff options
| -rw-r--r-- | chroma/cuda/geometry.h | 11 |
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) { |