Difficulty: Intermediate | Prerequisites: C/C++ pointers, basic understanding of CPU architecture
CUDA is NVIDIA's parallel computing platform that lets you run code on a GPU instead of (or alongside) the CPU. GPUs contain hundreds or thousands of small cores, each able to execute a thread. Where a CPU handles a few threads very quickly, a GPU handles thousands of threads simultaneously, which makes it suited to problems where the same operation is applied to many data elements at once. This topic sits at the intersection of systems programming and high-performance computing. You should already be comfortable with C pointers, memory allocation, and function declarations before diving in.
CUDA extends C/C++ so you can write functions (kernels) that run in parallel across GPU threads. You allocate memory on the GPU, copy data from the CPU to the GPU, launch a kernel, then copy results back. The programming model revolves around threads organised into blocks, and blocks organised into grids.
Kernel
A function written in CUDA that executes on the GPU. Declared with the global qualifier. When launched, it runs simultaneously across many threads. Think of it as a single set of instructions that every thread follows, each on its own slice of data.
Thread
The smallest unit of execution on the GPU. Each thread has a unique index (threadIdx.x) within its block, which you use to figure out which piece of data that thread should work on.
Block (thread block)
A group of threads that execute together and can share fast on-chip memory. You set the block size when you launch the kernel. Maximum block size is typically 1024 threads.
Grid
The collection of all blocks launched for a single kernel call. A grid can contain many blocks. In the simplest case, a grid has just one block.
Host
The CPU and its associated system memory. Your main() function and regular C/C++ code run on the host.
Device
The GPU and its associated memory. Kernel code runs on the device. Device memory is separate from host memory, so data must be explicitly copied between them.
cudaMalloc
Allocates memory on the GPU (device memory). Analogous to malloc on the CPU. Returns a pointer to device memory that cannot be dereferenced from host code.
cudaMemcpy
Copies data between host and device memory. You specify the direction: cudaMemcpyHostToDevice (CPU to GPU) or cudaMemcpyDeviceToHost (GPU to CPU). In simple terms, this is how you shuttle data back and forth across the bus.
cudaFree
Frees memory previously allocated on the device with cudaMalloc. Same principle as free() in standard C.
cudaDeviceSynchronize
Blocks the host (CPU) until the GPU has finished executing the current kernel. Without this, the CPU would carry on immediately and might try to read results before the GPU has finished writing them.
cudaGetLastError
Returns the error status of the most recent CUDA runtime call. Used immediately after a kernel launch to catch configuration errors (wrong block size, out-of-memory, etc.).
Stride loop
A pattern where each thread processes multiple elements by stepping through the array in increments equal to the total number of threads. This lets a small number of threads cover a large array.
blockDim.x
A built-in CUDA variable that tells each thread how many threads are in its block (the block size in the x-dimension).
threadIdx.x
A built-in CUDA variable giving the thread's index within its block, starting from 0.
The CPU (host) and GPU (device) have separate memory spaces.
You cannot directly read or write GPU memory from CPU code, or vice versa.
The typical workflow is: allocate device memory, copy input data from host to device, run the kernel, copy results from device to host, free device memory.
A kernel is declared with global void functionName(...).
It returns void and is called from host code using the triple-angle-bracket syntax: functionName<<.
The first number inside <<<>>> is the number of blocks in the grid.
The second number is the number of threads per block.
Example: addKernel<<<1, 256>>>(dev_c, dev_a, dev_b, size) launches one block of 256 threads.
Inside a kernel, threadIdx.x gives the thread's position within its block (0 to blockDim.x - 1).
With a single block, threadIdx.x is sufficient to index into your data.
With multiple blocks, the global index is typically blockIdx.x * blockDim.x + threadIdx.x.
When the array is larger than the block size, each thread handles more than one element.
The thread starts at index threadIdx.x and steps forward by blockDim.x (the stride) on each iteration.
This ensures all elements are covered without needing the block size to match the array size.
Formula: for (int i = threadIdx.x; i < size; i += blockDim.x)
Nearly every CUDA runtime function returns a cudaError_t value.
You compare this against cudaSuccess to check for failure.
After a kernel launch, call cudaGetLastError() to detect launch-configuration errors.
After that, call cudaDeviceSynchronize() to catch errors that occurred during kernel execution.
Robust CUDA code checks the return value of every API call.
Stride loop index calculation:
thread's elements = { threadIdx.x, threadIdx.x + blockDim.x, threadIdx.x + 2*blockDim.x, ... }
For a block of 8 threads processing an array of 5000 elements, thread 0 handles indices 0, 8, 16, 24, ..., 4992. Thread 7 handles indices 7, 15, 23, ..., 4999.
Kernel launch configuration:
kernel<<GPU-accelerated vector operations are the building blocks of scientific simulation, machine learning training, image processing, and financial modelling. The same allocate-copy-launch-copy-free pattern you see in a toy vector addition scales up to neural network layers operating on tensors with millions of elements.
Students often think GPU memory and CPU memory are shared. They are not. You must explicitly copy data between them using cudaMemcpy.
Students sometimes assume a kernel launch is synchronous (blocks until done). It is not. The CPU continues immediately, which is why cudaDeviceSynchronize exists.
A block size of 1024 threads is the hardware maximum, not a recommendation. Smaller block sizes are common and sometimes perform better depending on the workload.
Forgetting to check cudaGetLastError() after a kernel launch does not mean the kernel succeeded. Errors are silent unless you ask for them.
⚠️ The allocate-copy-launch-copy-free lifecycle is a staple exam question. Be able to list the steps in order.
⚠️ Know the difference between cudaMemcpyHostToDevice and cudaMemcpyDeviceToHost and when each is used.
⚠️ Understand why cudaDeviceSynchronize is necessary and what happens if you skip it.
⚠️ Be ready to trace through a stride loop and determine which elements a given thread processes.
True or false: A kernel function can return an integer value to the host. False. Kernels are declared void and return nothing. Results are written to device memory and copied back.
Fill in the blank: To copy data from the CPU to the GPU, the direction flag is ______. cudaMemcpyHostToDevice
True or false: threadIdx.x is unique across all blocks in the grid. False. threadIdx.x is unique only within a single block. For a global index across blocks, you combine blockIdx.x * blockDim.x + threadIdx.x.
Fill in the blank: The maximum number of threads per block on most NVIDIA GPUs is ____. 1024
Q: What is the purpose of cudaMalloc and how does it differ from standard malloc?
A: cudaMalloc allocates memory on the GPU (device). Unlike malloc, the returned pointer lives in device address space and cannot be dereferenced by host code. Data must be moved to and from it via cudaMemcpy.
Q: A kernel is launched with <<<1, 64>>> on an array of 1000 elements using a stride loop. How many elements does thread 0 process?
A: Thread 0 processes elements at indices 0, 64, 128, ..., 960. That is ceil(1000/64) = 16 elements.
Q: Why is cudaGetLastError() called immediately after a kernel launch, and why is cudaDeviceSynchronize() called after that?
A: cudaGetLastError() catches configuration errors in the launch itself (e.g., invalid block size). cudaDeviceSynchronize() waits for the kernel to finish and catches runtime errors that occur during execution on the GPU.
Q: Describe the five-step CUDA workflow for processing data on the GPU.
A: (1) Allocate device memory with cudaMalloc. (2) Copy input data from host to device with cudaMemcpy. (3) Launch the kernel. (4) Copy results from device to host with cudaMemcpy. (5) Free device memory with cudaFree.
This material connects directly to parallel algorithm design, where you consider how to decompose a problem so many threads can work on it simultaneously. It also ties into computer architecture (understanding why GPUs have many small cores vs. a CPU's few powerful ones) and operating systems concepts around memory spaces and DMA transfers.
CUDA, GPU programming, parallel computing, kernel function, global, threadIdx, blockDim, blockIdx, cudaMalloc, cudaMemcpy, cudaFree, cudaDeviceSynchronize, cudaGetLastError, host-device model, stride loop, grid-stride loop, vector addition, GPGPU, NVIDIA CUDA, thread block, warp, SIMT, device memory, host memory, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, HPC, high-performance computing