Welcome to my little cuda ray tracing tutorial, and first  a warning: Ray tracing is both fun and contagious, there is a fair chance that you will end up coding different variants of your ray tracer just to see those beautiful images. So if you value your spare time stop reading

I am not an ray-tracing expert, and i also value my spare time. But i really like to synthisize images on my computer, and i cannot explain why, other than it is really cool what a few lines of code can produce. But in ray tracing, it often takes a long time to create a decent result. There is of course a lot you can do to speed up thing:

1. first buy a faster computer, there more cores the better.
2. implement an acceleration structure.
3. dive into the world of optimizations, simd stuff, cache coherency ...
4. And if you are really talented you might end up doing some hairy math, that will sent your ray tracer into warp speed.

Or if you are like me, use the graphics card to speed up things a bit. This little tutorial will not state that the GPU thing is the only way to go or try to preach something. But it will hopefully help you to get started using CUDA to accelerate ray tracing. I have been writing a lot of shaders for other applications so it was very natural for me to experiment with a GPU implementation. So the natural choice would be to use Nvidia's CG or glsl. But i decided to try Cuda, because i wanted to learn what kind of beast it is and second it should be easier to code. For most parts cuda is a quite pleasant experience, for instance you can make emulation debug and print out values!! But again some times it crashes for real due to some invalid pointer, leaving you with a smoldering Vista-XP blue screen of death.

But let get started. This tutorial will focus on the CUDA stuff like how do you upload some triangles to your graphics card and how to implement a very simple ray tracing algorithm that runs on the GPU. It will be assumed you know Open GL, and some theory about ray tracing.

Anyways just to initialize some ray tracing knowledge pointer in your brain here is the ray tracing 101 (if you already is familiar with this GOTO 10):

To create a ray traced image like this:

1. for each pixel create a ray.
2. follow the ray into the scene and find the closest object that is intersected.
3. if the object is made of glass, mirror  or something similar find the reflected/refracted ray and trace on.
4. else computer the color of the pixel based on light sources and material information.

label : 10 // you made it this far: 100 xp award
Ok lets code:

The code has been written with one main design criteria. Keep it pretty simple! and dont include a bunch of unnecessary stuff. By my opinion a lot of tutorials i have seen, forces people to include complete frameworks and i think it makes the code less comprehensive. So the code is only one  .cpp (main.cpp) and one .cu (raytracer.cu) source file.

The ray tracer will be very simple, and it will be left as an exercise to include more advanced features like refraction, acceleration structures and cool materials. I will focus on sending a ray into a loaded triangulated model, computing the first hit position and do simple light calculations.

Above is a youtube video of the tutorial. Compared with the first image there is of course a lot of things missing, but hey its a start

To start we need to load some geometry and we use the .obj format to load the 3d model shown in the video above. After loading the geometry we need to send the triangle information to the graphic-card memory before we can start tracing rays. To do this we use a CUDA 1D texture that is declared in the top of ray-tracer.cu like this:

texture'<'float4, 1, cudaReadModeElementType'>' triangle_texture;

The CUDA texture memory provides a default caching so it is a bit faster than using global memory. The triangles is simple packed into the 1D texture in a way that each float4 contains the either the first triangle-vertex or two triangle edges.

Like this: ( float4(vertex.x,vertex.y,vertex.z, 0), float4 (egde1.x,egde1.y,egde1.z,0),float4 (egde2.x,egde2.y,egde2.z,0)) for each triangle.

This information is copied into a cuda device pointer and bound to the declared texture. The texture lookup is set to POINT mode to avoid linear interpolation.

The code for binding the device pointer to the 1D texture

void bindTriangles(float *dev_triangle_p, unsigned int number_of_triangles)
{
triangle_texture.normalized = false;                      // access with normalized texture coordinates
triangle_texture.filterMode = cudaFilterModePoint;        // Point mode, so no
triangle_texture.addressMode[0] = cudaAddressModeWrap;    // wrap texture coordinates
 
size_t size = sizeof(float4)*number_of_triangles*3;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();
cudaBindTexture(0,triangle_texture,dev_triangle_p,channelDesc,size);
}

