CUDA Kernel Execution: Memory Tracking

1. The CUDA Model: A Quick Refresher

Before we trace, let's briefly recall the CUDA programming model. A CUDA program involves two distinct entities:

• The Host (the CPU) which runs the main C++ part of your program.
• The Device (the GPU) which executes massively parallel functions called kernels.

The host is responsible for managing the overall workflow: allocating memory on the GPU, transferring data to and from the GPU, and launching kernels on the GPU.

The kernel launch creates a grid of thread blocks. Each block contains a group of threads. This hierarchical structure maps onto a corresponding memory hierarchy.

• Global Memory: A large memory space accessible by all threads in the grid, as well as the host. Data must be transferred from the host to global memory before the kernel can use it. It's powerful but has higher latency.
• Shared Memory: A smaller, very fast memory space that is private to each thread block. It's a key tool for performance optimization, allowing threads within a block to communicate and share data efficiently. It acts as a user-managed cache.

Our tracing will focus on tracking data as it moves between and is modified within these two memory spaces.

2. Tracing with Global Memory: Who Does What?

The first step in tracing any CUDA kernel is to determine the unique identity of each thread. Since every thread in the grid executes the same kernel code, we need a way to assign different work to different threads. This is done using built-in variables like threadIdx, blockIdx, and blockDim.

For a 1D grid of 1D blocks, the standard formula to calculate a unique global index for each thread is:
$\text{global\_idx}=\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}$

• threadIdx.x: The thread's index within its block (from 0 to blockDim.x - 1).
• blockIdx.x: The block's index within the grid.
• blockDim.x: The number of threads in each block.

This video provides an excellent visual explanation of how these variables are used to compute this unique index.

Tracing Example: Vector Addition

Let's apply this to trace a classic vector addition kernel. Each thread will be responsible for adding one pair of elements.

Kernel Code:

__global__ void vectorAdd(const float* A, const float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

The if (idx < N) check is a standard safety measure to prevent threads from accessing memory out of bounds if the number of elements N is not an exact multiple of the grid size.

Scenario:

• A (global memory) = [10, 20, 30, 40]
• B (global memory) = [1, 2, 3, 4]
• C (global memory) = [0, 0, 0, 0] (initially)
• Kernel Launch: vectorAdd<<<gridDim=2, blockDim=2>>>(A, B, C, 4);
  • This launch creates a grid with 2 blocks.
  • Each block has 2 threads.
  • Total threads = 2 * 2 = 4.

Let's trace the execution, focusing on the value of idx and the write to C for each thread.

Final State of C in Global Memory: [11, 22, 33, 37]

As you can see, even though all four threads run the same code, the indexing formula ensures each one operates on a different element of the arrays.

3. Tracing with Shared Memory: Cooperation and Synchronization

Shared memory introduces a new layer to our tracing. It's a powerful feature for performance, but it requires careful coordination between threads. The typical pattern is for threads in a block to:

1. Cooperatively load a tile of data from global memory into the block's shared memory.
2. Synchronize to ensure all threads in the block have finished loading before any thread starts computing.
3. Perform computations on the data in fast shared memory.
4. Write the final results back to global memory.

The key to this coordination is the __syncthreads() intrinsic. When a thread reaches this function, it pauses until every other thread in the same block also reaches it. This creates a barrier, preventing race conditions where one thread might try to read a value from shared memory before another thread has written it.

This next video explains these concepts in the context of a 1D stencil problem, demonstrating the motivation for shared memory and the critical role of synchronization.

Tracing Example: Array Reversal with Shared Memory

Let's trace a kernel that reverses an array. It uses shared memory to avoid inefficient memory access patterns when writing back to global memory. This example is adapted from the NVIDIA Developer Blog post "Using Shared Memory in CUDA C/C++".

Kernel Code:

__global__ void staticReverse(int* d, int n) {
    // Statically allocated shared memory array, private to each block
    __shared__ int s[64]; 

    // Local indices
    int t = threadIdx.x;      // This thread's original index
    int tr = n - 1 - t;       // This thread's reversed index

    // 1. Load from global memory 'd' into shared memory 's'
    s[t] = d[t];

    // 2. Synchronize all threads in the block
    __syncthreads();

    // 3. Write from shared memory 's' back to global memory 'd'
    d[t] = s[tr];
}

Scenario:

• d (global memory) = [10, 20, 30, 40]
• Kernel Launch: staticReverse<<<gridDim=1, blockDim=4>>>(d, 4);
  • 1 block, 4 threads.

Let's trace the state of global (d) and shared (s) memory step-by-step.

1. Initial State:

• d (global): [10, 20, 30, 40]
• s (shared): [?, ?, ?, ...] (uninitialized)

2. Execution before __syncthreads():
Each thread t executes s[t] = d[t]. We assume all these writes complete before the barrier.

• Thread 0 (t=0): s[0] = d[0], so s[0] becomes 10.
• Thread 1 (t=1): s[1] = d[1], so s[1] becomes 20.
• Thread 2 (t=2): s[2] = d[2], so s[2] becomes 30.
• Thread 3 (t=3): s[3] = d[3], so s[3] becomes 40.
• State of s after this phase: [10, 20, 30, 40, ...]

3. The __syncthreads() Barrier:
All 4 threads pause here. The program only proceeds once all 4 threads have reached this point. This guarantees that s is fully populated with the initial data.

4. Execution after __syncthreads():
Each thread t executes d[t] = s[tr].

• Thread 0 (t=0, tr=3): d[0] = s[3], which is 40.
• Thread 1 (t=1, tr=2): d[1] = s[2], which is 30.
• Thread 2 (t=2, tr=1): d[2] = s[1], which is 20.
• Thread 3 (t=3, tr=0): d[3] = s[0], which is 10.

5. Final State of Global Memory d:

• d: [40, 30, 20, 10]

The array has been successfully reversed. Shared memory acted as a temporary staging area, and __syncthreads() was the critical coordinator that made the communication safe.

4. A Glimpse into Professional Tracing: CUDA-GDB

While manual tracing is essential for learning and exams, in the real world you would use a debugger. NVIDIA provides CUDA-GDB, an extension of the GDB debugger specifically for GPUs. It allows you to do exactly what we just did manually:

• Set breakpoints inside a kernel.
• Switch between threads (cuda thread (1,0,0)).
• Inspect the values of variables (print i).
• Examine the contents of global and shared memory (print my_array[0]@4).

This turns the abstract concept of tracing into a concrete, interactive process. While we won't do a full tutorial, knowing such tools exist connects your theoretical knowledge to professional practice.

Conclusion

In this lesson, we transitioned from the distributed-memory model of MPI to the massively parallel device model of CUDA. We saw that tracing a CUDA kernel is a systematic process of tracking state across thousands of threads and multiple memory spaces.

Key Takeaways:

• Thread Indexing is Key: The first step is always to calculate the unique global index for each thread to understand what data it is responsible for.
• Global vs. Shared Memory: Tracing requires you to maintain a mental model of at least two separate memory spaces: the large, persistent global memory and the small, fast, per-block shared memory.
• Synchronization is Mandatory: When using shared memory for inter-thread communication, __syncthreads() is not an optimization—it is essential for correctness. It acts as a barrier, ensuring that operations happen in the correct logical order.
• Tracing is a Step-by-Step Process: Follow the code line by line, updating the state of memory for each group of threads as they execute instructions, paying special attention to synchronization points.

In our next lesson, we will synthesize what we've learned about OpenMP, MPI, and CUDA. You will learn to select and justify the most appropriate parallel model for various problem scenarios, moving from analyzing existing code to making high-level design decisions.