diff options
| author | Stan Seibert <stan@mtrr.org> | 2013-01-28 21:03:00 -0500 |
|---|---|---|
| committer | tlatorre <tlatorre@uchicago.edu> | 2021-05-09 08:42:39 -0700 |
| commit | 6706f67febc8fb234630986894df67fd948970ac (patch) | |
| tree | 53001ab3b00d32490bce82ae67767d0adca10ec9 | |
| parent | 1f1bf83bc4b38a36dc0ac6f877f952b103e4de0c (diff) | |
| download | chroma-6706f67febc8fb234630986894df67fd948970ac.tar.gz chroma-6706f67febc8fb234630986894df67fd948970ac.tar.bz2 chroma-6706f67febc8fb234630986894df67fd948970ac.zip | |
Remove unused function to read a word bypassing L1 cache. Seems to be incorrect asm syntax.
| -rw-r--r-- | chroma/cuda/geometry.h | 9 |
1 files changed, 0 insertions, 9 deletions
diff --git a/chroma/cuda/geometry.h b/chroma/cuda/geometry.h index 6a930e7..c81d9b9 100644 --- a/chroma/cuda/geometry.h +++ b/chroma/cuda/geometry.h @@ -10,15 +10,6 @@ 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__ uint4 get_packed_node(Geometry *geometry, const unsigned int &i) { |