After uploading and binding triangle information, we need a way to communicate between openGL and CUDA, because we use openGL to visualize the ray tracing result. Fortunately this is easily done by using a PBO (pixel buffer object), that serves as communication between the two. If you have the CUDA SDK this is described in the openGL PostProcess example.

Now the fun part begins... ray tracing the triangles. So the first problem is spawning the rays. Here we create a virtual camera and by setting a few parameters rotation, distance to origo and elevation, we can calculate the a vector basis for the camera. Given this basis we can compute the ray start position and direction. The camera calculation is done in the updateCamera() function.

Now we have a ray(a start position and a direction) for each pixel, and the we would like to know the closest triangles hit if any. But before that we do a little optimization that is equivilant to frustum culling. By calculating a axis aligned bounding box that encapsulates the entire scene, we first check wether the ray intersect this rectangle. If not we don't have to worry and just set the pixel color to the background value.

The ray tracing is all done in the CUDA kernel found in raytracer.cu file

__global__ void raytrace(...)

Next stop ray triangle intersection. To compute the ray triangle intersection is a science by itself because this is where all the action is going on. So a slight optimization will really show off. For our tutorial we use the method developed by Thomas Muller and Ben Trumbore, which is very fast and easy to implement in CUDA.

To find the closest triangle hit point we run linear through all the triangles in our scene and compute the ray triangle intersection value. To get the smallest intersection value we store the smallest t value yet found in a hit-record. Below is the code that calculates the ray triangle intersection values and finds the closest.

for(int i = 0; i < number_of_triangles; i++)
 {
 float4 v0 = tex1Dfetch(triangle_texture,i*3);
 float4 e1 = tex1Dfetch(triangle_texture,i*3+1);
 float4 e2 = tex1Dfetch(triangle_texture,i*3+2);
 
 float t = RayTriangleIntersection(r, make_float3(v0.x,v0.y,v0.z),make_float3(e1.x,e1.y,e1.z), make_float3(e2.x,e2.y,e2.z));
 
 if(t < hit_r.t && t > 0.001)
 {
 hit_r.t = t; // found a closer hit! store it.
 hit_r.hit_index = i;
 }

If we hit something we calculate a face normal using a cross product, and given a light position and color we do a Phong lighting computation.

Shadows and reflections.

By using ray tracing it is very easy to calculate shadows and reflections, so this should of course be included in the tutorial. The way we determine if a pixel is in the shadow is done by creating a new ray starting from the triangle hit-point and pointing towards the light source. by intersecting all triangles like we did before we can determine if there is triangles blocking the path to the light. If yes we multiply the color by some value < 1.0 to darken the color.

Reflection is also simple to do, in a CPU ray tracer you could create a recursive call, but here we just use a while loop to follow the ray through the scene. We employ a little hack to create a reflective object in the center of the scene. This is achieved by storing a extra value in the triangle float4 values, because only the first three is used we can store extra information in the last component.

Suggested improvement.

As stated before there is a some of improvements that need to be done, before this little ray tracer can produce nice pictures. Here is some suggestions that you might want to play with.

1. Refraction, transparent materials is the core of ray tracing, impress your mom by rendering household equipment!
2. Anti-aliasing, one ray pr pixel is not enough to produces decent images, experiment with random and jittered sampling.
3. Acceleration structures, there is a lot of speed to be found by implementing some sort of tree structure (KD trees, Bounding volumes, octtrees)
4. More advanced materials.
5. write to me and i will add your improvement to the list

I really hope that you enjoy this little tutorial to start doing some ray tracing, and email me anytime if you have any ideas that might improve this tutorial or questions.

peter.trier {at} alexandra(.)dk

By the way thanks for hanging on to the end 1000 xp bonus!

Source code download

Tutorial source and windows executable download cuda 2.3

Tutorial source and windows executable download_cuda_2.1

Links:

Cool ray tracing forum

CUDA development zone

http://www.cs.virginia.edu/~gfx/Courses/2003/ImageSynthesis/papers/Acceleration/Fast%20MinimumStorage%20RayTriangle%20Intersection.pdf