Rendering CUDA

Kernel form

render-cuda accepts a program in kernel form: every computation sits inside a global loop at the top level, and each kernel uses only buffers, literals and its own ranges. A value computed outside a kernel would live on the CPU, and the GPU couldn't see it. A program that isn't in kernel form is still valid: run and render take it, and only render-cuda exits 2.

One loop becomes one kernel

Take the reverse from lesson 3, 2 blocks of 4. Here is what the reference prints for it, without the helper functions at the top:

__global__ void k0(int32_t *__restrict__ b0, int32_t *__restrict__ b1) {
  int32_t r6 = blockIdx.x;
  __shared__ int32_t a7[4];
  for (int32_t z_ = threadIdx.x; z_ < 4; z_ += blockDim.x) a7[z_] = 0;
  __syncthreads();
  int32_t v8 = (int32_t)((uint32_t)r6 * (uint32_t)4);
  {
    int32_t r9 = threadIdx.x;
    int32_t v11 = (int32_t)((uint32_t)v8 + (uint32_t)r9);
    int32_t v13 = b0[v11];
    a7[r9] = v13;
  }
  __syncthreads();
  {
    int32_t r16 = threadIdx.x;
    ...
    int32_t v22 = a7[v20];
    b1[v17] = v22;
  }
  __syncthreads();
}

void kernel(int32_t *b0, int32_t *b1) {
  UIR_LAUNCH(k0, 2, 4, b0, b1);
}

Line by line:

• __global__ marks a function the CPU launches and the GPU runs, once per thread. Its parameters are the buffers the kernel uses, which you find the same way as week 7's free names.
• The global loop's index is blockIdx.x. There's no loop: the GPU starts every block.
• The shared alloc is a __shared__ array. The threads clear it together, each taking every blockDim.x-th element, and wait.
• Code in the global loop but outside a local loop, like v8, runs on every thread. That's fine, since it only computes values. Stores outside local loops are forbidden for exactly this reason.
• Each local loop becomes a block of code where the index is threadIdx.x, followed by a barrier. The braces keep one phase's names apart from the next.
• The host function kernel takes the program's buffers, which are already on the GPU, and launches each kernel in program order with UIR_LAUNCH(name, blocks, threads, args...).

Helpers like f32_max get __device__, so GPU code can call them. Everything between the braces is the C your week 7 renderer already writes for a loop body.

UIR_LAUNCH

Your output never writes the <<<blocks, threads>>> launch syntax itself. UIR_LAUNCH comes from the grader, and that one choice is what lets the same file run in two places:

• With real CUDA, it's k0<<<2, 4>>>(b0, b1). The grader builds with nvcc, copies the buffers to the GPU, calls your kernel, and copies them back.
• In the emulator, it's a loop over blocks. Each block's threads run as real threads, __shared__ becomes a static array they share, and __syncthreads() is a real barrier.

A program with several global loops launches several kernels, one after another. Each launch finishes before the next starts, so a later kernel may read anything an earlier one wrote. That's the barrier between blocks that lesson 1 said a kernel doesn't have. c_two_kernels checks it.

For the CPU, nothing new

render still has to accept these programs. A global loop renders like week 7's thread loop: an outlined function run by uir_parallel_for, with the shared memory declared inside it so each block gets its own, zeroed. Local loops are plain loops, so phases are consecutive loops and need no barrier. The tests r_reverse_shared_cpu and r_gemm_tile_per_block_cpu check that.

Bits

Every thread does the same operations, in the same order, as the interpreter does for that iteration. The C you emit is unchanged. Two compiler flags keep the GPU's arithmetic exact: -fmad=false stops nvcc from fusing a multiply and an add, and -prec-div=true keeps division correctly rounded. The math helpers from week 6 are written with +, *, / and comparisons, so they come out the same too.