GitHub - houslast3/How-I-Squeezed-31.9-TFLOPS-out-of-an-RTX-4060-Beating-cuBLAS-by-14- · GitHub
/" data-turbo-transient="true" />
Skip to content
Search or jump to...
Search code, repositories, users, issues, pull requests...
-->
Search
Clear
Search syntax tips
Provide feedback
--><br>We read every piece of feedback, and take your input very seriously.
Include my email address so I can be contacted
Cancel
Submit feedback
Saved searches
Use saved searches to filter your results more quickly
-->
Name
Query
To see all available qualifiers, see our documentation.
Cancel
Create saved search
Sign in
/;ref_cta:Sign up;ref_loc:header logged out"}"<br>Sign up
Appearance settings
Resetting focus
You signed in with another tab or window. Reload to refresh your session.<br>You signed out in another tab or window. Reload to refresh your session.<br>You switched accounts on another tab or window. Reload to refresh your session.
Dismiss alert
{{ message }}
houslast3
How-I-Squeezed-31.9-TFLOPS-out-of-an-RTX-4060-Beating-cuBLAS-by-14-
Public
Notifications<br>You must be signed in to change notification settings
Fork
Star
main
BranchesTags
Go to file
CodeOpen more actions menu
Folders and files<br>NameNameLast commit message<br>Last commit date<br>Latest commit
History<br>1 Commit<br>1 Commit
scripts
scripts
src
src
README.md
README.md
beating-cublas-rtx4060.md
beating-cublas-rtx4060.md
View all files
Repository files navigation
CUTLASS/WMMA GEMM Benchmarks — NVIDIA RTX 4060 (Ada Lovelace)
Systematic benchmark of FP16→FP32 GEMM optimization strategies on a consumer-grade Ada Lovelace GPU, comparing cuBLAS, CUTLASS device-level API (custom pipeline depths and tile sizes), and a hand-tuned WMMA Tensor Core kernel.
Key Results
Configuration<br>TFLOPS<br>% Peak<br>vs cuBLAS
CUTLASS S=3, 256×128, epi4<br>31.9<br>52.7%<br>+1.3%
CUTLASS S=3, 256×128, epi1<br>31.4<br>51.9%<br>–0.3%
CUTLASS S=4, 256×128, epi4<br>31.1<br>51.4%<br>–1.3%
CUTLASS S=4, 128×128, epi4<br>31.0<br>51.2%<br>–1.6%
cuBLAS (profiler ref)<br>~31.5<br>52.0%<br>100%
WMMA custom (N=4096)<br>25.1<br>41.5%<br>~89%
Central finding: Deepening the pipeline from S=3 to S=4 decreases throughput by 2.5% (31.9→31.1 TFLOPS), contradicting simple analytical models. The cause is occupancy collapse (2→1 block/SM) driven by shared memory increase (72→96 KB) and register pressure.
Repository Structure
repo/<br>├── src/<br>│ ├── cutlass_custom.cu — CUTLASS device-level API benchmark harness<br>│ ├── gemm_ultimate.cu — Custom WMMA Tensor Core kernel + validation<br>│ ├── final_benchmark.cu — Multi-kernel comparison (V, W, X kernels)<br>│ ├── cutlass_validate.cu — Cross-validation across problem sizes<br>│ ├── validation_gemm.cu — Numerical validation vs cuBLAS<br>│ ├── explore_cutlass.py — Exploration of 1000+ CUTLASS configs<br>│ └── cutlass_eq.py — Analytical model calibration<br>├── scripts/<br>│ ├── build.bat — Windows compilation script<br>│ └── run.bat — Benchmark execution script<br>└── README.md
Requirements
NVIDIA GPU with Ada Lovelace architecture (sm_89) or compatible
CUDA 13.1+ (with cuBLAS)
MSVC 2022 BuildTools (or Visual Studio 2022)
CUTLASS 4.6.0
CMake + Ninja (for CUTLASS build)
Python 3.8+ (for exploration scripts)
Setup
1. Install CUTLASS
git clone https://github.com/NVIDIA/cutlass.git<br>cd cutlass<br>mkdir build && cd build<br>cmake .. -G Ninja -DCUTLASS_NVCC_ARCHS=89<br>ninja<br>set CUTLASS_ROOT=%cd%\..
2. Build Benchmarks
cd repo\scripts<br>set CUTLASS_ROOT=C:\path\to\cutlass<br>build.bat
Usage
cd repo\scripts
# All benchmarks<br>run.bat
# Individual benchmarks<br>run.bat cutlass_custom # CUTLASS API (N=8192)<br>run.bat gemm_ultimate # WMMA kernel (N=4096)<br>run.bat final_benchmark # Multi-kernel comparison (N=8192)
# Custom sizes<br>run.bat cutlass_custom 4096 # CUTLASS at N=4096<br>run.bat gemm_ultimate 2048 # WMMA at N=2048
CUTLASS Device-Level API
The harness in cutlass_custom.cu instantiates GEMM kernels via cutlass::gemm::device::Gemm with compile-time template parameters:
Pipeline depth (Stages): 2, 3, 4
Threadblock tile (CtaM, CtaN): 128×128, 256×128
Epilogue vectorization: 1, 4
Instruction shape: GemmShape (s16816)
Warp shape: GemmShape
Custom WMMA Kernel
gemm_ultimate.cu implements a 128×128 threadblock tile with 512 threads (16 warps), using the WMMA API directly. It includes a three-stage validation framework:
Launch error detection (cudaGetLastError)
Asynchronous error detection (cudaDeviceSynchronize)
Element-wise comparison against cuBLAS reference
Analytical Model
The calibrated performance model:
GFLOPS(S, M_CTA, N_CTA) = 60,550 × S/(S+2.56) × min(1, M_CTA·N_CTA/32,768) × 0.96
Accurate for S=2 (+0.4%) and S=3 (–1.5%), overestimates S=4 by +13.8%. See cutlass_eq.py for details.
Known Issues
Windows TDR-like crash: Interleaving CUTLASS and cuBLAS calls in the same process causes access violation 0xC0000005 after ~10 iterations. Run benchmarks in separate invocations.
S=4 constraint: Max shared memory per SM is 100 KB (opt-in). S=4 with 256×128 tile requires 96 KB, limiting to 1...