diff options
| author | Stan Seibert <stan@mtrr.org> | 2012-02-16 22:44:42 -0500 |
|---|---|---|
| committer | tlatorre <tlatorre@uchicago.edu> | 2021-05-09 08:42:38 -0700 |
| commit | 3cd18a51d4e73d9148fc721527880c8dbbd08871 (patch) | |
| tree | 4e0fec56aaf7125bf57e1f9e0ad6e7f1a628d08f /chroma/cuda | |
| parent | b509d3bda9f735a48a77d98c24c9c48a276ebfb0 (diff) | |
| download | chroma-3cd18a51d4e73d9148fc721527880c8dbbd08871.tar.gz chroma-3cd18a51d4e73d9148fc721527880c8dbbd08871.tar.bz2 chroma-3cd18a51d4e73d9148fc721527880c8dbbd08871.zip | |
Collapse chains of BVH nodes with single children.
Diffstat (limited to 'chroma/cuda')
| -rw-r--r-- | chroma/cuda/bvh.cu | 15 |
1 files changed, 15 insertions, 0 deletions
diff --git a/chroma/cuda/bvh.cu b/chroma/cuda/bvh.cu index 0c64157..882c22e 100644 --- a/chroma/cuda/bvh.cu +++ b/chroma/cuda/bvh.cu @@ -522,4 +522,19 @@ extern "C" node[b_index] = temp4; } + __global__ void collapse_child(unsigned int start, unsigned int end, + uint4 *node) + { + unsigned int thread_id = blockDim.x * blockIdx.x + threadIdx.x; + unsigned int stride = gridDim.x * blockDim.x; + + for (unsigned int i=start+thread_id; i < end; i += stride) { + uint4 this_node = node[i]; + unsigned int nchild = this_node.w >> CHILD_BITS; + unsigned int child_id = this_node.w & ~NCHILD_MASK; + if (nchild == 1) + node[i] = node[child_id]; + } + } + } // extern "C" |