Skip to main content
Back to path

Welcome to a Grasp lesson

Create a free account to keep studying and create your own study path.

Lesson 30

Optimizing Data Access with Shared Memory

Recommended resources

Introduction

In our last lesson, you learned how to optimize access to global memory by ensuring your threads' requests are coalesced. This is a critical first step, as it maximizes the efficiency of the available memory bandwidth. However, even at 100% efficiency, global memory remains the primary performance bottleneck in many kernels due to its high latency. The most effective optimization is often not just to access global memory efficiently, but to access it less frequently.

This lesson introduces the key hardware feature that allows us to do just that: shared memory. We will explore how this small, fast, on-chip memory space acts as a programmer-controlled cache. By the end of this lesson, you will be able to declare and use shared memory within your kernels to explicitly cache data, enabling data reuse among threads in a block and dramatically reducing your reliance on slow global memory.

What is Shared Memory and Why is it So Fast?

The performance characteristics of a memory system are dictated by physics and economics. Memory that is physically closer to the processing units has lower latency (it takes less time to access) and can be engineered for higher bandwidth (more data can be moved per second).

This diagram shows the architecture of a single Streaming Multiprocessor (SM). Notice that the "Shared Memory / L1 Cache" is located directly on the chip, alongside the CUDA Cores. This physical proximity is the source of its high performance, in stark contrast to the off-chip Global Memory (DRAM).

Because it is located on the SM, shared memory latency is roughly 100 times lower than uncached global memory. It provides a level of performance much closer to that of registers.

Let's watch a short segment that contrasts the physical implementation and performance of global and shared memory.

02 CUDA Shared Memory

This video, from the "cat blue" channel, provides a clear explanation of where shared and global memory reside physically and why that affects their performance.

Watch the section from Shared Memory Introduction. Pay attention to the distinction between on-chip (shared) and off-chip (global) memory and the performance implications (higher bandwidth, lower latency).

A crucial point mentioned in the video is that shared memory is user-managed. This distinguishes it from an automatic cache like the L1 cache. You, the programmer, have explicit control over what data is loaded into shared memory, when it is loaded, and when it is evicted. This gives you tremendous power to optimize data access patterns for your specific algorithm, a concept that aligns perfectly with your interest in low-level systems programming.

This diagram gives a programmer's perspective on the memory hierarchy. It clearly shows the vast difference in access latency, with shared memory being orders of magnitude faster than global memory.

Declaring and Scoping Shared Memory

Before we can use shared memory, we need to understand its two defining logical characteristics: its scope and its declaration syntax.

  1. Scope: Per-Block Resource
    Shared memory is allocated on a per-thread-block basis. This means all threads within a single block have access to the same region of shared memory. Conversely, threads in different blocks have their own, completely independent shared memories. One block cannot see or modify the shared memory of another. This scope makes it the perfect tool for communication and data sharing among threads that are cooperating on a specific sub-problem.

  2. Declaration: Static vs. Dynamic
    You can declare shared memory in two ways, depending on whether its size is known at compile time.

    • Static Allocation: If the amount of shared memory you need is a compile-time constant, you can declare it directly inside your kernel using the __shared__ keyword.

      __global__ void myKernel(...) { __shared__ float tile[256]; ... }

      Here, tile is an array of 256 floats that will be available in shared memory for each block that executes myKernel.

    • Dynamic Allocation: If the amount of memory needed depends on runtime parameters, you must declare it as an unsized extern array and specify the size during the kernel launch.

      Kernel Code:
      __global__ void myDynamicKernel(...) { extern __shared__ float tile[]; ... }

      Host Code (Kernel Launch):
      int size_in_bytes = num_elements * sizeof(float);
      myDynamicKernel<<<gridDim, blockDim, size_in_bytes>>>(...);

      The third parameter in the <<<...>>> launch configuration tells the CUDA runtime how many bytes of shared memory to allocate for each block.

To see these declaration methods in practice, let's turn to a helpful video.

02 CUDA Shared Memory

We'll return to the "cat blue" video to see how the per-block logical scope works and then see a live coding demonstration of static versus dynamic declaration.

