CUDA PROGRAMMING
CUDA in a NUTSHELL
Hierarchy of computations
→ Threads
→ Blocks
→ Grids
Corresponding Memory Spaces
→ Local
→ Shared
→ Global
Synchronisation Primitives
→ Implicit & Explicit Barriers
→ Thread Synchronisation
CUDA (Compute Unified Device Architecture) programming is a parallel
computing platform and application programming interface model created by
Nvidia. It allows developers to harness the power of Nvidia GPUs (Graphics
Processing Units) to accelerate computing tasks. CUDA programming is
widely used in fields such as scientific computing, machine learning, and
computer graphics, and has become an essential skill for developers looking
to optimize their applications for high-performance computing. With CUDA,
developers can take advantage of the massive parallel processing capabilities
of GPUs to significantly speed up their applications and achieve better
performance.
G80 Architecture
10 series Architecture
Data Movement
Data is moved from the host (CPU) to the GPU and vice versa through a
process called data transfer or data copying. This process involves copying
data from the main memory of the host to the memory of the GPU, and vice
versa. This is typically done using specialised functions provided by the GPU
programming framework, such as CUDA or OpenCL.
In CUDA, for example, data can be moved from the host to the GPU using the
cudaMemcpy function, and from the GPU to the host using the cudaMemcpy
function as well. These functions take as parameters the source and
destination memory addresses, the size of the data to be copied, and the
direction of the transfer (host to device or device to host).
Execution Model
The execution model of CUDA involves the use of parallel processing on the
GPU. In CUDA, a program is divided into a series of parallel threads, each of
which is executed on a separate core of the GPU. These threads are
organised into blocks, and the blocks are organised into a grid. The grid is
then executed on the GPU, with each block being assigned to a separate
streaming multiprocessor (SM) on the GPU.
The execution of the grid is managed by the CUDA runtime, which schedules
the blocks for execution on the available SMs. The threads within each block
are then executed in parallel on the SM, with each thread performing a
specific task. This allows for massive parallelism and high performance, as the
GPU can execute thousands of threads simultaneously.
Key parallel Abstractions in CUDA
1. Thread parallelism: CUDA allows for the parallel execution of multiple
threads on the GPU, enabling the processing of multiple data elements
simultaneously.
2. Block parallelism: CUDA organizes threads into blocks, which can be
executed in parallel on the GPU, allowing for efficient utilization of the GPU's
resources.
3. Grid parallelism: CUDA further organizes blocks into grids, enabling the
parallel execution of multiple blocks, which can be beneficial for handling
larger datasets and complex computations.
4. Warp parallelism: CUDA executes threads in groups called warps, which
allows for the simultaneous execution of multiple threads within a warp,
improving the overall performance of the GPU.
What are Kernels?
In CUDA, a kernel is a function that is executed in parallel by multiple threads
on the GPU. Kernels are defined with the `__global__` qualifier and are called
from the CPU to be executed on the GPU. Kernels are designed to perform a
specific task in parallel, such as image processing, matrix multiplication, or
any other computationally intensive operation.
Function Qualifiers
Function qualifiers in CUDA are used to specify the execution environment
and behavior of a function. There are several function qualifiers in CUDA,
including:
1. `__global__`: This qualifier is used to define a kernel function that will be
executed on the GPU.
2. `__device__`: This qualifier is used to define a function that will be executed
on the GPU and can be called from a kernel or another device function.
3. `__host__`: This qualifier is used to define a function that will be executed
on the CPU and can be called from the host code.
4. `__constant__`: This qualifier is used to define a variable that resides in
constant memory on the GPU and is accessible by all threads.
5. `__shared__`: This qualifier is used to define a variable that resides in
shared memory on the GPU and is accessible by all threads within a block.
These function qualifiers help to specify the behavior and execution
environment of functions in CUDA, allowing for efficient parallel processing on
the GPU.
Note:
→ As a programmer, All you have to worry about is giving the GPU a
big pile of thread
blocks and the GPU will take care of assigning them to run on the hardware
SMs
→ GPU is responsible for allocating blocks to SMs
Threads and Blocks
Threads are the smallest unit of execution in a parallel program. They are
individual sequences of instructions that can be executed independently. In
GPU programming, threads are organized into groups called blocks.
Blocks are a collection of threads that can be executed together on a GPU.
Blocks are used to divide the work into smaller units that can be executed in
parallel.
In GPU programming, threads and blocks are accessed using the threadIdx
and blockIdx parameters. These parameters allow the programmer to identify
the specific thread or block being executed and to use that information to
perform specific tasks or calculations.
Explicit and implicit synchronization
Explicit synchronization refers to the use of synchronization primitives, such
as barriers or locks, to coordinate the execution of threads or blocks. This
allows the programmer to control the order in which threads or blocks are
executed and to ensure that certain operations are completed before others
begin.
Implicit synchronization, on the other hand, refers to the automatic
synchronization that occurs in certain situations, such as when a kernel (a
function that runs on the GPU) is launched. In these cases, the GPU
automatically handles the synchronization of threads and blocks, without the
need for explicit synchronization primitives.
Thread synchronization
Thread synchronization is the process of coordinating the execution of
multiple threads to ensure that they do not interfere with each other and that
they can safely access shared resources. This is important in multi-threaded
applications where multiple threads are running concurrently and accessing
the same data or resources.
There are several mechanisms for achieving thread synchronization, including
locks, semaphores, and monitors. These mechanisms help to prevent race
conditions, where multiple threads try to access or modify the same resource
at the same time, leading to unpredictable behavior and potential data
corruption.
One common approach to thread synchronization is the use of locks, which
allow threads to acquire exclusive access to a resource. When a thread
acquires a lock, it prevents other threads from accessing the resource until the
lock is released. This ensures that only one thread can access the resource at
a time, preventing conflicts and ensuring data integrity.
Another approach is the use of semaphores, which are a more general
synchronization mechanism that can be used to control access to a resource
based on a specified limit. Semaphores can be used to coordinate the
execution of multiple threads and ensure that only a certain number of threads
can access a resource at a time.
Monitors are another synchronization mechanism that provides a higher-level
abstraction for coordinating access to shared resources. Monitors allow
threads to wait for a certain condition to be met before accessing a resource,
and they provide built-in support for mutual exclusion and condition variables.
Program Flow with CUDA
The program flow in CUDA follows a similar structure to traditional C
programs, with the addition of parallel execution on the GPU.
1. Host code: The program starts with the host code, which runs on the CPU.
This code is responsible for initializing the GPU, allocating memory on the
GPU, and launching the kernel functions that will be executed on the GPU.
2. Memory allocation: The host code allocates memory on the GPU using
CUDA-specific functions such as cudaMalloc() and cudaMemcpy(). This
allows data to be transferred from the CPU to the GPU for processing.
3. Kernel launch: The host code launches one or more kernel functions on the
GPU using the <<<...>>> syntax. These kernel functions are the parallel code
that will be executed on the GPU.
4. Kernel execution: The GPU executes the kernel functions in parallel on
multiple threads. Each thread processes a different portion of the data,
allowing for massive parallelism and increased performance.
5. Data transfer: After the kernel functions have finished executing, the host
code can transfer the results back from the GPU to the CPU using
cudaMemcpy().
6. Cleanup: Finally, the host code deallocates the memory on the GPU using
cudaFree() and performs any necessary cleanup before the program exits.
Overall, the program flow in CUDA involves initializing the GPU, transferring
data to the GPU, launching parallel kernel functions, executing the kernel
functions in parallel on the GPU, transferring the results back to the CPU, and
cleaning up the GPU resources.
Indexing With in Grid
In CUDA, indexing within a grid refers to the way in which threads are
organized and accessed within a grid of blocks. The grid is a collection of
thread blocks, and each block is a collection of threads.
The indexing within a grid is typically done using a two-dimensional grid of
blocks, where each block is identified by a unique two-dimensional index
(blockIdx.x, blockIdx.y). Within each block, the threads are organized in a one-
dimensional grid, and each thread is identified by a unique one-dimensional
index (threadIdx.x).
When accessing elements within the grid, the programmer can use the block
and thread indices to calculate the global index of each thread. This global
index can then be used to access the corresponding data elements in the
input and output arrays.
Overall, indexing within a grid in CUDA allows for efficient organization and
access of threads within a grid of blocks, enabling parallel processing of data
in a highly efficient manner.
Memory model
Memory model in GPUs refers to the organization and management of
memory resources within a graphics processing unit. GPUs typically have
several types of memory, including global memory, shared memory, and
constant memory.
Global memory is the largest and slowest type of memory in a GPU, and it is
used to store data that is accessible to all threads in a kernel. Shared
memory, on the other hand, is a smaller and faster type of memory that is
shared among threads within a thread block. It is used for communication and
data sharing between threads.
Constant memory is a type of read-only memory that is used to store data that
is constant throughout the execution of a kernel. It is optimized for high
bandwidth and low latency access.
The memory model in GPUs also includes the management of memory
access patterns, such as coalesced memory access, which refers to the
efficient access of memory by threads in a warp.
Overall, the memory model in GPUs is designed to optimize memory access
and utilization for parallel processing, allowing for efficient and high-
performance execution of graphics and compute workloads.
Constant Memory
Constant memory in GPUs refers to a type of memory that is used for storing
data that remains constant throughout the execution of a kernel (a function
that runs on the GPU). This type of memory is typically used for storing read-
only data that is accessed frequently by the kernel, such as lookup tables,
constants, or other data that does not change during the execution of the
kernel.
Constant memory is optimized for high-speed access and is typically smaller
in size compared to other types of memory in the GPU, such as global
memory. It is also cached to improve access speed, making it ideal for storing
data that is repeatedly accessed by the kernel.
One of the key advantages of using constant memory in GPUs is that it can
significantly improve the performance of the kernel by reducing memory
access latency and improving overall throughput. By storing frequently
accessed read-only data in constant memory, the GPU can access this data
more quickly, leading to faster execution of the kernel.
Overall, constant memory in GPUs plays a crucial role in optimizing the
performance of GPU kernels by providing a high-speed, read-only memory
space for storing frequently accessed data
Example code
```cpp
#include <iostream>
#include <cuda_runtime.h>
__global__ void kernel(int* array, int size) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < size) {
array[tid] *= 2;
}
}
int main() {
const int size = 10;
int hostArray[size] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10};
int* deviceArray;
cudaMalloc((void**)&deviceArray, size * sizeof(int));
cudaMemcpy(deviceArray, hostArray, size * sizeof(int),
cudaMemcpyHostToDevice);
int blockSize = 256;
int gridSize = (size + blockSize - 1) / blockSize;
kernel<<<gridSize, blockSize>>>(deviceArray, size);
cudaMemcpy(hostArray, deviceArray, size * sizeof(int),
cudaMemcpyDeviceToHost);
for (int i = 0; i < size; i++) {
std::cout << hostArray[i] << " ";
}
std::cout << std::endl;
cudaFree(deviceArray);
return 0;
}
```
PyCUDA
PyCUDA is a Python library that provides a way to use Nvidia's CUDA
framework for parallel computing directly from Python. It allows developers to
write high-performance GPU-accelerated code using Python, while still having
access to the full power and flexibility of CUDA. With PyCUDA, users can
easily create and manipulate GPU arrays, write custom CUDA kernels, and
execute them on the GPU, all from within Python. This makes it a valuable
tool for anyone looking to harness the power of GPU computing for their
applications.
This tutorial will guide you through the process of setting up PyCUDA and
running a simple CUDA kernel.
Step 1: Install CUDA Toolkit
Before you can use PyCUDA, you need to have the CUDA Toolkit installed on
your system. You can download the CUDA Toolkit from Nvidia's website and
follow the installation instructions for your operating system.
Step 2: Install PyCUDA
Once you have the CUDA Toolkit installed, you can install PyCUDA using pip:
```
pip install pycuda
```
Step 3: Verify Installation
To verify that PyCUDA is installed correctly, you can run the following Python
code:
```python
import pycuda
print([Link])
```
If PyCUDA is installed correctly, you should see the version number printed to
the console.
Step 4: Write a CUDA Kernel
Now that PyCUDA is installed, you can write a simple CUDA kernel in Python.
Here's an example of a simple kernel that adds two arrays together:
```python
import [Link]
import [Link] as cuda
from [Link] import SourceModule
mod = SourceModule("""
__global__ void add(int *a, int *b, int *c) {
int idx = threadIdx.x;
c[idx] = a[idx] + b[idx];
}
""")
add = mod.get_function("add")
```
Step 5: Allocate Memory and Run Kernel
Next, you need to allocate memory on the GPU and run the kernel. Here's an
example of how to do this:
```python
import numpy as np
a = [Link]([1, 2, 3, 4]).astype(np.int32)
b = [Link]([5, 6, 7, 8]).astype(np.int32)
c = np.zeros_like(a)
a_gpu = cuda.mem_alloc([Link])
b_gpu = cuda.mem_alloc([Link])
c_gpu = cuda.mem_alloc([Link])
cuda.memcpy_htod(a_gpu, a)
cuda.memcpy_htod(b_gpu, b)
add(a_gpu, b_gpu, c_gpu, block=(4, 1, 1))
cuda.memcpy_dtoh(c, c_gpu)
print(c)
```
This code allocates memory on the GPU for the input and output arrays,
copies the input arrays to the GPU, runs the CUDA kernel, and then copies
the result back to the CPU. Finally, it prints the result.
That's it! You've now successfully written and run a simple CUDA kernel using
PyCUDA. You can use this as a starting point to explore more advanced
CUDA programming with PyCUDA.