V4. warp tensorcore - HuajingLin/Explore-GPU-performance GitHub Wiki
Basic Concepts
- Tensor Cores are computing units within NVIDIA GPUs designed specifically for matrix multiplication and related tensor operations.
- Tensor Cores can execute these types of MMA (Matrix Multiply-Accumulate) operations very efficiently.
- $\color{red}{\text{WMMA}}$ (Warp Matrix Multiply and Accumulate) can be understood as a relatively high-level Tensor Core programming interface provided by NVIDIA CUDA.
the Kernel
1. run_wmma_gemm
dim3 blockDim(128, 4);
dim3 gridDim((M + (WMMA_M * (blockDim.x / 32)) - 1) / (WMMA_M * (blockDim.x / 32)),
(N + (WMMA_N * blockDim.y) - 1) / (WMMA_N * blockDim.y));
wmma_naive_gemm_kernel<<<gridDim, blockDim>>>(d_A, d_B, d_C, M, N, K);
128 × 4 = 512 threads
512 / 32 = 16 warps
(1) Every 32 threads form a warp:
warp 0 : y=0, x=0 ~ 31
warp 1 : y=0, x=32 ~ 63
warp 2 : y=0, x=64 ~ 95
warp 3 : y=0, x=96 ~ 127
warp 4 : y=1, x=0 ~ 31
warp 5 : y=1, x=32 ~ 63
warp 6 : y=1, x=64 ~ 95
warp 7 : y=1, x=96 ~ 127
warp 8 : y=2, x=0 ~ 31
...
warp 15 : y=3, x=96 ~ 127
So it looks like this:
blockDim.x = 128
┌──────┬──────┬──────┬──────┐
y = 0 │ W0 │ W1 │ W2 │ W3 │
├──────┼──────┼──────┼──────┤
y = 1 │ W4 │ W5 │ W6 │ W7 │
├──────┼──────┼──────┼──────┤
y = 2 │ W8 │ W9 │ W10 │ W11 │
├──────┼──────┼──────┼──────┤
y = 3 │ W12 │ W13 │ W14 │ W15 │
└──────┴──────┴──────┴──────┘
↑
4 warps in x direction
(2) grid
dim3 gridDim((M + (WMMA_M * (blockDim.x / 32)) - 1) / (WMMA_M * (blockDim.x / 32)),
(N + (WMMA_N * blockDim.y) - 1) / (WMMA_N * blockDim.y));
# equal:
gridDim.x = ceil(M / 64);
gridDim.y = ceil(N / 64);
A block is responsible for a 64×64 tile of C.
V4.1 shared memory
Output example:
[verify naive wmma kernel] and [verify tiled wmma kernel] both PASS
[compare] Improvement of tiled over naive: $\color{red}{\text{1.65x}}$
[compare] tiled kernel achieved $\color{red}{\text{9.5%}}$ of cuBLAS performance
=== V4 WMMA Tensor Core GEMM (FP16 input, FP32 accumulation) ===
M=2048 N=2048 K=2048
[verify] N > 512; skipping full CPU verification (too slow).
[bench] Custom WMMA kernel: avg_time=7.832 ms, GFLOPS=2193.51
[bench] tiled WMMA kernel: avg_time=4.885 ms, GFLOPS=3516.89
[compare] tiled over naive up: 1.60x
[bench] cuBLAS (TensorCore): avg_time=0.419 ms, GFLOPS=40992.53
[compare] naive kernel achieved 5.4% of cuBLAS performance
[compare] tiled kernel achieved 8.6% of cuBLAS performance
(Acceptance criteria: >= 70% for N>=1024 is considered passing)
V4.2 tiled+dbuf WMMA kernel
/content/gpu_perf# ./build/v4_wmma 2048
=== V4 WMMA Tensor Core GEMM (FP16 input, FP32 accumulate) ===
M=2048 N=2048 K=2048
[verify] N>512 skip
[bench] naive WMMA kernel: avg_time=11.558 ms, GFLOPS=1486.36
[bench] tiled WMMA kernel: avg_time=9.722 ms, GFLOPS=1767.03
[compare] tiled over naive promote: 1.19x
[bench] tiled+dbuf WMMA kernel: avg_time=6.726 ms, GFLOPS=2554.20
[compare] tiled+dbuf over tiled promote: 1.45x
[bench] cuBLAS(TensorCore): avg_time=0.427 ms, GFLOPS=40256.97
[compare] naive kernel achieving cuBLAS performance 3.7%
[compare] tiled kernel achieving cuBLAS performance 4.4%
[compare] tiled+dbuf kernel achieving cuBLAS performance 6.3%
(Acceptance Criteria: N>=1024 >= 70% Deemed to meet the standard.)
V4.3 Increase TILE_K from 16 to 64.
=== V4 WMMA Tensor Core GEMM (FP16 input, FP32 accumulation) ===
M=2048 N=2048 K=2048
[verify] N > 512; skipping full CPU verification (too slow).
[bench] Custom WMMA kernel: avg_time=7.832 ms, GFLOPS=2193.51
[bench] tiled WMMA kernel: avg_time=4.885 ms, GFLOPS=3516.89
[compare] tiled over naive promote: 1.60x
[bench] cuBLAS (TensorCore): avg_time=0.419 ms, GFLOPS=40992.53
[compare] naive kernel achieved 5.4% of cuBLAS performance
[compare] tiled kernel achieved 8.6% of cuBLAS performance
(Acceptance criteria: >= 70% for N>=1024 is considered passing)
change TILE_K from 64 to 32.
- The results show: TILE_K=16 is currently the optimal configuration.
#bench and plot result
chmod +x bench/bench.sh
bench/bench.sh
chmod +x ./bench/plot_reuslts.py
./bench/plot_results.py
V4.4 wmma + warp + tile
/content/gpu_perf# ./build/v4_wmma 1024
=== V4 WMMA Tensor Core GEMM (FP16 input, FP32 accumulate) ===
M=1024 N=1024 K=1024
[verify] N>512 skip
[bench] naive WMMA kernel: avg_time=0.988 ms, GFLOPS=2172.67
[bench] tiled WMMA kernel: avg_time=0.938 ms, GFLOPS=2290.48
[compare] tiled over naive promote: 1.05x
[bench] tiled+dbuf WMMA kernel: avg_time=1.192 ms, GFLOPS=1800.89
[compare] tiled+dbuf over tiled promote: 0.79x
- [bench] warptile WMMA kernel: avg_time=0.535 ms, GFLOPS=4010.62
[compare] warptile over tiled+dbuf promote: 1.75x
[bench] cuBLAS(TensorCore): avg_time=0.092 ms, GFLOPS=23458.08
[compare] naive kernel achieving cuBLAS performance 9.3%
[compare] tiled kernel achieving cuBLAS performance 9.8%
[compare] tiled+dbuf kernel achieving cuBLAS performance 7.7%
- [compare] warptile kernel achieving cuBLAS performance 17.1%
(Acceptance Criteria: N>=1024 >= 70% Deemed to meet the standard.)
WARP_TILE_M change to 4
/content/gpu_perf# ./build/v4_wmma 1024
=== V4 WMMA Tensor Core GEMM (FP16 input, FP32 accumulate) ===
M=1024 N=1024 K=1024
[verify] N>512 skip
[bench] naive WMMA kernel: avg_time=1.496 ms, GFLOPS=1435.21
[bench] tiled WMMA kernel: avg_time=1.382 ms, GFLOPS=1553.94
[compare] tiled over naive promote: 1.08x
[bench] tiled+dbuf WMMA kernel: avg_time=1.761 ms, GFLOPS=1219.31
[compare] tiled+dbuf over tiled promote: 0.78x
- [bench] warptile WMMA kernel: avg_time=0.658 ms, GFLOPS=3263.46
[compare] warptile over tiled promote: 2.10x
[bench] cuBLAS(TensorCore): avg_time=0.133 ms, GFLOPS=16116.44
[compare] naive kernel achieving cuBLAS performance 8.9%
[compare] tiled kernel achieving cuBLAS performance 9.6%
[compare] tiled+dbuf kernel achieving cuBLAS performance 7.6%
- [compare] warptile kernel achieving cuBLAS performance 20.2%
(Acceptance Criteria: N>=1024 >= 70% Deemed to meet the standard.)
Conclusion:
Measurement noise has become so significant that the question of "which configuration is superior" is itself rendered unreliable.