V1. Shared Memory Tiling - HuajingLin/Explore-GPU-performance GitHub Wiki
1. Shared memory in CUDA operates at the block level.
(1) When shared memory is allocated, an instance of that memory (of the same size) is created within each block.
(2) The __syncthreads() operation ensures synchronization only among threads within that specific block and synchronizes the shared memory within that block.
2. An important use of storing data in shared memory
is to allow multiple threads within a block to reuse data, thereby reducing the number of accesses to global memory.
3. Improvement over V0:
Introduces tiling by loading sub-blocks of A and B into shared memory for reuse, thereby eliminating a large number of redundant global memory reads. Each block is responsible for computing a TILE_SIZE x TILE_SIZE sub-block of C, iterating along the K dimension to load tiles and accumulate results.
4. Core processes
[!TIP]
It uses the y-coordinates of the block and thread to determine which row of matrix A the thread is to process.
It uses the x-coordinates of the block and thread to determine which column of matrix B the thread is to process.
It also determine the number at which row and column in matrix C the thread is to process.
(1) Construct a 2D array of threads corresponding to matrix C, where each thread corresponds to a single element.
(2) A single element of matrix C = a row of matrix A Γ a column of matrix B.
(3) Use shared memory to reduce the number of accesses.
$\color{red}{\text{Each thread copies one number;}}$ with 32 Γ 32 = 1,024 threads, 1,024 numbers can be moved simultaneously β
synchronize (Copying complete) β
compute (using data in shared memory) β
synchronize (computing complete) β
write results back to C.
A (m Γ k) B (k Γ n) C (m Γ n)
βββββββββββββ βββββββββββββ βββββββββββββ
β β β β β β
β β β β β β
β β Γ β β = β β
β β β β β β
β β β β β β
βββββββββββββ βββββββββββββ βββββββββββββ
β β β β β β
m k k n m n
rows columns rows columns rows columns
5. Sample output (Google Colab, T4)
# ./build/v1_naive32 1024
=== V1 Shared Memory Tiling GEMM (TILE_SIZE=32) ===
M=1024 N=1024 K=1024
shared memory per block = 8192 bytes (limit 49152B)
[verify] N > 512, skipping full CPU verification (too slow); suggest verifying correctness with a smaller scale separately
[bench] avg_time=0.526 ms, GFLOPS=566.21
6. benchmark data (T4)
N | V0 | V1-tile8 | V1-tile16 | V1-tile32 | Optimal speedup
----|-----|----------|-----------|-----------|---------------
512 | 428 | 256 | 496 | 554 | 1.29Γ
1024 | 414 | 324 | 528 | 566 | 1.36Γ
2048 | 420 | 374 | 565 | 695 | 1.65Γ
4096 | 420 | 379 | 572 | 695 | 1.65Γ
# # chmod +x ./bench/bench.sh
# # ./bench/bench.sh
# # cat results.csv
version,tile_size,N, avg_ms, gflops
V0, NA, 512, 0.771, 348.13
V0, NA, 1024, 6.114, 351.23
V0, NA, 2048, 41.142, 417.57
V0, NA, 4096, 318.617,431.36
V1, 8, 512, 0.504, 532.18
V1, 8, 1024, 4.153, 517.10
V1, 8, 2048, 45.863, 374.59
V1, 8, 4096, 371.252,370.20
V1, 16, 512, 0.347, 773.34
V1, 16, 1024, 2.831, 758.69
V1, 16, 2048, 31.018, 553.87
V1, 16, 4096, 247.247,555.88
V1, 32, 512, 0.341, 788.17
V1, 32, 1024, 2.690, 798.25
V1, 32, 2048, 24.531, 700.32
V1, 32, 4096, 194.393,707.01
best speedup: 2.27x
7. Nsight Compute
ncu is the command-line interface (CLI) tool for NVIDIA Nsight Compute.
(1) capture key indicators
# V0 Key indicators under N=2048
sudo /usr/local/cuda/bin/ncu --metrics \
l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum,\
dram__throughput.avg.pct_of_peak_sustained_elapsed,\
lts__t_sector_hit_rate.pct,\
sm__warps_active.avg.pct_of_peak_sustained_active,\
smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct \
./build/v0_naive 2048
# V1-tile32 Key metrics under N=2048, plus bank conflict
sudo /usr/local/cuda/bin/ncu --metrics \
l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum,\
dram__throughput.avg.pct_of_peak_sustained_elapsed,\
lts__t_sector_hit_rate.pct,\
sm__warps_active.avg.pct_of_peak_sustained_active,\
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum \
./build/v1_tile32 2048
(2) capture full metrics
sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH" \
/usr/local/cuda/bin/ncu --set full -o v0_profile ./build/v0_naive 2048
# Save to a file, retrieve the first 100 lines.
sudo env "PATH=$PATH" "LD_LIBRARY_PATH=/usr/lib64-nvidia:$LD_LIBRARY_PATH" \
/usr/local/cuda/bin/ncu --import v1_profile.ncu-rep --page details > v1_report.txt
head -100 v1_report.txt
$\color{red}{\text{Focus on these fields:}}$
dram__throughput...pct_of_peak β DRAM bandwidth utilization
lts__t_sector_hit_rate.pct β L2 cache hit rate
sm__warps_active...pct_of_peak β occupancy
l1tex__data_bank_conflicts... β shared memory bank conflict count
(3) $\color{red}{\text{Interpretation of Key Figures}}$ for V0
Memory-related
L1/TEX Hit Rate = 87.26%οΌL2 Hit Rate = 50.73%:
ββ There are a large number of cache hits; the hardware cache has already helped reuse a significant amount of data.
Memory Throughput = 58.01 GB/s, Mem Busy = 45.95%:
ββ The T4 has a theoretical memory bandwidth of approximately 320 GB/s; at 58 GB/s, the figure is less than one-fifth of the peakβfar from saturating the bandwidthβindicating that V0 is by no means "bandwidth-bound."
Compute-related
SM Busy = 20.45%οΌIssue Slots Busy = 17.80%
ββ The computing units are also significantly underutilized.
$\color{red}{\text{The most crucial point:}}$ Every scheduler...issues an instruction every 5.6 cycles, Eligible Warps Per Scheduler = 0.73οΌ8 is maxοΌ
ββ Most of the time is spent waiting.
(4) $\color{red}{\text{Interpretation of Key Figures}}$ for V1
L1/TEX Cache Throughput = 89.78%οΌCompute (SM) Throughput = 89.26%
ββ The L1/TEX pipeline handling shared memory is heavily saturated (approaching 90% utilization).
DRAM Throughput is 9.60%οΌL2 Cache Throughput is 4.52%
ββ global memory / L2 has become virtually idle, demonstrating that shared memory tiling successfully achieved its design objective: shifting memory access pressure from global memory to shared memory.
But even so, achieved 10% of fp32 peak performance
ββ Computing power utilization remains very low.
The most crucial clueβquoting the report directly:
On average, each warp spends 26.9 cycles stalled while waiting for space to become available in the MIO (memory input/output) instruction queue; such stalls typically occur when the MIO pipeline (including shared memory accesses) is oversubscribed.
(5) Comparison of Key metrics (N=2048)
# Retrieve the complete data using the command:
/content/gpu_perf# grep -A 12 "GPU Speed Of Light Throughput" v0_report.txt
grep -A 12 "GPU Speed Of Light Throughput" v1_report.txt
Section: GPU Speed Of Light Throughput
Metric Name Metric Unit Metric Value
----------------------- ----------- -------------
DRAM Frequency Ghz 5.00
SM Frequency Mhz 585.00
Elapsed Cycles cycle 43,843,535
Memory Throughput % 61.28
DRAM Throughput % 18.13
Duration ms 74.95
L1/TEX Cache Throughput % 91.93
L2 Cache Throughput % 6.67
SM Active Cycles cycle 43,806,988.88
...
metrics | V0 (naive) | V1 (tile32) | change
---------------------- | ---------- | ----------- | ---------------
Memory Throughput | 61.3% | 89.2% | β Obviously busier
DRAM Throughput | 18.1-18.6% | 9.6-9.8% | β Dropped by nearly half
L1/TEX Cache Throughput | ~92% | ~89.8% | Basically on par (both very high)
L2 Cache Throughput | 6.7-6.9% | 4.5-4.6% | β Lower
Duration | 74.9-75ms | 53.3-54.7ms | β Obviously faster
Conclusion
V1 resolved the issue (shared memory reuse eliminated most global memory stalls) $\color{red}{\text{but introduced a new bottleneck:}}$ the density of shared memory access instructions saturated the MIO queue, causing warps to stall while waiting for space in the shared memory instruction queue rather than for global memory data.