Introduction to the CUDA Programming Model

1. The Architectural Motivation: CPU vs. GPU

Before diving into the programming model, it's crucial to understand why we use GPUs for certain computational problems. CPUs and GPUs are designed with fundamentally different philosophies. While we've touched on this in the module on parallel architectures, let's look at it specifically in the context of CUDA.

A CPU is optimized for low-latency access to memory and executing a sequence of tasks as fast as possible. A significant portion of its silicon is dedicated to sophisticated control logic and large caches. A GPU, on the other hand, is a throughput machine. It dedicates most of its silicon to a massive number of simpler arithmetic logic units (ALUs), making it ideal for problems that can be broken down into thousands of identical, independent tasks.

To see a great visual explanation of this trade-off, please watch the first part of the following video.

This difference in design leads directly to the "heterogeneous" computing model, where the CPU and GPU work together, each handling the tasks for which it is best suited.

2. The Core CUDA Concepts: Host, Device, and Kernels

The CUDA programming model formalizes this heterogeneous approach by defining distinct roles for the CPU and GPU, and a specific way for the CPU to assign work to the GPU.

• Host: The CPU and its memory (system RAM). The host code is standard C++ that runs on the CPU. It manages the overall application flow, I/O, and serial parts of the code.
• Device: The GPU and its memory (on-board VRAM). The device code runs on the GPU's many cores.
• Kernel: A function written in CUDA C++ that is executed on the device. When the host calls (or "launches") a kernel, it is executed in parallel by many GPU threads.

The following image provides a high-level visual summary of this relationship. The host launches kernels, which execute on the device as a grid of threads.

Let's deepen this understanding by watching the next segment of the video, which introduces these core components.

Since you are proficient in C++, you can think of a kernel as a special C++ function, marked with __global__, that the compiler targets for the GPU architecture. The host code calls this function, but instead of executing it once on the CPU, the GPU executes it potentially millions of times in parallel.

An important aspect of this model is that kernel launches are asynchronous. When the host launches a kernel, it doesn't wait for the kernel to finish. It immediately continues executing the next lines of host code. This allows the CPU and GPU to work in parallel.

3. Organizing Parallelism: The Thread Hierarchy

The true power of CUDA lies in how it organizes the massive number of threads that execute a kernel. Instead of a flat, one-dimensional pool of threads, CUDA uses a three-level hierarchy:

1. Grid: The entire collection of threads launched for a single kernel. A kernel is launched as exactly one grid.
2. Block: A grid is composed of multiple thread blocks. All threads within a block can cooperate by sharing a special, fast on-chip memory (which we'll cover later) and can synchronize their execution.
3. Thread: A block is composed of multiple threads. This is the fundamental unit of execution that runs the kernel function.

Grids and blocks can be organized into one, two, or three dimensions, which is useful for mapping threads to data in problems with inherent dimensionality, such as image processing or matrix computations.

The following video provides an excellent and detailed explanation of this hierarchy.

4. Mapping Threads to Data

Now that we understand the hierarchy, the next logical question is: how does a specific thread know which piece of data it's supposed to work on?

CUDA provides built-in variables, accessible only within kernel code, that allow each thread to identify its unique position within the grid. These are the most important ones for indexing:

• threadIdx: A 3D vector (.x, .y, .z) specifying a thread's index within its block.
• blockIdx: A 3D vector (.x, .y, .z) specifying the block's index within the grid.
• blockDim: A 3D vector (.x, .y, .z) specifying the dimensions of the block (i.e., the number of threads in each dimension).
• gridDim: A 3D vector (.x, .y, .z) specifying the dimensions of the grid (i.e., the number of blocks in each dimension).

For a one-dimensional array of data, you can calculate a unique global index for each thread using the following idiomatic pattern:

Let's break this down:

• blockIdx.x * blockDim.x: This calculates the starting index of the current block. For example, if each block has 256 threads, block 0 starts at index 0, block 1 starts at index 256, block 2 starts at index 512, and so on.
• + threadIdx.x: This adds the thread's local index within its block to find its unique global position.

To see this in action, please read the following sections from two excellent articles.

5. The Complete Program Flow

Let's tie everything together. A typical CUDA program follows a standard pattern that reflects the host-device model:

1. Allocate memory on the device (GPU) for the data you want to process.
2. Copy data from the host (CPU memory) to the device (GPU memory).
3. Launch the kernel on the device to process the data in parallel. The host specifies the grid and block dimensions during the launch.
4. Wait for the kernel to complete if necessary.
5. Copy results from the device back to the host.
6. Free memory allocated on the device.

The following video segment walks through a simple but complete code example that demonstrates this entire flow.

Conclusion

In this lesson, we have unpacked the fundamental concepts of the CUDA programming model. You should now have a solid mental model of how a CUDA application is structured.

Here are the key takeaways:

• Heterogeneous Model: CUDA programs use both the CPU (host) and GPU (device), assigning serial work to the host and parallel work to the device.
• Kernels: These are C++ functions, marked with __global__, that are executed in parallel by many threads on the device.
• Thread Hierarchy: Threads are organized into a three-level hierarchy of grids, blocks, and threads. This structure is key to managing and scaling the parallelism.
• Thread Indexing: You can use built-in variables (threadIdx, blockIdx, blockDim) within a kernel to compute a unique global index for each thread, allowing it to operate on a specific element of your data.

In our next lesson, we will focus on the practical aspects of steps 1, 2, 5, and 6 of the program flow. You will learn how to explicitly manage GPU memory by allocating and deallocating it, and how to transfer data between the host and device using CUDA API functions like cudaMalloc and cudaMemcpy.