In this series we explore the wonderful world of the bounding volume hierarchy: the mysterious data structure that enables real-time ray tracing on recent GPUs, in complex, animated scenes. Despite the broad adoption of ray tracing in games, in-depth understanding of the underlying technology seems to be reserved to a select few. In this article we show how a proper BVH is implemented without support from DXR / RTX / OptiX / Vulkan, in easy-to-read ‘sane C++’. Whether this is actually useful for you, or just interesting as background information, depends on your needs, obviously.
This series consist of ten articles. For each article, full source code is available on Github, in easy-to-read ‘Sane C++’. Contents:
- Basics: data structure, simple construction and ray traversal.
- Faster rays: the Surface Area Heuristic and ordered traversal.
- Faster construction: binned BVH building.
- BVHs for animated models: refitting and rebuilding.
- BVHs for animated scenes: the top-level BVH.
- TLAS & BLAS part 2 – all together now.
- Consolidating, part 1: ray tracing textured and shaded meshes.
- Consolidating, part 2: Whitted-style ray tracing.
- OpenCL: GPU ray tracing, part 1 (this article).
- OpenCL: GPU ray tracing, part 2 (series final).
We approach the end of this series. After building the BVH on the CPU, and rendering an animated scene with textured and reflective teapots on the CPU, we take the rendering process to the GPU, chasing visions of better performance. The GPU bits in this article will be implemented in OpenCL. The GPU code will not rely on any hardware ray tracing; in fact, we’ll not use any of that. Of course we will do some performance measurements, and compare against RTX – in article 10.
GPU
Compared to a CPU, a GPU is a strange beast. It sits pretty isolated in a computer: separated, by a slow connection, from CPU, RAM, user input, disk I/O and other commodities. The only thing it has quick access to is the display. In terms of processing it’s an oddity as well: a large group of cores chews on an even larger number of threads, which ideally all do the same thing, just on different data. Still, the prospect of vast bandwidth and great compute power is tempting, and justifies some pain and effort to get this beast under control.
In this article we will use OpenCL to access the GPU. This is a somewhat arbitrary choice; e.g. CUDA (NVIDIA only) or compute shaders would also do the job. Regardless of the chosen GPU programming framework, we’ll have to deal with an API, which can be a bit cumbersome, especially compared to the quite pure C++ we have been using so far. To make our lives a bit easier, templates exist, and the template we’ve been using so far comes with some decent OpenCL functionality. If you are not using the template, please copy the OpenCL specifics from the template to your own code; it’s a pretty thin wrapper.
The OpenCL program we’re going to run is pretty simple:
|
1 2 3 4 5 6 7 8 9 |
__kernel void render( __global uint* target, int offset ) { // plot a pixel into the target array in GPU memory int threadIdx = get_global_id( 0 ); int x = threadIdx % 640; int y = threadIdx / 640; int red = x / 3 + offset, green = y / 3; target[x + y * 640] = (red << 16) + (green << 8); } |
The code would produce a a nice gradient, but we’re not quite there yet. Someone needs to start this code, ideally from some C++. To do that, we store the OpenCL code in a file, kernels.cl, and load it from the C++ code that we have been developing. To ::Init() we add:
|
1 2 3 |
// prepare OpenCL tracer = new Kernel( "cl/kernels.cl", "render" ); target = new Buffer( 640 * 640 * 4 ); |
The first line loads the OpenCL source file from subdirectory cl, compiles it, and looks for a kernel function called render. The second line creates a Buffer of 640 * 640 * 4 bytes. This ‘Buffer’ is something special: the bytes it stores can reside in regular RAM, or on the GPU, or both. The two locations are not automatically in sync, but we can copy from VRAM to RAM (using Buffer::CopyFromDevice() ) or from RAM to VRAM (using Buffer::CopyToDevice() ). So now we have a way to provide the render kernel with data, and we have a way to obtain results, back to the CPU, where we can operate on them more easily.
We can now execute the kernel, e.g. in the ::Tick() function:
|
1 2 |
tracer->SetArguments( target, 0 ); tracer->Run( 640 * 640 ); |
The kernel is expecting two arguments: the first is a __global uint*, the second a simple int. The call to SetArguments(...) passes the buffer we created, and 0 for the integer, to OpenCL; the second line starts the kernel on the GPU, using 640×640 threads.
Once the kernel is done (which will be in an instant), we can get the results by copying the buffer from the GPU to RAM:
|
1 |
target->CopyFromDevice(); |
The buffer now contains 640×640 unsigned integers, which we can now simply copy to the screen surface. We could use a nested for loop over x and y for that, but a memcpy does the job more quickly:
|
1 |
memcpy( screen->pixels, target->GetHostPtr(), target->size ); |
And now we can start the program, and admire its glorious GPU-calculated gradient.
With that out of the way, we can now focus on getting some ray tracing done on the GPU. We just created GPU code that produces colors that we can display; all that we need to do now is calculate these colors using a Trace function – running on the GPU. The Trace function will need data: a skybox, triangles, and of course: a BVH.
GPGPU Best Practices
From here, we will port the existing CPU ray tracer to the GPU, one bit at a time. A few ‘best practices’ will help to make this process (and probably other GPGPU projects) go smooth:
First of all, we start with working CPU code. Debugging on the GPU is hard; it really helps if we don’t complicate things further by building new things directly in GPU code.
Secondly, start with something simple and build from there. We just did that: getting anything to the screen (‘first image‘) is an important first step, from a practical and psychological point of view. From there, small, verifiable steps will ultimately lead to a finished project.
Third, functionality comes before speed. “The most important optimization is the one that turns a not working program into a working one,” and: “Premature optimization is the root of all evil.” Applied to GPGPU: make it work first, in the most intuitive manner; tweaking is for later.
And finally, let the GPU do what it does best. In our case: Building a BVH is not easy to do on a massive parallel architecture. So: don’t! Tracing rays on the other hand is a perfect fit: We basically execute the same task, for thousands of pixels.
Trace
Before we proceed, let’s make some small changes to the render kernel, so that it can work with a different window size than the hardcoded 640×640 we used so far.
|
1 2 3 4 5 6 7 8 9 |
__kernel void render( __global uint* target, int offset ) { // plot a pixel into the target array in GPU memory int threadIdx = get_global_id( 0 ); int x = threadIdx % SCRWIDTH; int y = threadIdx / SCRWIDTH; int red = x / 3 + offset, green = y / 3; target[x + y * 640] = (red << 16) + (green << 8); } |
This requires two additional changes. The constants SCRWIDTH and SCRHEIGHT (used later) are defined in a file named common.h, which can be found in the template folder. We thus include this file at the top of kernels.cl:
|
1 |
#include "template/common.h" |
Note that common.h will now be included by the CPU code as well as the GPU code. This is a great way to share some functionality, and to ensure that both use the same settings (such as screen size, in this case).
The second change is a replacement of 640 * 640 by SCRWIDTH * SCRHEIGHT in the CPU code. This occurs in two places: once for the allocation of the target buffer, and once for the invocation of the trace kernel, in ::Tick.
Time to do some Tracin’. In the render kernel we can use pixel coordinates x and y to generate a primary ray:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 |
__kernel void render( __global uint* target, ... ) { // plot a pixel into the target array in GPU memory int threadIdx = get_global_id( 0 ); int x = threadIdx % SCRWIDTH; int y = threadIdx / SCRWIDTH; // create a primary ray for the pixel struct Ray ray; ray.O = camPos; float3 pixelPos = ray.O + p0 + (p1 - p0) * ((float)x / SCRWIDTH) + (p2 - p0) * ((float)y / SCRHEIGHT); ray.D = normalize( pixelPos - ray.O ); ray.hit.t = 1e30f; // 1e30f denotes 'no hit' // trace the primary ray ... } |
Setting up the ray requires a camera position camPos and three screen corners p0, p1 and p2. On the CPU we would calculate those outside the tile/pixel loop, because they are the same for all pixels. For GPU code, the only way to do something ‘outside the loop’ is to do it on the CPU. We thus pass the four positions as arguments of the kernel:
|
1 2 3 4 |
__kernel void render( __global uint* target, float3 camPos, float3 p0, float3 p1, float3 p2 ) { ... } |
On the CPU-side, the change is straight-forward:
|
1 |
tracer->SetArguments( target, camPos, p0, p1, p2 ); |
Setting up a ray requires a Ray class. We could move the CPU-side Ray definition to common.h to use the same struct on CPU and GPU, but that has issues: The CPU Ray definition uses __m128 variables for SIMD processing, which is not supported on GPUs. It is simpler to just re-define Ray in kernels.cl, and, with it, Intersection:
|
1 2 3 4 5 6 7 8 9 10 11 12 |
struct Intersection { float t; // intersection distance along ray float u, v; // barycentric coordinates of the intersection uint instPrim; // instance index (12 bit) and primitive index (20 bit) }; struct Ray { float3 O, D, rD; // in OpenCL, each of these will be padded to 16 bytes struct Intersection hit; // total ray size: 64 bytes }; |
Note the strong similarity between C code and OpenCL code. Some differences exist as well though; where C++ lets us get away with just using Intersection as a type for member hit in struct Ray, OpenCL insists on calling it struct Intersection hit. Likewise, Ray ray in the kernel must be labeled struct Ray ray. Another notable difference that we encounter early on: you can’t pass anything ‘by reference’, so when we call the Trace function, we either pass the ray ‘by value’ or using a pointer.
|
1 2 3 4 |
float3 Trace( struct Ray* ray ) { return (float3)( 1, 1, 1 ); } |
The initial version of function Trace (which just returns a white pixel) reveals one more nuisance. A float3 is not constructed using float3( 1, 1, 1 ), but instead using (float3)( 1, 1, 1 ), as shown in the code snippet.
Function Trace returns a float3, like we did in the CPU code. Since the buffer that the kernel operates on is in fact an unsigned int buffer, we need to convert the final pixel value from float3 to uint:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 |
inline uint RGB32FtoRGB8( float3 c ) { int r = (int)(min( c.x, 1.f ) * 255); int g = (int)(min( c.y, 1.f ) * 255); int b = (int)(min( c.z, 1.f ) * 255); return (r << 16) + (g << 8) + b; } __kernel void render( __global uint* target, float3 camPos, float3 p0, float3 p1, float3 p2 ) { ... // trace the primary ray float3 color = Trace( &ray ); // plot the result target[x + y * SCRWIDTH] = RGB32FtoRGB8( color ); } |
We now have a kernel that renders a white window. But let’s look on the bright side: The white pixels are produced using correct primary rays, using a Trace function that should be easy to extend to what we had on the CPU.
Sky
The easiest way to prove that the primary rays are in fact correct is by converting the skydome functionality.
In the CPU code, in ::Tick, we create a new buffer:
|
1 2 |
skyData = new Buffer( skyWidth * skyHeight * 3 * sizeof( float ), skyPixels ); skyData->CopyToDevice(); |
The buffer takes the existing HDR skydome, stored in skyPixels. The second line copies the data to the GPU. We need to do this only once: Skydomes generally don’t change over time, so there will be no per-frame data transport for this data. To use the data in the kernel we do however need to pass a pointer to it. We add an argument to the kernel:
|
1 2 3 4 |
__kernel void render( __global uint* target, float* skyData, float3 camPos, float3 p0, float3 p1, float3 p2 ) { ... } |
In the CPU code, in ::Tick, we pass the pointer to the data on the GPU device.
|
1 |
tracer->SetArguments( target, skyData, camPos, p0, p1, p2 ); |
To make the data accessible in the Trace function, we simply pass it along:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 |
float3 Trace( struct Ray* ray, float* skyPixels ) { ... } __kernel void render( __global uint* target, float* skyPixels, float3 camPos, float3 p0, float3 p1, float3 p2 ) { ... // trace the primary ray float3 color = Trace( &ray, skyPixels ); // plot the result target[x + y * SCRWIDTH] = RGB32FtoRGB8( color ); } |
Inside Trace, we now have a ray, and skyPixels. We can test both by pretending that we never hit any geometry; this should yield a rotating skydome without teapots. Code:
|
1 2 3 4 5 6 7 8 9 |
float3 Trace( struct Ray* ray, float* skyPixels ) { // sample sky float phi = atan2( ray->D.z, ray->D.x ); uint u = (uint)(3200 * (phi > 0 ? phi : (phi + 2 * PI)) * INV2PI - 0.5f); uint v = (uint)(1600 * acos( ray->D.y ) * INVPI - 0.5f); uint skyIdx = (u + v * 3200) % (3200 * 1600); return 0.65f * (float3)(skyPixels[skyIdx * 3], skyPixels[skyIdx * 3 + 1], skyPixels[skyIdx * 3 + 2]); } |
That requires some explanation. First of all, the size of the skydome texture is now hardcoded. To make that more generic, we really should really pass skyWidth and skyHeight along with the skyPixels pointer. Secondly, INVPI and INV2PI are conveniently used from common.h: Including it pays off again. And finally: some math functions have different names in OpenCL. Here: atan2f and acosf become atan2 and acos. Some functions also have slightly different behavior. Here: atan2 can yield a negative number that we need to handle.
The kernel now produces a considerably more interesting image:
BVH Data
Let’s take a step back, and see what other data we need on the GPU to complete the Trace function.
- A single mesh, with Tri and TriEx arrays.
- A BVH for the mesh, consisting of a triangle index array and a set of nodes.
- A texture: again, just one.
- An array of BVHInstances, which includes transforms.
- A TLAS, with instance indices in the leaves.
To get all this data to the GPU, we start with creating seven buffers, using the data we have created on the CPU side, at the end of the ::Init() function:
|
1 2 3 4 5 6 7 8 |
triData = new Buffer( 1024 * sizeof( Tri ), mesh->tri ); triExData = new Buffer( 1024 * sizeof( TriEx ), mesh->triEx ); Surface* tex = mesh->texture; texData = new Buffer( tex->width * tex->height * sizeof( uint ), tex->pixels ); instData = new Buffer( 256 * sizeof( BVHInstance ), bvhInstance ); tlasData = new Buffer( tlas.nodesUsed * sizeof( TLASNode ), tlas.tlasNode ); bvhData = new Buffer( mesh->bvh->nodesUsed * sizeof( BVHNode ), mesh->bvh->bvhNode ); idxData = new Buffer( 1024 * sizeof( uint ), mesh->bvh->triIdx ); |
Buffers triData and triExData still assume the hardcoded triangle limit of 1024 triangles in the Mesh class, which is enough for the teapot object. If you already relaxed this constraint, you can just replace 1024 by the actual number of triangles in the mesh.
Once all the buffers are created, we copy them to the GPU:
|
1 2 3 4 5 6 |
triData->CopyToDevice(); triExData->CopyToDevice(); texData->CopyToDevice(); instData->CopyToDevice(); bvhData->CopyToDevice(); idxData->CopyToDevice(); |
One buffer is missing from this list: the tlasData buffer. This buffer changes every frame when we animate objects and update the TLAS; the call to CopyToDevice() will thus go to the ::Tick() function rather than the ::Init() function.
To get pointers to all this data in the kernel code, we pass all the pointers to the kernel, like we did with the skydome texture data:
|
1 2 3 4 5 |
tracer->SetArguments( target, skyData, triData, triExData, texData, tlasData, instData, bvhData, idxData, camPos, p0, p1, p2 ); |
In kernels.cl we update the definition of the render kernel:
|
1 2 3 4 5 6 7 8 9 10 |
__kernel void render( __global uint* target, __global float* skyPixels, __global struct Tri* triData, __global struct TriEx* triExData, __global uint* texData, __global struct TLASNode* tlasData, __global struct BVHInstance* instData, __global struct BVHNode* bvhNodeData, __global uint* idxData, float3 camPos, float3 p0, float3 p1, float3 p2 ) { ... } |
OpenCL allows us to pass pointers to undefined structs like Tri and TriEx, but to use these, we’ll need to properly define them. This however reveals an important complication.
Alignment
The problem is the float3 type. On the CPU side, we expect it to be 12 bytes in size; in fact, we rely on this, e.g. in the (rather important) TLASNode class, which fits snuggly in 32 bytes:
|
1 2 3 4 5 6 7 8 |
struct TLASNode { float3 aabbMin; uint leftRight; // 2x16 bits float3 aabbMax; uint BLAS; ... }; |
In OpenCL, things are different. A GPU can read 16 bytes from memory more quickly than 12 bytes, and therefore, a float3 is padded with a dummy field, to make it 16 bytes. To enforce the correct layout of the TLASNodes that we receive from the CPU, we can define the struct as follows:
|
1 2 3 4 5 6 7 |
struct TLASNode { float minx, miny, minz; uint leftRight; // 2x16 bits float maxx, maxy, maxz; uint BLAS; }; |
This struct is 32 bytes in size, and has the same layout as the CPU version. Copying data into this should thus yield the correct result.
The other structs that use float3 variables can be modified in a similar fashion:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 |
struct Tri { float v0x, v0y, v0z; float v1x, v1y, v1z; float v2x, v2y, v2z; float cx, cy, cz; }; struct TriEx { float2 uv0, uv1, uv2; float N0x, N0y, N0z; float N1x, N1y, N1z; float N2x, N2y, N2z; float dummy; }; struct BVHNode { float minx, miny, minz; int leftFirst; float maxx, maxy, maxz; int triCount; }; |
The last structure that we need is the BVHInstance class. The CPU version of this class has some data and methods that we don’t need for rendering on the GPU: on the GPU end we just need the two matrices it holds. Still, synchronization is easier if we keep the same layout, so we devise a struct that does the job:
|
1 2 3 4 5 6 7 8 |
struct BVHInstance { uint dummy1, dummy2; uint idx; float16 transform; float16 invTransform; // inverse transform uint dummy[6]; }; |
The variables dummy1 and dummy2 together take up 8 bytes, just like the BVH* in the original class. The last member variable, aabb bounds, is 6 times 4 bytes in size, just like the uint dummy array at the end of the proposed struct. Note the odd float16 data type: this will represent the mat4 we used on the CPU.
All data is in position, time to start rendering.
Tea Time
To render a teapot using ray tracing on the GPU, we need an OpenCL version of the BVH::Intersect function, which implements BLAS traversal. Once this works, we can convert TLAS traversal and shading, but all that will not produce anything verifiable without proper BLAS traversal. Let’s start with the CPU code, and tweak it until it works:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 |
void BVH::Intersect( Ray& ray, uint instanceIdx ) { BVHNode* node = &bvhNode[0], *stack[64]; uint stackPtr = 0; while (1) { if (node->isLeaf()) { for (uint i = 0; i < node->triCount; i++) { uint instPrim = (instanceIdx << 20) + triIdx[node->leftFirst + i]; IntersectTri( ray, mesh->tri[instPrim & 0xfffff /* 20 bits */], instPrim ); } if (stackPtr == 0) break; else node = stack[--stackPtr]; continue; } BVHNode* child1 = &bvhNode[node->leftFirst]; BVHNode* child2 = &bvhNode[node->leftFirst + 1]; float dist1 = IntersectAABB( ray, child1->aabbMin, child1->aabbMax ); float dist2 = IntersectAABB( ray, child2->aabbMin, child2->aabbMax ); if (dist1 > dist2) { swap( dist1, dist2 ); swap( child1, child2 ); } if (dist1 == 1e30f) { if (stackPtr == 0) break; else node = stack[--stackPtr]; } else { node = child1; if (dist2 != 1e30f) stack[stackPtr++] = child2; } } } |
The first change we make is the function header. OpenCL is not an object-oriented language, so BVH::Intersect will have to be a function at the global scope. We also need to make changes to the data we pass into it: The Ray& must either become a struct Ray*, or simply a struct Ray (i.e., by value). In this case the pointer is the better option, primarily because we will change the ray that is handed to us. And finally, we’ll need some additional data that we accessed via the BVH class so far: The triangle data, and of course the BVH data, consisting of nodes and triangle indices.
|
1 2 3 4 5 |
void BVHIntersect( struct Ray* ray, uint instanceIdx, struct Tri* tri, struct BVHNode* bvhNode, uint* triIdx ) { ... } |
After adding the struct keyword on lines 3, 17 and 18 (before each BVHNode* declaration), line 7 shows up in the error log. Because strict methods are not allowed in OpenCL, we replace the condition by:
|
1 |
if (node->triCount > 0) // isLeaf() |
Next up is the call to IntersectTri(...). This function doesn’t exist yet, so let’s add it and convert it to OpenCL:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 |
void IntersectTri( struct Ray* ray, struct Tri* tri, const uint instPrim ) { float3 v0 = (float3)(tri->v0x, tri->v0y, tri->v0z); float3 v1 = (float3)(tri->v1x, tri->v1y, tri->v1z); float3 v2 = (float3)(tri->v2x, tri->v2y, tri->v2z); float3 edge1 = v1 - v0, edge2 = v2 - v0; float3 h = cross( ray->D, edge2 ); float a = dot( edge1, h ); if (a > -0.00001f && a < 0.00001f) return; // ray parallel to triangle float f = 1 / a; float3 s = ray->O - v0; float u = f * dot( s, h ); if (u < 0 || u > 1) return; float3 q = cross( s, edge1 ); float v = f * dot( ray->D, q ); if (v < 0 || u + v > 1) return; float t = f * dot( edge2, q ); if (t > 0.0001f && t < ray->hit.t) ray->hit.t = t, ray->hit.u = u, ray->hit.v = v, ray->hit.instPrim = instPrim; } |
Note that after obtaining the three float3‘s for the vertices (lines 3-5, highlighted), the code is quite straight-forward and close to the C++ original: Even the cross-product on line 7 is supported out-of-the-box by OpenCL. We call the IntersectTri from BVHIntersect, using a pointer to a single triangle:
|
1 |
IntersectTri( ray, &tri[instPrim & 0xfffff /* 20 bits */], instPrim ); |
One more remaining issue is now the AABB intersection, called from lines 19 and 20. Luckily we kept the non-SIMD version, which is easily converted to OpenCL. Original:
|
1 2 3 4 5 6 7 8 9 10 |
inline float IntersectAABB( const Ray& ray, const float3 bmin, const float3 bmax ) { float tx1 = (bmin.x - ray.O.x) * ray.rD.x, tx2 = (bmax.x - ray.O.x) * ray.rD.x; float tmin = min( tx1, tx2 ), tmax = max( tx1, tx2 ); float ty1 = (bmin.y - ray.O.y) * ray.rD.y, ty2 = (bmax.y - ray.O.y) * ray.rD.y; tmin = max( tmin, min( ty1, ty2 ) ), tmax = min( tmax, max( ty1, ty2 ) ); float tz1 = (bmin.z - ray.O.z) * ray.rD.z, tz2 = (bmax.z - ray.O.z) * ray.rD.z; tmin = max( tmin, min( tz1, tz2 ) ), tmax = min( tmax, max( tz1, tz2 ) ); if (tmax >= tmin && tmin < ray.hit.t && tmax > 0) return tmin; else return 1e30f; } |
And the OpenCL version, which uses the ‘three-float notation’ for the extends of the bounding box stored in a BVHNode:
|
1 2 3 4 5 6 7 8 9 10 |
float IntersectAABB( struct Ray* ray, struct BVHNode* node ) { float tx1 = (node->minx - ray->O.x) * ray->rD.x, tx2 = (node->maxx - ray->O.x) * ray->rD.x; float tmin = min( tx1, tx2 ), tmax = max( tx1, tx2 ); float ty1 = (node->miny - ray->O.y) * ray->rD.y, ty2 = (node->maxy - ray->O.y) * ray->rD.y; tmin = max( tmin, min( ty1, ty2 ) ), tmax = min( tmax, max( ty1, ty2 ) ); float tz1 = (node->minz - ray->O.z) * ray->rD.z, tz2 = (node->maxz - ray->O.z) * ray->rD.z; tmin = max( tmin, min( tz1, tz2 ) ), tmax = min( tmax, max( tz1, tz2 ) ); if (tmax >= tmin && tmin < ray->hit.t && tmax > 0) return tmin; else return 1e30f; } |
The final problem in BVHIntersect is a bit silly: OpenCL doesn’t have a ‘swap’ function. So, when we find that the second child is closer than the first, we swap their pointers and distances manually:
|
1 2 3 4 5 |
if (dist1 > dist2) { float d = dist1; dist1 = dist2; dist2 = d; struct BVHNode* c = child1; child1 = child2; child2 = c; } |
The full BVHIntersect function, ported to OpenCL:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 |
void BVHIntersect( struct Ray* ray, uint instanceIdx, struct Tri* tri, struct BVHNode* bvhNode, uint* triIdx ) { struct BVHNode* node = &bvhNode[0], *stack[32]; uint stackPtr = 0; while (1) { if (node->triCount > 0) // isLeaf() { for (uint i = 0; i < node->triCount; i++) { uint instPrim = (instanceIdx << 20) + triIdx[node->leftFirst + i]; IntersectTri( ray, &tri[instPrim & 0xfffff /* 20 bits */], instPrim ); } if (stackPtr == 0) break; else node = stack[--stackPtr]; continue; } struct BVHNode* child1 = &bvhNode[node->leftFirst]; struct BVHNode* child2 = &bvhNode[node->leftFirst + 1]; float dist1 = IntersectAABB( ray, child1 ); float dist2 = IntersectAABB( ray, child2 ); if (dist1 > dist2) { float d = dist1; dist1 = dist2; dist2 = d; struct BVHNode* c = child1; child1 = child2; child2 = c; } if (dist1 == 1e30f) { if (stackPtr == 0) break; else node = stack[--stackPtr]; } else { node = child1; if (dist2 != 1e30f) stack[stackPtr++] = child2; } } } |
All done. We can call this function from Trace, to see if the rather massive amount of new code and data actually works. At the start of the Trace function we call the new BVHIntersect function:
|
1 2 3 |
// see if we hit a teapot BVHIntersect( ray, 0, triData, bvhNodeData, idxData ); if (ray->hit.t < 1e30f) return (float3)(1,1,1); |
This should produce white teapot, without any shading, and skydome pixels for all rays that do not hit the teapot. And lo and behold!:
TODO
Now, we’re obviously not done. It’s just one teapot, it is not animated, there’s no reflection and no texture. But really, that is for another day, the article is long enough as it is. But from here things should go smooth: OpenCL is now under our command, and the data is in position. Next time we really complete the project, with some large-scale instanced scenery.
Closing Remarks
The complete code for this project, in ten convenient incremental stages, can be found on GitHub (for Windows / Visual Studio).
Enjoyed the article? Spotted an error?
Ping me on Twitter (@j_bikker), comment on this article, or send an email to bikker.j@gmail.com.
You can also find me on Github: https://github.com/jbikker.
Or read some more articles: https://jacco.ompf2.com. Some favorites:
- Wavefront path tracing
- Probability Theory for Physically Based Rendering – part 1
- Probability Theory for Physically Based Rendering – part 2
- OPT#3: SIMD (part 1 of 2) – also check the other optimization blog posts!
