Repository navigation
V1. Shared Memory Tiling
(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.
is to allow multiple threads within a block to reuse data, thereby reducing the number of accesses to global memory.
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.
Tip
It uses the y-coordinates of the block and thread to determine which row of matrix A 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.
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
# ./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 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.01best speedup: 2.27x
ncu is the command-line interface (CLI) tool for NVIDIA Nsight Compute.
# 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 2048sudo 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
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
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."
SM Busy = 20.45%，Issue Slots Busy = 17.80%
—— The computing units are also significantly underutilized.
—— Most of the time is spent waiting.
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.
# 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 fasterV1 resolved the issue (shared memory reuse eliminated most global memory stalls)