diff options
| author | Stan Seibert <stan@mtrr.org> | 2012-02-17 18:26:41 -0500 |
|---|---|---|
| committer | tlatorre <tlatorre@uchicago.edu> | 2021-05-09 08:42:38 -0700 |
| commit | cc386a0dbad610f4c2ed6a79ed1cfc2af57aade3 (patch) | |
| tree | 017fb4cc93f526dfc61951f282fb9f26526fd478 /chroma/cuda | |
| parent | e2bf16de5d94a7b438e9c1af52f0bb4dc632b35f (diff) | |
| download | chroma-cc386a0dbad610f4c2ed6a79ed1cfc2af57aade3.tar.gz chroma-cc386a0dbad610f4c2ed6a79ed1cfc2af57aade3.tar.bz2 chroma-cc386a0dbad610f4c2ed6a79ed1cfc2af57aade3.zip | |
Add an argsort_direction() function to chroma.tools and use it to
group photons so that they take similar paths on the GPU.
argsort_direction() morton-orders an array of normalized direction
vectors according to their spherical coordinates. Photons sorted in
this way tend to follow similar paths through a detector geometry,
which enhances cache locality. As a result, get_node() uses the GPU
L1 cache again, with good results.
Diffstat (limited to 'chroma/cuda')
| -rw-r--r-- | chroma/cuda/geometry.h | 2 |
1 files changed, 1 insertions, 1 deletions
diff --git a/chroma/cuda/geometry.h b/chroma/cuda/geometry.h index 0735792..3db6dad 100644 --- a/chroma/cuda/geometry.h +++ b/chroma/cuda/geometry.h @@ -22,7 +22,7 @@ __device__ uint4 read_skip_l1(uint4 *ptr) __device__ Node get_node(Geometry *geometry, const unsigned int &i) { - uint4 node = read_skip_l1(geometry->nodes + i); + uint4 node = geometry->nodes[i]; Node node_struct; if (node.x == 0) { |