Difficulty: Intermediate | Prerequisites: CUDA Programming Fundamentals study notes, C/C++ pointers, familiarity with goto in C
This set of notes walks through a complete, working CUDA program that adds two integer vectors element-by-element on the GPU. The code also experiments with different block sizes (8, 16, 32, ..., 1024) to observe how thread configuration affects execution. If you have not yet reviewed the CUDA fundamentals notes covering kernels, memory management, and thread indexing, start there. This walkthrough assumes you understand those concepts and focuses on how they come together in real code.
The program initialises two arrays of 5000 integers on the CPU, sends them to the GPU, launches a kernel that adds them in parallel using a stride loop, copies the result back, and repeats the whole process at increasing block sizes from 8 up to 1024 (doubling each time). Every CUDA API call is wrapped in an error check.
block_seed
The starting block size in this program's experimentation loop. Set to 8. The loop then doubles the block size on each iteration (8, 16, 32, 64, 128, 256, 512, 1024).
cudaSetDevice
Selects which GPU to use when the system has more than one. cudaSetDevice(0) picks the first GPU. In simple terms, you are telling CUDA which graphics card to run your code on.
cudaDeviceReset
Destroys all allocations and resets the device state. Called to ensure profiling tools (like Nsight) capture complete trace data. Also cleans up between iterations in this code.
goto Error (cleanup pattern)
A C idiom for centralised error cleanup. When any CUDA call fails, execution jumps to the Error: label where all device memory is freed. This avoids duplicating cleanup code after every possible failure point. Think of it as a poor man's exception handler in C.
Includes: cuda_runtime.h and device_launch_parameters.h provide the CUDA API. Standard headers (stdio.h, iostream, fstream) handle I/O.
Forward declaration: addWithCuda is declared before main so the compiler knows its signature.
Kernel: addKernel is the global function that runs on the GPU.
main(): Sets up host data and calls addWithCuda.
addWithCuda(): Handles all device work: allocation, copying, launching, synchronisation, cleanup.
__global__ void addKernel(int *c, const int *a, const int *b, const int size)
{
int stride = blockDim.x;
for (int i = threadIdx.x; i < size; i += stride) {
c[i] = a[i] + b[i];
}
}
Each thread starts at its own index (threadIdx.x) and steps by the block size.
With a block of 8 threads and an array of 5000, thread 0 covers indices 0, 8, 16, ..., 4992.
The const qualifiers on a and b tell the compiler (and the reader) these are read-only inputs.
const int arraySize = 5000;
int a[arraySize], b[arraySize], c[arraySize] = { 0 };
for (int ii = 0; ii < arraySize; ii++) {
a[ii] = 2 * ii;
b[ii] = 0.6 * ii;
}
a gets even numbers: 0, 2, 4, 6, ...
b gets 0.6 * ii, but because b is an int array, the result is truncated to an integer. So b[1] = 0 (0.6 truncated), b[2] = 1 (1.2 truncated), b[3] = 1 (1.8 truncated), b[5] = 3 (3.0).
c is zero-initialised and will hold the output.
This function wraps the entire GPU workflow inside a loop that doubles the block size on each pass.
Device pointer declaration: dev_a, dev_b, dev_c are host-side pointers that will point to device memory.
Block size loop: for (block_size = block_seed; block_size <= 1024; block_size *= 2) iterates through block sizes 8, 16, 32, 64, 128, 256, 512, 1024.
Within each iteration:
cudaSetDevice(0) selects the GPU.
Three cudaMalloc calls allocate space for dev_a, dev_b, and dev_c on the device, each size * sizeof(int) bytes.
Two cudaMemcpy calls copy a and b from host to device (cudaMemcpyHostToDevice).
The kernel launches: addKernel<<<1, block_size>>>(dev_c, dev_a, dev_b, size).
One block, block_size threads.
The stride loop inside the kernel ensures all 5000 elements are processed regardless of how many threads there are.
cudaGetLastError() checks for launch errors.
cudaDeviceSynchronize() waits for the kernel to finish and checks for execution errors.
cudaMemcpy copies the result from dev_c back to host array c (cudaMemcpyDeviceToHost).
cudaDeviceReset() resets the device for the next iteration.
The Error: label frees all three device buffers with cudaFree.
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// ... more operations ...
Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
Every CUDA call checks its return value.
On failure, control jumps to Error:, which frees whatever was allocated.
cudaFree on a null pointer is safe (does nothing), so this works even if allocation failed partway through.
This pattern is common in C-based CUDA code. In C++, you would more likely use RAII or smart pointers.
The loop tests block sizes 8, 16, 32, 64, 128, 256, 512, and 1024.
All produce the same result (the addition is correct regardless of thread count), but execution time varies.
Smaller blocks mean each thread does more work (more iterations of the stride loop). Larger blocks mean more threads share the work.
In practice, block sizes of 128 or 256 often perform well on modern GPUs, but the optimal choice depends on the hardware and workload.
Elements per thread (stride loop):
elements_per_thread = ceil(arraySize / blockSize)
For arraySize = 5000:
Block size | Threads | Elements per thread (approx.) |
|---|---|---|
8 | 8 | 625 |
64 | 64 | ~79 |
256 | 256 | ~20 |
1024 | 1024 | ~5 |
Memory allocated on device per vector:
bytes = arraySize * sizeof(int) = 5000 * 4 = 20,000 bytes
Total device memory for all three vectors: 60,000 bytes (roughly 60 KB).
The pattern of "loop over block sizes and measure performance" is standard GPU benchmarking methodology. Engineers tuning deep learning kernels, physics simulations, or video encoding pipelines use the same approach: sweep a configuration parameter and profile each run to find the sweet spot for their particular hardware.
Students often think increasing the block size always makes the code faster. It does not. Beyond a certain point, GPU occupancy, register pressure, and memory bandwidth become the bottleneck, not thread count.
The goto Error pattern looks alarming if you have been told "never use goto." In C (especially systems and CUDA code), it is an accepted idiom for centralised cleanup. It is not the same as spaghetti code.
Students sometimes assume cudaDeviceReset() is only needed once at program exit. In this code, it is called inside the loop to reset state between block-size experiments, which is a valid (if aggressive) approach.
Calling cudaMalloc and cudaFree inside a loop is not free. Each call has overhead. Production code would allocate once and reuse buffers.
⚠️ Be able to trace through the kernel and determine the output for a given input and block size. The stride loop is the key mechanism.
⚠️ Know the order of operations inside addWithCuda: set device, allocate, copy to device, launch, check errors, synchronise, copy to host, free.
⚠️ Understand why cudaFree is safe to call on a null or already-freed pointer, and why the cleanup label works.
⚠️ If asked about the b[ii] = 0.6 * ii line, note the implicit float-to-int truncation. This is a subtle bug or design choice worth commenting on in an exam.
⚠️ The kernel uses <<<1, block_size>>>, meaning one block with varying thread counts. Be ready to explain why this still processes all 5000 elements (the stride loop).
True or false: The kernel addKernel<<<1, 8>>> with a stride loop can process an array of any size. True. The stride loop ensures each thread steps through the entire array, processing one element every blockDim.x indices.
Fill in the blank: In the expression b[ii] = 0.6 ii, when ii is 3, the value stored in b[3] is ____. 1 (because 0.6 3 = 1.8, which truncates to 1 when assigned to an int).
True or false: The Error: label in the cleanup section is only reached when an error occurs. False. In this code, cudaFree calls under the Error: label execute on every iteration of the loop, whether an error occurred or not, because there is no jump past them on the success path.
Q: Why does the program call cudaSetDevice(0) inside the loop rather than once before it?
A: Because cudaDeviceReset() is called at the end of each iteration, which destroys the CUDA context. The device must be re-selected afterwards. Placing cudaSetDevice inside the loop ensures a fresh context each time.
Q: What would happen if the kernel were launched as <<<1, 5000>>> instead of using a stride loop with a smaller block size?
A: It would fail. The maximum threads per block is 1024 on most GPUs. A block size of 5000 would cause a launch error, which cudaGetLastError() would catch.
Q: The program allocates device memory inside the loop and frees it at the Error: label. What performance issue does this create?
A: Repeated allocation and deallocation of device memory incurs overhead on each iteration. A more efficient approach would allocate once before the loop and free once after.
Q: Explain what happens, step by step, when addKernel<<<1, 8>>> is launched on an array of size 5000.
A: Eight threads are created in a single block. Thread 0 starts at index 0, thread 1 at index 1, and so on up to thread 7 at index 7. Each thread then increments its index by 8 (the stride) and processes the next element. This continues until all 5000 elements have been covered. Thread 0 handles indices 0, 8, 16, ..., 4992. The last thread to process index 4999 is thread 7 (since 4999 mod 8 = 7).
This code connects to software engineering practices around error handling and resource management, particularly the RAII pattern in C++ and how C code compensates without it. It also ties into performance engineering: choosing block sizes, understanding GPU occupancy, and profiling, topics you will encounter in parallel computing or HPC courses. The implicit type truncation in b[ii] = 0.6 * ii links back to type conversion rules in C/C++.
CUDA vector addition, addKernel, addWithCuda, cudaMalloc, cudaMemcpy, cudaFree, cudaSetDevice, cudaDeviceReset, cudaGetLastError, cudaDeviceSynchronize, stride loop, grid-stride loop, block size, thread block, kernel launch syntax, triple angle brackets, <<<>>>, GPU memory allocation, host to device copy, device to host copy, CUDA error handling, goto cleanup pattern, block size tuning, GPU benchmarking, implicit type conversion, float to int truncation, CUDA profiling, Nsight, Visual Profiler