GPU Programming Concepts¶
This episode introduces the fundamental concepts behind GPU programming and explains why GPUs are highly effective at accelerating data-parallel applications. It begins by discussing different types of parallelism, including task parallelism, instruction-level parallelism, and data parallelism, and emphasizes that GPUs are primarily designed to exploit massive data parallelism. This episode also introduces the concept of exposing parallelism, which involves restructuring algorithms so that many independent operations can be executed simultaneously, enabling efficient utilization of thousands of GPU cores.
The episode then presents the GPU execution model and the core terminology used in GPU programming. It explains how computations are organized into CUDA threads, which are grouped into warps, and how multiple warps form a thread block. It also introduces the hierarchical organization of threads, blocks, and grids, providing the foundation for understanding how GPU kernels are executed and how work is distributed across the hardware. Together, these concepts establish the essential vocabulary and programming model needed to develop efficient GPU applications.
Questions
What types of parallelism can be exploited in parallel computing?
What is the difference between data parallelism and task parallelism, and how are they used in parallel computing?
How are computations parallelized and executed on GPUs? What are the key considerations for writing efficient GPU programs?
What are the general principles for writing efficient GPU code?
Objectives
Explain the different types of parallelism, including data parallelism and task parallelism, and identify when each is appropriate.
Describe the GPU execution model, including the roles of threads, warps, thread blocks, and grids.
Explain how computational work is exposed, organized, and mapped onto GPU hardware for parallel execution.
Apply the fundamental terminology of GPU programming, such as kernels, threads, warps, blocks, and occupancy, when discussing GPU applications.
Identify the key characteristics of algorithms and applications that lead to efficient execution on GPU architectures.
Instructor note
30 min teaching
10 min exercises/discussion
Different Types of Parallelism¶
Before exploring GPU programming, it is important to understand the fundamental parallel computing models used in modern computer architectures. This section introduces the differences between distributed-memory and shared-memory architectures and explains how parallel programs are organized using processes and threads. These concepts provide the foundation for understanding how computational work is divided, coordinated, and executed efficiently across CPUs and GPUs.
Distributed-memory vs. shared-memory architecture¶
Most computational problems are not trivially parallelizable, meaning that subtasks need to periodically exchange intermediate results or share data during execution. How this communication is performed depends on the underlying hardware architecture.
In a distributed-memory system, each processing unit operates independently with its own private memory and cannot directly access the memory of other processing units. Data exchange or communication between processing units occurs through message passing over a network, and each processing unit runs its own instance of the operating system. In contrast, a shared-memory system allows all processing units to access the same memory space, enabling them to read and modify shared variables directly.
Figure: Distributed-memory vs shared-memory parallel computing.¶
Processes and threads¶
The underlying hardware architecture, distributed or shared memory, determines the appropriate parallel programming model. Correspondingly, there are two primary forms of parallelism: process-based parallelism and thread-based parallelism, as shown in the figure below.
In distributed-memory systems, a process-based parallel programming model is employed. Each process is an independent execution unit with its own private memory address space. Processes are created when the parallel program starts and remain active until the program terminates. Because processes cannot directly access each other’s memory, they communicate explicitly by exchanging messages using libraries such as MPI (Message Passing Interface).
In shared-memory architectures, a thread-based parallel programming model can be used. Threads are lightweight execution units that can be created and terminated with relatively low overhead. Each thread maintains its own state information, while all threads share the same memory address space. Communication between threads is performed through shared memory, allowing them to access and modify common data when needed.
For applications running on shared-memory systems, common parallel programming libraries include OpenMP, which provides a simple directive-based approach for parallelizing existing code, and Pthreads, which offers low-level control over thread creation, management, and synchronization. Other modern approaches, such as C++ threads and Intel oneTBB, provide higher-level abstractions for developing efficient shared-memory parallel applications.
Note
In the context of GPU programming, it is worth noting that CUDA and HIP also use a thread-based execution model. However, unlike traditional CPU shared-memory programming models, they are specifically designed to target GPU architectures and exploit massive parallelism across thousands of GPU cores.
Exposing Parallelism¶
There are two main types of parallelism that can be exploited in parallel computing: data parallelism and task parallelism.
Data parallelism focuses on distributing data across multiple computational units.
Task parallelism involves executing different tasks concurrently.
Data parallelism¶
Data parallelism occurs when a dataset can be divided among multiple computational units that execute in parallel, as shown in the figure below. Each unit performs the same or similar operations on different portions of the data. A common example is applying a blur filter to an image, where the same function is independently applied to each pixel. This type of parallelism maps naturally to GPUs, where the same instruction sequence can be executed simultaneously by many threads.
Data parallelism can often be exploited effectively on GPUs. A common starting point is to identify loops that operate over large numbers of data elements and transform them into GPU kernels, allowing different iterations to execute concurrently on multiple GPU threads. If the dataset contains a sufficiently large number of elements (tens or hundreds of thousands), GPUs can achieve significant performance improvements.
While this straightforward approach may not deliver peak GPU performance, it is often the first and most practical step when porting applications. Achieving maximum performance requires a deeper understanding of GPU architecture, including factors such as memory access patterns, thread organization, and hardware utilization.
Note
Data parallelism does not automatically guarantee GPU acceleration. Applications must expose sufficient parallel work, minimize CPU–GPU data transfers, and use efficient memory access patterns to achieve good performance. Simply moving a loop to a GPU kernel may provide limited speedup if the workload is too small, memory-bound, or contains significant overhead.
Figure: Data parallelism and task parallelism. The data parallelism is when the same operation applies to multiple data (e.g., multiple elements of an array are transformed). The task parallelism implies that there are more than one independent task that, in principle, can be executed in parallel.¶
Task parallelism¶
Another type of parallelism is task parallelism. It occurs when an application consists of multiple independent tasks that perform different operations on the same or different data and can be executed concurrently. An everyday example of task parallelism is cooking: slicing vegetables and grilling can be performed simultaneously because they are separate tasks.
It is important to note that different tasks may require different computational resources. For example, one task may be better suited for a CPU, while another may benefit from execution on a GPU or other specialized hardware. This flexibility allows applications to take advantage of heterogeneous computing resources to improve performance.
Data/task parallelism vs process/thread parallelism¶
Data/task parallelism and process/thread parallelism are related, but they describe different levels of parallel computing concepts.
Data parallelism vs. task parallelism describes how computational work is divided.
Process-based vs. thread-based parallelism describes how parallel work is executed and how computational units communicate.
They are complementary concepts and can be combined.
Data parallelism + threads: A common approach for shared-memory systems. Multiple threads perform the same operation on different portions of data.
Example: Using OpenMP threads to divide a large array among CPU cores, where each thread processes a subset of elements.
Thread 1 → process array elements 0–999 Thread 2 → process array elements 1000–1999 Thread 3 → process array elements 2000–2999
This is also the dominant model for GPUs, where thousands of GPU threads execute the same operation on different data elements.
Data parallelism + processes: Common in distributed-memory systems. Different processes operate on different portions of a large dataset, usually with communication through MPI.
Example: A weather simulation where each MPI process computes a different region of the simulation domain.
Process 1 → compute region A Process 2 → compute region B Process 3 → compute region C
Task parallelism + threads: Multiple threads execute different tasks concurrently within the same memory space.
Example: One thread reads input data, another performs computation, and another writes results.
Thread 1 → read and preprocess input data Thread 2 → perform numerical computation Thread 3 → write and store results
Task parallelism + processes: Different processes execute independent tasks, often on different nodes.
Example: A scientific workflow where one process performs data preprocessing, another runs simulations, and another analyzes results.
Process 1 → perform data preprocessing Process 2 → run scientific simulation Process 3 → analyze simulation results
Below are representative software packages and programming models for the four types of parallelism:
Parallelism type |
Representative software |
Example applications |
|---|---|---|
Distributed-memory + data parallelism |
MPI (Message Passing Interface), PETSc, HYPRE, Lattice Boltzmann Method (LBM) frameworks |
Large-scale scientific simulations where data domains are divided across multiple compute nodes |
Shared-memory + data parallelism |
OpenMP, oneAPI Threading Building Blocks (oneTBB), C++ |
Parallel loops over arrays, matrix operations, image processing, and numerical kernels on multicore CPUs |
Thread-based task parallelism |
OpenMP tasks, oneTBB task scheduler, C++ |
Applications where different threads perform independent activities, such as pipeline processing, I/O, and computation |
Process-based task parallelism |
MPI, Dask, Apache Spark, workflow engines such as Apache Airflow |
Workflows where independent programs or tasks execute concurrently across nodes or machines |
In short
Computing problems can be parallelized using distributed-memory or shared-memory architectures.
In distributed-memory systems, each computing unit operates independently with its own private memory, and data exchange between nodes is performed through explicit communication.
In shared-memory systems, multiple computing units access the same memory space and communicate by reading and modifying shared variables.
Parallel programming can be implemented using process-based or thread-based models, depending on the underlying memory architecture.
Process-based parallelism uses independent processes with separate memory spaces, requiring explicit communication mechanisms such as message passing.
Thread-based parallelism uses lightweight threads that share the same memory space and communicate through shared data.
Parallel programming can also be implemented using data parallelism and task parallelism, depending on the structure of application and how computational workload can be divided.
Data parallelism distributes data across multiple computational units, where each unit performs the same or similar operations on different data elements.
Task parallelism divides an application into multiple independent tasks that can execute concurrently, allowing different tasks to be assigned to the most suitable computational resources.
GPU Execution Model¶
To achieve maximum performance, it is important to understand how GPUs execute programs. As mentioned previously, a CPU is a flexible processor designed for general-purpose computing. It is fast and versatile, capable of running operating systems and a wide variety of applications. CPUs include many features, such as sophisticated control logic, large caches, and cache coherence mechanisms, that support complex workloads but are not directly related to raw computational throughput. They optimize execution primarily by minimizing latency through techniques such as extensive caching and branch prediction.
Figure: Analogy of CPU and GPU execution. A narrow road represents a CPU, optimized for low latency but limited throughput, while a wide road represents a GPU, optimized for high throughput by processing many tasks in parallel.¶
In contrast, GPUs dedicate a relatively small fraction of their transistors to control logic and caching, while a much larger fraction is devoted to performing mathematical operations. Because GPU cores are designed primarily for tasks such as 3D graphics rendering, they can be made significantly simpler than CPU cores, allowing a much larger number of them to be integrated into a single device. Modern GPUs contain thousands of CUDA cores.
GPU performance is achieved through a high degree of parallelism. Large numbers of threads are launched and executed concurrently, and achieving good performance typically requires using several times more threads than the number of available CUDA cores. GPU threads are much lighter-weight than traditional CPU threads and incur very little overhead during context switching. This allows the GPU to efficiently hide memory latency: while some threads are waiting for memory operations (such as reads or writes) to complete, other threads can continue executing instructions.
Threads, Warps/Wavefronts, Blocks, and Grids¶
To understand the GPU execution model, here we consider the well-known axpy operation.
This operation consists of multiplying each element of an input array by a scalar value and adding the result to the corresponding element of another array.
On a single CPU core, this computation is performed sequentially using a for/do loop that iterates over each element of the array.
For every index id, the CPU executes the operation y[id] = y[id] + a * x[id] one element at a time until the entire array has been processed.
void axpy_(int n, double a, double *x, double *y)
{
for(int id=0;id<n; id++) {
y[id] += a * x[id];
}
}
To perform the same operation on a GPU, the program launches a function called a kernel.
Note
A GPU kernel is conceptually similar to a subroutine in Fortran or a function in C: it is a block of code that performs a specific computation and can be invoked by a program.
Unlike a CPU implementation, where the operation is executed sequentially, a GPU kernel is executed concurrently by a large number of threads. Tens of thousands of lightweight threads can be launched and scheduled across the available GPU cores, allowing many elements of the array to be processed simultaneously. Each thread typically performs the computation for a single element (or a small subset of elements), enabling the GPU to achieve high throughput through massive parallelism.
GPU_K void ker_axpy_(int n, double a, double *x, double *y, int id)
{
y[id] += a * x[id]; // id<n
}
Programmers control how many instances of ker_axpy_ are launched and must ensure that all elements of the data are processed while avoiding any out-of-bounds memory accesses.
Threads¶
Threads (for both NVIDIA and AMD GPUs) are much lighter-weight than traditional CPU threads and incur very little overhead when switching between execution contexts. By “over-subscribing” the GPU — launching more threads than can execute simultaneously — the GPU can effectively hide memory latency. While some threads are waiting for memory operations, such as reads or writes, to complete, other threads can continue executing instructions, keeping the computational units occupied. This allows the hardware to remain busy and improves overall throughput.
The figure below illustrates the relationship between CUDA threads and CUDA cores. It is important to note that a CUDA core is not the same as a CUDA thread. Although they are closely related, they represent different concepts.
A CUDA core is a hardware execution unit inside a GPU.
It is a physical arithmetic unit capable of performing operations such as integer and floating-point calculations.
A CUDA thread is a software execution context created when a kernel is launched.
It represents a logical instance of the kernel code operating on a particular piece of data.
A GPU may contain thousands of CUDA cores but can have millions of threads launched. Threads are not permanently assigned one-to-one to CUDA cores. Instead, the GPU scheduler groups threads into warps (typically 32 threads on NVIDIA GPUs) or wavefronts (typically 64 threads on AMD GPUs), and warps/wavefronts are scheduled onto the available execution units. When one group of threads is waiting for memory, another group can be scheduled to use the CUDA cores.
Every CUDA thread is assigned a unique intrinsic index that can be used to determine which elements of an array it should process. Each thread has its own execution context and a private set of local variables. Although all threads can access the GPU’s global memory, there is no general mechanism for synchronizing all threads during the execution of a kernel. Consequently, if some threads need to read data from global memory that has been modified by other threads, the computation must be divided into multiple kernel launches. Kernel execution acts as a global synchronization point, ensuring that all writes to global memory are completed before the next kernel begins.
Warps/Wavefronts¶
Besides being much more lightweight than CPU threads, GPU threads differ in several other important ways. GPU threads are organized into groups called warps (for NVIDIA GPU) or wavefronts (for AMD GPU), a grouping that is performed automatically by the hardware.
On some GPU architectures, all threads within a warp/wavefront execute the same instruction simultaneously, a model known as lock-step execution.
This improves performance but makes conditional branches (if statements) potentially expensive.
If threads within the same warp/wavefront take different execution paths, a phenomenon known as warp divergence occurs.
The hardware executes each branch sequentially, with only the relevant threads active, reducing parallel efficiency.
On newer architectures without lock-step execution, such as NVIDIA Volta / Turing (e.g., GeForce 16xx-series) and later, independent thread scheduling reduces the cost of warp divergence, although minimizing divergence remains an important optimization.
Blocks¶
The GPU execution hierarchy includes another level above threads: blocks. Threads are grouped into blocks, and each block is assigned to a single Streaming Multiprocessor (SMP) unit. An SMP contains one or more SIMT (Single Instruction, Multiple Threads) execution units, schedulers, and a small amount of very fast on-chip memory. Part of this on-chip memory is exposed to programmers as shared memory.
Shared memory can be used to “cache” data that is accessed by multiple threads, reducing the number of accesses to the much slower global memory. It is also useful for reorganizing memory accesses to improve efficiency. For example, in a matrix transpose, only one of the two memory operations (read or write) is naturally coalesced. By first loading a tile of the matrix into shared memory using coalesced reads, the tile can be transposed locally and then written back using coalesced writes, significantly improving performance.
Shared memory is also widely used for block-level reductions and other collective operations. Since all threads within a block can synchronize, they can safely cooperate by exchanging data through shared memory. When shared memory is used for communication between threads, explicit synchronization is required to ensure that all writes are completed before the data is read, guaranteeing correct program execution.
Finally, a thread block cannot be split across multiple SMPs; all threads within a block are executed on the same SM. For good performance, each block should contain multiple warps, enabling the SMP to switch between warps and hide memory latency while some are waiting for memory operations to complete. If sufficient hardware resources are available, an SMP can execute multiple blocks concurrently.
Grids¶
Threads are organized into blocks, and the complete collection of blocks launched for a kernel is called a grid. Blocks are independent of one another and may be scheduled on any available SMP units, allowing the GPU to efficiently distribute the workload across the hardware.
It should be noted that blocks within the same grid cannot communicate directly through shared memory, because shared memory is allocated separately for each block and is only visible to the threads belonging to that block. If data exchange or synchronization between blocks is required, it must be performed through global memory, typically by separating the computation into multiple kernel launches. A kernel completion provides the global synchronization point needed to ensure that all memory operations from the previous kernel have completed.
Mapping threads, blocks, and grids to array elements¶
Below is an example illustrating how threads and blocks in a grid are mapped to specific elements of an array.
All threads in the preceding figure belong to a grid containing 4096 threads.
The threads are organized into blocks of 256 threads each, with representative threads highlighted in orange, light blue, dark blue, and green.
The orange thread has a local index of 3 within block 2 (highlighted in dark blue), resulting in a global thread index of 515.
For a vector addition example, this global index is used to access the corresponding array element c[index] = a[index] + b[index].
Terminology¶
Currently, there are three major GPU manufacturers: NVIDIA, Intel, and AMD. Although the fundamental concepts behind their GPU architectures are similar, each company uses different terminology to describe the various hardware components. In addition, several software environments for GPU programming exist, including vendor-specific frameworks and third-party solutions, each with its own terminology. The table below provides a brief compilation of commonly used terms across different GPU platforms and programming environments.
CUDA |
HIP |
OpenCL |
SYCL |
|---|---|---|---|
grid of threads |
NDRange |
||
block |
work-group |
||
warp |
wavefront |
sub-group |
|
thread |
work-item |
||
registers |
private memory |
||
shared memory |
local data share |
local memory |
|
threadIdx.{x,y,z} |
get_local_id({0,1,2}) |
nd_item::get_local({2,1,0}) |
|
blockIdx.{x,y,z} |
get_group_id({0,1,2}) |
nd_item::get_group({2,1,0}) |
|
blockDim.{x,y,z} |
get_local_size({0,1,2}) |
nd_item::get_local_range({2,1,0}) |
|
Warning
In SYCL, the thread indexing is inverted. In a 3D grid, physically adjacent threads have consecutive X (0) index in CUDA, HIP, and OpenCL, but consecutive Z (2) index in SYCL. In a 2D grid, CUDA, HIP, and OpenCL still has contiguous indexing along X (0) dimension, while in SYCL it is Y (1). Same applies to block dimensions and indexing.
Summary of GPU Execution Hierarchy¶
Taking full advantage of GPUs requires algorithms that can divide a problem into many small independent tasks that can be executed concurrently. Computations are offloaded to the GPU by launching thousands of threads that execute the same function (kernel) with each thread processing a different part of the problem. Threads are organized into groups called blocks, with each block assigned to a Streaming Multiprocessor (SMP). The threads within a block are further divided into warps/wavefronts, which are executed by SIMT units. Threads within a warp execute instructions together, and memory accesses are performed collectively at the warp level. Threads can synchronize and exchange data through shared memory at the block level, while some architectures also provide limited data exchange mechanisms at the warp level.
To hide latency and maximize hardware utilization, it is generally beneficial to “over-subscribe” GPU by launching many more threads and blocks than can execute simultaneously. In particular, there should be significantly more blocks than SMPs available on the device. Similarly, maintaining multiple active warps on each SMP helps keep the computational units occupied. While some warps are waiting for memory operations to complete, other warps can continue executing instructions, effectively hiding latency and improving overall throughput.
In addition to basic execution model, modern GPUs provide architecture-specific features that can further improve performance. Warp-level operations are hardware-supported primitives that enable efficient communication and synchronization between threads within a warp without requiring explicit block-level synchronization. These operations, together with the hierarchical organization of threads into blocks and grids, allow the implementation of efficient parallel algorithms. More advanced features, such as CUDA cooperative groups, provide finer control over thread organization by allowing threads within a block to form smaller groups that can synchronize and exchange data independently.
Exercises¶
What are threads in the context of shared-memory architectures?
A. Independent execution units with their own private memory address spaces.
B. Lightweight execution units that share a common memory address space.
C. Communication mechanisms between separate memory units.
D. Programming models for distributed-memory systems
Solution
Correct answer: B). Threads are lightweight execution units that run concurrently while sharing access to a common memory space. This allows threads within the same program to efficiently exchange data and collaborate on parallel tasks.
What is data parallelism?
A. Distributing data across computational units that execute in parallel, applying the same or similar operations to different data elements.
B. Distributing tasks across computational units that execute in parallel, applying different operations to the same data elements.
C. Distributing data across computational units that execute sequentially, applying the same operation to all data elements.
D. Distributing tasks across computational units that execute sequentially, applying different operations to different data elements.
Solution
Correct answer: A). Data parallelism divides a large dataset into smaller portions that can be processed simultaneously by multiple computational units. Each unit performs the same or similar operation on different data elements, increasing performance through parallel execution.
What type of parallelism is most natural for GPUs?
A. Task parallelism
B. Data parallelism
C. Both data and task parallelism
D. Neither data nor task parallelism
Solution
Correct answer: B). GPUs are designed to execute the same operation on many data elements simultaneously, making them highly efficient for data-parallel workloads. Their architecture uses thousands of lightweight threads that perform identical computations on different portions of the data.
What is a kernel in the context of GPU execution?
A. A specific section of the CPU used for memory operations.
B. A specific section of the GPU used for memory operations.
C. A type of thread that operates on the GPU.
D. A function that is executed simultaneously by thousands of threads on GPU cores.
Solution
Correct answer: D). A GPU kernel is a function written to run on the GPU and is executed concurrently by many threads. Each thread typically processes a different portion of the data, enabling massive parallelism and high computational throughput.
What is the function of shared memory in the context of GPU execution?
A. It is used to store global memory data.
B. It is used to store all threads within a block.
C. It can be used to cache data accessed by multiple threads, reducing repeated reads from global memory.
D. It is used to store all CUDA cores.
Solution
Correct answer: C). Shared memory is a fast on-chip memory space shared by threads within the same block. It allows threads to reuse data efficiently and reduces the need to repeatedly access slower global memory.
What is the significance of over-subscribing the GPU?
A. It reduces the overall performance of the GPU.
B. It ensures that there are more blocks than SMs available on the device, helping to hide latency and maintain high GPU occupancy.
C. It causes a memory overflow in the GPU.
D. It ensures that there are more SMs than blocks available on the device.
Solution
Correct answer: B). Over-subscribing the GPU by launching more blocks than available SMs allows the hardware to keep additional work ready for execution. While some threads are waiting for memory operations to complete, other threads can execute, helping to hide latency and improve GPU utilization.
Keypoints
Parallel computing can be classified into distributed-memory and shared-memory architectures.
Two main types of parallelism that can be exploited are data parallelism and task parallelism.
GPUs are a type of shared-memory architecture that are particularly well suited for data-parallel workloads.
GPUs achieve massive parallelism by organizing threads into warps and blocks.
GPU optimization involves techniques such as efficient use of shared memory, maintaining high thread and warp occupancy.
In addition, architecture-specific features such as warp-level operations and cooperative groups can be used to further improve performance.