Programming a ray tracer in CUDA Part 1.2
This is an intermediate chapter to show you the basic structure of a CUDA program for path tracing. It covers uploading the necessary data to the GPU (models, textures, etc.), rendering to a framebuffer, and downloading the result back to the CPU. Think of it as learning how to write custom fragment shaders in pure C, but running on the GPU.
This part is intended as both a short introduction to CUDA and a basic mental map of how path tracing will work, before we dive into the actual lighting math.
CUDA initialization
One of the nicest things about CUDA is that it initializes on demand with no boilerplate. We literally just need:
#include <cuda.h>
#include <cuda_runtime.h>
int main() {
// Start using CUDA immediately
}
That's it. Remember to compile with nvcc instead of gcc or clang.
Allocating a framebuffer
First we need a framebuffer: a flat array of bytes representing the color of each pixel. We'll use standard RGB (one byte per channel, packed continuously in memory).
int width = 1024;
int height = 1024;
int fbo_size = width * height * 3; // 3 bytes per pixel (R, G, B)
uint8_t* fbo_h = (uint8_t*)malloc(fbo_size); // CPU memory
Since we're rendering from the GPU, we also need GPU-accessible memory (VRAM). CUDA provides functions for allocating, freeing, and transferring GPU memory in a style similar to C's malloc and free. It's common practice to suffix GPU buffer names with _d (for device) to distinguish them from regular CPU memory.
uint8_t* fbo_d = nullptr;
cudaMalloc(&fbo_d, fbo_size); // GPU memory
Writing the kernel
A kernel is the entry point of a CUDA workgroup that we launch from the CPU. Once launched, the GPU assigns resources to execute it in parallel across many threads, then returns control to the CPU.
Thread grid layout
CUDA launches threads in a two-level hierarchy: threads are grouped into blocks, and blocks form a grid.
- A warp is the actual hardware execution unit (typically 32 threads) and always comes from a single block.
- We choose how many threads per block and how many blocks in the grid.
For rendering, the natural choice is one thread per pixel with an 8x8 block:
int numThreads = 8;
kernel<<<
dim3(width / numThreads, height / numThreads), // grid dimensions
dim3(numThreads, numThreads) // block dimensions
>>>(fbo_d, width, height);
CUDA kernels use the familiar C function call syntax, with the addition of the <<<grid, block>>> specifier between the function name and its arguments.
The kernel itself
__global__ void kernel(uint8_t* framebuffer, int width, int height) {
// Determine which pixel this thread is responsible for
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
// Guard against out-of-bounds threads (when image size is not
// a perfect multiple of block size)
if (x >= width || y >= height) return;
// Normalize to [0, 1]
float u = x / float(width);
float v = y / float(height);
// Write a simple gradient to the framebuffer
framebuffer[(y * width + x) * 3 + 0] = 255; // R
framebuffer[(y * width + x) * 3 + 1] = uint8_t(u * 255); // G
framebuffer[(y * width + x) * 3 + 2] = uint8_t(v * 255); // B
}
We get our 2D pixel coordinate by combining the block position in the grid (blockIdx) with the thread position within the block (threadIdx). The blockDim tells us how many threads are in each block dimension.
A cleaner alternative is to use a small struct for the color:
typedef struct {
uint8_t r, g, b;
} Color;
__global__ void kernel(Color* framebuffer, int width, int height) {
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
if (x >= width || y >= height) return;
float u = x / float(width);
float v = y / float(height);
framebuffer[y * width + x] = { 255, uint8_t(u * 255), uint8_t(v * 255) };
}
Much cleaner for the write side. The memory layout is identical either way.
Downloading the result and writing to disk
Once the GPU is done, we copy the framebuffer back to CPU memory and write it to disk:
// Wait for the kernel to finish, then copy result to CPU
cudaMemcpy(fbo_h, fbo_d, fbo_size, cudaMemcpyDeviceToHost);
// Write to PNG using the STB image write library
stbi_write_png("result.png", width, height, 3, fbo_h, width * 3);
// Free both CPU and GPU memory
free(fbo_h);
cudaFree(fbo_d);
The cudaMemcpyDeviceToHost flag tells CUDA to copy from GPU to CPU memory (as opposed to cudaMemcpyHostToDevice which goes the other way).
Putting it all together
Here's a complete minimal program that renders a gradient to a PNG file:
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdint.h>
#include <stdlib.h>
#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "stb_image_write.h"
typedef struct { uint8_t r, g, b; } Color;
__global__ void kernel(Color* fb, int width, int height) {
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
if (x >= width || y >= height) return;
float u = x / (float)width;
float v = y / (float)height;
fb[y * width + x] = { 255, (uint8_t)(u * 255), (uint8_t)(v * 255) };
}
int main() {
int width = 1024, height = 1024;
int num_pixels = width * height;
Color* fbo_h = (Color*)malloc(num_pixels * sizeof(Color));
Color* fbo_d = nullptr;
cudaMalloc(&fbo_d, num_pixels * sizeof(Color));
int tile = 8;
dim3 blocks(width / tile, height / tile);
dim3 threads(tile, tile);
kernel<<<blocks, threads>>>(fbo_d, width, height);
cudaMemcpy(fbo_h, fbo_d, num_pixels * sizeof(Color), cudaMemcpyDeviceToHost);
stbi_write_png("result.png", width, height, 3, fbo_h, width * 3);
free(fbo_h);
cudaFree(fbo_d);
return 0;
}
Compile with:
nvcc main.cu -o main
./main
# result.png contains a 1024x1024 gradient image
This is the foundation for everything we build from here. The demonstration in Part 1 was coded in C to keep it minimal. For the rest of the series I'll switch to C++ so we can have cleaner abstractions (the Scene class, proper vector math types, etc.) without having to fight with raw pointer arithmetic everywhere.
What's next
In the next part we will build on this foundation to add:
- A proper camera with configurable field of view and orientation
- Ray-sphere intersection tests
- A basic diffuse material using random sampling
- Accumulation across multiple frames for progressive rendering
The resulting image won't be anything fancy yet, but it will be the first version of our renderer that actually simulates how light bounces around the scene.