10.2e Shared memory and L1 cache within the SM: the scratchpad memory model

📦 Nvidia GPU Architecture 📖 GPU Microarchitecture

🧭 Context Introduction

When you write a program for an NVIDIA GPU, your code runs on Streaming Multiprocessors (SMs). Each SM has its own memory hierarchy, and two of the most important — and most confusing — components are shared memory and L1 cache. Together, they form what is called the scratchpad memory model.

Think of it this way: L1 cache is like a small, fast, automatic notepad that the hardware manages for you. Shared memory is like a whiteboard that you, the programmer, control explicitly. Both live in the same physical on-chip memory space inside the SM, and you can decide how much of that space is used for each purpose.

⚙️ What is the Scratchpad Memory Model?

The scratchpad memory model refers to the programmer-managed, on-chip memory inside an SM. Unlike a traditional CPU cache (which is fully hardware-managed), the scratchpad gives engineers direct control over data placement and reuse.

  • Scratchpad memory is a small, fast memory region located physically on the GPU chip.
  • It is software-managed — you explicitly load data into it and read from it.
  • It is shared among all threads running in the same thread block (also called a Cooperative Thread Array, or CTA).
  • It provides low latency (a few clock cycles) and high bandwidth compared to global memory (the large off-chip DRAM).

🧠 Shared Memory vs. L1 Cache: The Key Difference

Both shared memory and L1 cache occupy the same physical on-chip SRAM inside the SM. However, they behave very differently.

Feature Shared Memory L1 Cache
Who controls it? Programmer (explicit) Hardware (automatic)
Visibility All threads in a thread block Single thread (per SM)
Latency ~5–10 clock cycles ~20–30 clock cycles
Size (typical) Configurable (e.g., 0–48 KB per SM) Configurable (e.g., 0–48 KB per SM)
Use case Data sharing, reduction, tiling Automatic data reuse, latency hiding
Coherency Not coherent across thread blocks Coherent within a single thread

🛠️ How the Partitioning Works

NVIDIA GPUs allow you to partition the on-chip memory between shared memory and L1 cache. This is done at kernel launch time.

  • Example partitioning options (per SM):
  • 48 KB shared memory + 16 KB L1 cache
  • 32 KB shared memory + 32 KB L1 cache
  • 16 KB shared memory + 48 KB L1 cache

  • 0 KB shared memory + 64 KB L1 cache (on some architectures)

  • The partitioning is set using a compiler hint or API call (e.g., cudaFuncSetAttribute).

  • The choice depends on your workload:
  • Use more shared memory if your kernel does a lot of explicit data sharing between threads.
  • Use more L1 cache if your kernel benefits from automatic caching of frequently accessed global memory data.

📊 Visual Representation: SM Shared Memory and L1 Cache Allocation

This diagram displays how local SRAM space within each Streaming Multiprocessor is partitioned between hardware-managed L1 cache and programmer-managed Shared Memory.

flowchart LR SRAM["Physical SM SRAM"] --> L1["L1 Cache (Hardware Managed)"] SRAM --> Shared["Shared Memory (Programmer Managed / Thread block local)"] Shared --> Banks["32-Bank Structure (Avoid Bank Conflicts)"] classDef cpu fill:#eafaf1,stroke:#76b900,stroke-width:2px,rx:6px,ry:6px; classDef memory fill:#f0f7ff,stroke:#3498db,stroke-width:1.5px,rx:4px,ry:4px; classDef system fill:#f1f5f9,stroke:#64748b,stroke-width:1.5px; class SRAM cpu; class L1,Shared memory; class Banks system;

🕵️ Why This Matters for AI Workloads

In AI inference and training, the scratchpad memory model is critical for performance.

  • Matrix multiplication (GEMM): AI kernels like those in cuBLAS or custom fused kernels use shared memory to hold tiles of input matrices. This reduces global memory traffic by a factor of 10x or more.
  • Convolution: Shared memory stores input feature map tiles and filter weights, enabling massive data reuse.
  • Attention mechanisms: In transformer models, shared memory is used to store query, key, and value blocks during attention computation.

Without explicit use of shared memory, your AI kernel would be constantly reading from slow global memory, and performance would drop dramatically.

📊 A Simple Mental Model

Imagine you are cooking in a large kitchen (the GPU).

  • Global memory is the pantry — huge, but slow to walk to.
  • L1 cache is your countertop — the hardware automatically puts ingredients you use often there.
  • Shared memory is a small prep table that you explicitly load with ingredients you know all your sous-chefs (threads) will need.

The scratchpad memory model is the prep table — you decide what goes on it, when, and for how long.

🔍 Key Takeaways for New Engineers

  • Shared memory and L1 cache share the same physical on-chip memory inside the SM.
  • Shared memory is programmer-managed; L1 cache is hardware-managed.
  • You can partition the on-chip memory between them at kernel launch.
  • Using shared memory correctly is one of the most important optimization techniques for GPU programming.
  • For AI workloads, shared memory enables data reuse and reduced global memory traffic, which directly translates to faster training and inference.

📚 Further Exploration

To practice using the scratchpad memory model, try writing a simple matrix transpose or vector reduction kernel in CUDA. Use shared memory to store a tile of data, synchronize threads with __syncthreads(), and observe the performance improvement over a naive global memory-only implementation.

Remember: the scratchpad is your friend — but only if you use it wisely.