Tiled CUDA GEMM: shared memory vs naive global memory

Shared Memory Workbench
B16×16 global memory (K × N)
BARRIER
__syncthreads()
idle
s_B

No shared memory. Without a __shared__ array, each load goes from global memory (via the caches) straight into one thread's register.

s_A
k
A16×16 global memory (M × K)
C16×160/64 blocks
A row strip B column strip active K-tile selected block register acc written to C __syncthreads()

Naive: every thread walks global memory

One thread per output element. Thread (ty, tx) streams a whole row of A and column of B straight from global memory, the slow “main pantry”.


      
    
      
    It uses no shared memory: without a __shared__ array each value goes from global memory into one thread's register and is dropped after its multiply-add. The counts here are per-thread load requests. The hardware softens the redundancy: loads of the same address within a warp merge into one memory transaction, and the L1 and L2 caches can serve repeats. On Volta and later, L1 is the same on-chip SRAM as shared memory. But what stays cached is up to the hardware, so the reuse is not guaranteed. NVIDIA's Best Practices Guide measures 120 GB/s for a kernel like this on a V100 and 196 GB/s once both tiles are staged in shared memory.
    
    
    
    
    
      
    Tiled: load, sync, compute, sync
    
      
    The block fetches a 2×2 tile of A and of B once into shared memory (the “workbench”), then every thread computes from there.
    
      
      
      
      
        
      Global memory (DRAM)Visible to every block. Hundreds of cycles per access.
      
        
      Shared memory (SRAM)Per block, on-chip. ~20–30 cycles per access.
      
      
      
    
      
    
      
      
      Reuse grows with the tile size
      
      
      
      
      
      
      Shared memory per block is 2·T²·4 B: 2 KB at T = 16, 8 KB at T = 32. For this one-thread-per-output kernel the real cap is threads per block (T² ≤ 1024, so T ≤ 32); register-blocked kernels use bigger tiles, where shared-memory capacity and occupancy become the limit.