First, watch the segment on logical characteristics to solidify your understanding of the per-block scope. Then, jump to the end of the video for a practical demonstration of how to declare dynamic shared memory, contrasting it with the static approach.

The Canonical Pattern: Tiling for Data Reuse

Now we arrive at the core purpose of shared memory: reducing global memory traffic. The most common and powerful technique for this is called tiling or cache blocking.

The fundamental idea is to exploit data reuse. If you know that multiple threads in a block will need to access the same data from global memory, it is far more efficient to have them cooperatively load that data once into fast shared memory and then perform all subsequent reads from there.

This leads to a canonical three-step pattern in many high-performance kernels:

  1. Load: Threads in a block collectively load a "tile" of data from global memory into a shared memory array. Each thread is responsible for loading one or more elements.
  2. Synchronize: Crucially, all threads must wait until the entire tile has been loaded into shared memory before any thread can begin processing it. If a "fast" thread races ahead and tries to use data that a "slower" thread hasn't loaded yet, it will read garbage data, resulting in a race condition. This synchronization is achieved with the __syncthreads() intrinsic. This function acts as a barrier; no thread in a block can proceed past it until all threads in that same block have reached it.
  3. Compute: Once synchronized, all threads can safely access the data tile from fast shared memory to perform their calculations.

Let's examine a complete example that reverses an array. This seems simple, but it demonstrates the pattern perfectly.

Using Shared Memory in CUDA C/C++ | NVIDIA Technical Blog

This article from the official NVIDIA Developer Blog provides a complete, working example of both static and dynamic shared memory usage.

Please read the sections Shared Memory Example and Static Shared Memory. Focus on the staticReverse kernel. Notice the three-step pattern: Load: s[t] = d[t]; Synchronize: __syncthreads(); Compute/Store: d[t] = s[tr]; Also, read the brief section on Dynamic Shared Memory to see how the exact same logic is applied, just with a different declaration and launch syntax. The full code is provided for context.

The array reversal example is a good teaching tool, but its real-world benefit on modern GPUs is minor. A more impactful application, and one central to your goal of mastering AI kernels, is in matrix multiplication. Here, tiling allows a block of threads to load sub-matrices (tiles) of the input matrices A and B into shared memory. Then, they can perform the many dot products required for the output tile C, reusing the values in shared memory repeatedly, which dramatically reduces the total number of global memory reads.

How to Optimize a CUDA Matmul Kernel for cuBLAS-like ...

This article by Sibo Boehm provides an excellent, practical walkthrough of optimizing matrix multiplication. We will look at the first step where shared memory is introduced.

Read the section Kernel 3: Shared Memory Cache-Blocking. Don't worry about all the pointer arithmetic details for now. Focus on the code structure and how it maps to our canonical pattern: Load: As[...] = A[...]; and Bs[...] = B[...]; Synchronize: The first __syncthreads(); Compute: The for loop that calculates tmp. (Note the second __syncthreads(); is to prevent a race condition on the next iteration of the outer loop, a concept we will explore later).

Conclusion

In this lesson, you've moved beyond simply optimizing access to global memory and have learned how to actively avoid it. This is a fundamental leap in writing high-performance kernels. By understanding the GPU's memory hierarchy and taking explicit control of data placement with shared memory, you can unlock significant performance gains.

Key Takeaways:

  • Shared memory is a fast, on-chip, user-managed cache with low latency and high bandwidth.
  • It is a per-block resource, enabling efficient data sharing and communication between threads within a block.
  • It can be declared statically (if the size is a compile-time constant) or dynamically (if the size is determined at runtime).
  • The primary use case is tiling (or cache-blocking) to facilitate data reuse, which drastically reduces expensive global memory traffic.
  • The canonical pattern for using shared memory is: Load -> Synchronize -> Compute.
  • Synchronization using __syncthreads() is not optional; it is essential for preventing race conditions and ensuring correctness.

We have established that synchronization is critical when using shared memory. However, __syncthreads() is just one piece of the puzzle. In our next lesson, we will delve deeper into the topic of synchronization, examining exactly why threads can get out of sync and how __syncthreads() works as a block-level execution barrier to enforce order.