diff options
| author | Stan Seibert <stan@mtrr.org> | 2012-05-14 11:37:18 -0400 |
|---|---|---|
| committer | tlatorre <tlatorre@uchicago.edu> | 2021-05-09 08:42:39 -0700 |
| commit | c83e502705ae1312c6030a4e9de2b4327e2d1dd7 (patch) | |
| tree | 93237026aa0c70bdf97759a2fff6b4beddc4f5d2 /chroma/cuda | |
| parent | 6d515c3c8035c0899f7d27344c6ee545f3ce6f8a (diff) | |
| download | chroma-c83e502705ae1312c6030a4e9de2b4327e2d1dd7.tar.gz chroma-c83e502705ae1312c6030a4e9de2b4327e2d1dd7.tar.bz2 chroma-c83e502705ae1312c6030a4e9de2b4327e2d1dd7.zip | |
GPU geometry modification to permit the BVH node storage to be split
between GPU and CPU. This allows much more complex geometries to
be run on CUDA devices with less memory.
The GPUGeometry object now takes a min_free_gpu_mem parameter giving
the minimum number of bytes that can be free on the GPU after the BVH
is loaded. By default, this number is 300 MB. Cards with sufficient
memory will have the entire BVH on card, but those without enough
memory will have the BVH split such that the top of the hierarchy (the
most frequently traversed) is on the GPU.
Diffstat (limited to 'chroma/cuda')
| -rw-r--r-- | chroma/cuda/geometry.h | 7 | ||||
| -rw-r--r-- | chroma/cuda/geometry_types.h | 4 |
2 files changed, 9 insertions, 2 deletions
diff --git a/chroma/cuda/geometry.h b/chroma/cuda/geometry.h index b718655..4d21779 100644 --- a/chroma/cuda/geometry.h +++ b/chroma/cuda/geometry.h @@ -22,7 +22,12 @@ __device__ uint4 read_skip_l1(uint4 *ptr) __device__ Node get_node(Geometry *geometry, const unsigned int &i) { - uint4 node = geometry->nodes[i]; + uint4 node; + if (i < geometry->nprimary_nodes) + node = geometry->primary_nodes[i]; + else + node = geometry->extra_nodes[i - geometry->nprimary_nodes]; + Node node_struct; uint3 lower_int = make_uint3(node.x & 0xFFFF, node.y & 0xFFFF, node.z & 0xFFFF); diff --git a/chroma/cuda/geometry_types.h b/chroma/cuda/geometry_types.h index 6da1a47..7e8bd4b 100644 --- a/chroma/cuda/geometry_types.h +++ b/chroma/cuda/geometry_types.h @@ -57,11 +57,13 @@ struct Geometry uint3 *triangles; unsigned int *material_codes; unsigned int *colors; - uint4 *nodes; + uint4 *primary_nodes; + uint4 *extra_nodes; Material **materials; Surface **surfaces; float3 world_origin; float world_scale; + int nprimary_nodes; }; #endif |