CUDA Ray Tracer

Abhi S.

Noah K.

Jared G.

What is CUDA?

• CUDA (Compute Unified Device Architecture) is a parallel computing platform and API by NVIDIA.
• API for using GPUs to process tasks that were traditionally handled by CPUs, significantly speeding up computations.
• Used for GPGPU, using the GPU for applications other than graphics processing.

Ray Tracing

• Fire a ray from the camera position through a pixel in the near-plane.
• Collide the ray with every sphere in the scene, save the minimum distance and the center of that sphere.
• Accumulate the color of the sphere multiplied by the light contribution.
• Compute the normal by subtracting collision position with center of collided sphere.
• Reflect ray across normal, repeat collision and accumulation until max_iter or if missing all spheres. Add skybox contribution on miss.
• Divide accumulated color by number of bounces.

https://www.scratchapixel.com/lessons/3d-basic-rendering/introduction-to-ray-tracing/implementing-the-raytracing-algorithm.html

Ray Tracing

• Fire a ray from the camera position through a pixel in the near-plane.
• Collide the ray with every sphere in the scene, save the minimum distance and the center of that sphere.
• Accumulate the color of the sphere multiplied by the light contribution.
• Compute the normal by subtracting collision position with center of collided sphere.
• Reflect ray across normal, repeat collision and accumulation until max_iter or if missing all spheres. Add skybox contribution on miss.
• Divide accumulated color by number of bounces.

miss!

1

2

3

[ Camera ]

Comparison

Serial

CUDA

Performance with Varying Sphere Count

Performance with Varying Sphere Count

CUDA Memory Management

CUDA offers a "Unified memory" abstraction that syncs device and host address space.

Our problem was simple enough that we could manage device memory manually.

Steps:

1. Allocate input/output buffers on device and host
1. Inputs: Spheres, Materials, Camera, etc.
2. Outputs: Framebuffer (pixel data)

For each frame:

• Memcpy input data (host -> device)
• Run Kernel
• Call cudaDeviceSynchronize(); to block CPU until kernel finishes
• Memcpy framebuffer (device -> host)
• Render it!

CUDA Abstraction

• Threads are the smallest unit of execution.
• Threads are grouped into blocks.
• All the blocks form the grid.

​

Threads < Thread-Blocks < Grid

CUDA Architecture

• Blocks are executed on a Streaming Multiprocessor(SM).
• The device is a grid made up of many SMs.
• Threads on an SM share a L1 cache.
• All threads share a global L2 cache.

​

​

​

​

​

[Global L2 cache]

SM 1

​

[L1 cache]

SM 2

​

[L1 cache]

SM 3

​

[L1 cache]

SM 4

​

[L1 cache]

CUDA Architecture

In hardware, SMs are implemented with warps.

​

​

​

​

​

​

​

​

​

​

​

[SM 1]

W0

​

​

​

​

​

​

​

​

W2

​

​

​

​

​

​

​

​

W3

​

​

​

​

​

​

​

​

W1

​

​

​

​

​

​

​

​

T0

T1

T2

T3

...

T0

T1

T2

T3

...

T0

T1

T2

T3

...

T0

T1

T2

T3

A warp is a set of 32 threads within a block.

- All threads in a warp execute the same instruction.

- Context switches are done for whole warps.

- Threads in a warp that branch differently become serial.

(thread-divergent branches)

​

...

...

Thread-Divergent Branches

if (threadIdx.x % 2)

a();

else

b();

c();

​

a and b won't run in parallel on threads in the same warp.

Since the threads have diverged, c effectively runs in half the speed!!

Thread-divergent branches are bad for parallelism.

CUDA Kernel Configuration

• 1 Thread per pixel
• 1 Ray per thread
• No dependencies between rays.
• 50 Bounces per ray for reflections
• Ends early if ray “misses”
• All bounces are calculated sequentially.

​

• Want to group threads into blocks such that all threads in a block have similar runtime, to avoid thread-divergent branches.

​

​

miss!

1

2

3

[ Camera ]

CUDA Kernel Configuration

CUDA Kernel Configuration

Observation:

​

The orange square will have same runtime of 0 bounces.

CUDA Kernel Configuration

Observation:

​

Pixels within a square will generally have a similar # of bounces.

CUDA Kernel Configuration

Observation: Nearby pixels will

• Have similar number of bounces (runtime).
• Access nearby memory in the framebuffer and accumulation buffer.
• Better utilize SM's L1 cache.

Warps are assigned 32 threads,

So we make blocks of 8x8 chunks of pixels.

​

dim3 blocks(width / 8, height / 8);

dim3 threads(8, 8);

raytraceKernel<<<blocks, threads>>>(...);

Scatter

​

No scatter

Too perfect

With scatter

Too noisy

Accumulation

With Accumulation

Just right :)

Frame 1

Frame 2

Frame 3

+

+

+

Frame ...

No Accumulation

Too noisy

Ray Tracing

• Fire a ray from the camera position through a pixel in the near-plane.
• Collide the ray with every sphere in the scene, save the minimum distance and the center of that sphere.
• Accumulate the color of the sphere multiplied by the light contribution.
• Compute the normal by subtracting collision position with center of collided sphere.
• Reflect ray across normal, repeat collision and accumulation until max_iter or if missing all spheres. Add skybox contribution on miss.
• Divide accumulated color by number of bounces.

https://www.scratchapixel.com/lessons/3d-basic-rendering/introduction-to-ray-tracing/implementing-the-raytracing-algorithm.html

Collisions and Light Contribution

Light On intersection the ray takes a dot product between the color to calculate the light contribution

Specular after the intersect a new ray is cast from the direction reflected from the normal

Scattering a random uniformly distributed vector is added to the the subsequent ray direction according to the roughness

Intersect take the dot product of the difference between the ray origin and the sphere center. With the radius use pythagorean and subtract and add.

https://www.lighthouse3d.com/tutorials/maths/ray-sphere-intersection/

​

https://astrobites.org/2020/10/22/specular-reflections-titan/