Learning CUDA with PMPP, Part 3: CUDA C++ Program Structure

CUDA C++ is built on NVIDIA’s CUDA platform. The structure of a CUDA program includes both host (CPU) code and device (GPU) code. Device code is marked with a special keyword and includes functions, or kernels, that are executed in a data-parallel manner.

A CUDA C++ program begins with serial host-code execution. When the host launches a kernel, it creates a grid of device threads. Each grid consists of multiple thread blocks, and each block contains multiple threads. Once all threads in the grid finish executing, the kernel terminates and host execution continues, or another kernel is launched. The power of heterogeneous programming comes from using both the CPU and GPU effectively.

A thread is a sequential execution context consisting of program code, its current execution point, and the variables and data structures required for execution.

In CUDA, a program starts on the host as serial CPU code. Parallel execution begins when a kernel function—the parallel device code—is called. This launches a grid of device threads that process different portions of the data in parallel. Threads in a grid are organized into blocks. After execution terminates, the program continues on the host.

A simple vector addition

We can explore the CUDA programming model by writing a simple vector addition kernel.

Suppose we want to add two vectors: A = [1, 2, 3, 4, 5, 6, 7, 8, 9, 10] and B = [10, 20, 30, 40, 50, 60, 70, 80, 90, 100]. Solving this serially, our solution would be to loop through both arrays and add the corresponding elements into their corresponding element in vector C.

// vec_add.cpp
#include <iostream>

void vec_add(const float* A, const float* B, float* C, int n)
{
    for (int i = 0; i < n; i++)
    {
        C[i] = A[i] + B[i];
    }
}

int main()
{
    constexpr int n = 10;

    float A[n] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10};
    float B[n] = {10, 20, 30, 40, 50, 60, 70, 80, 90, 100};
    float C[n] = {};

    vec_add(A, B, C, n);
    std::cout << "C = [";

    for (int i = 0; i < n; i++) {
        if (i > 0) {
            std::cout << ", ";
        }
        std::cout << C[i];
    }
    std::cout << "]\n";

    return 0;
}

The vector addition code above is executed serially. Pointers to the ith elements of A and B are used to iteratively add each pair of elements and move to the next until the entire array is filled.

Moving the work to the device

A straightforward way to parallelize the vector addition function is to move its execution onto the device (GPU). This involves a few steps:

1. Allocate space on the device to hold copies of A, B, and C.
2. Call the vector addition kernel function to launch a grid of threads on the device to process the data.
3. Copy the data back to host memory and free the allocated device memory.

CUDA runtime functions

The API functions for allocating and freeing device global memory (DRAM) are:

cudaMalloc()

Requires the address of the pointer and the size of the object in bytes.

cudaFree()

Requires a pointer to the memory to be freed.

cudaMemcpy()

Performs a memory transfer and requires four parameters: a pointer to the destination, a pointer to the source, the size, and the transfer type or device direction.

• The final portion is calling a kernel to process the data.