optimize-shared-memory-tiling

Skill: Optimize Shared Memory Tiling

Purpose

Guide the agent through designing and tuning shared memory tiling strategies for CUDA kernels, covering bank conflict analysis and elimination, tile shape selection, double buffering with async copy, occupancy tradeoffs from shared memory allocation, and the decision of when smem tiling is worth the complexity.

Use this when

• A kernel repeatedly reads the same global memory data from multiple threads and would benefit from staging through a shared memory tile (GEMM, convolution, stencil)
• Profiling shows high shared memory bank conflict rates in Nsight Compute (l1tex__data_bank_conflicts_pipe_lsu_mem_shared)
• You are designing the smem layout for a GEMM or attention tiling kernel and need to choose tile dimensions and padding
• You are adding double buffering to overlap global memory loads with computation using cp.async on SM80+
• An existing kernel has smem usage that limits occupancy and needs to be restructured

Do not use this when

• The access pattern is truly random (no spatial reuse) and shared memory staging will not increase the reuse factor
• The data is small enough to fit in L1 cache across all accesses without explicit smem management: compiler-managed L1 may be sufficient
• The kernel is compute-bound and memory latency is not the bottleneck: smem tiling adds complexity without throughput benefit
• The kernel is a simple elementwise operation with no data reuse: smem staging provides no benefit

Inputs the agent should gather first

• Kernel type and access pattern: describe the data reuse structure — which threads reuse which elements, and along which dimension. E.g., in GEMM, each row of threads reuses a row of A and each column of threads reuses a column of B.
• Tile shape context: what are the thread block dimensions (BM, BN, BK) or equivalent? How many threads are in the block?
• Dtype and element size: fp32 (4 bytes), fp16 (2 bytes), int8 (1 byte) — affects how many elements map to a single bank.
• Hardware target: SM architecture. On all current NVIDIA GPUs (Kepler through Hopper): 32 banks, 4-byte bank width (by default; can be configured to 8-byte via cudaDeviceSetSharedMemConfig). SM80+ has up to 164 KB smem per SM in some configurations.
• Current smem usage: how many bytes of smem are currently allocated per block? How does this affect occupancy?
• Whether async copy is applicable: SM80+ with cp.async, SM90 with cp.async.bulk (TMA). Is the kernel targeting those architectures?