-
Notifications
You must be signed in to change notification settings - Fork 0
Expand file tree
/
Copy pathsummary.txt
More file actions
115 lines (94 loc) · 3.46 KB
/
Copy pathsummary.txt
File metadata and controls
115 lines (94 loc) · 3.46 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
llm-fusion-compiler
Author: Joao Felipe De Souza
Year: 2026
Overview
--------
Compiler MVP that detects fusion patterns in Transformer operations,
generates optimized CUDA kernels via Jinja2 templates, compiles them
with nvcc, and executes them on real GPU hardware.
The compiler pipeline is fully functional end-to-end:
IR graph -> pattern fusion -> CUDA codegen -> nvcc compile -> GPU execute
Hardware
--------
Target: SM75 (Turing) -- RTX 2070
Tensor Cores: FP16 input, FP32 accumulate (WMMA API)
Peak: 28.5 TFLOPs (FP16in/FP32acc)
Memory BW: 448 GB/s
Compiler Pipeline
-----------------
1. IR: TensorDesc, Op, Graph, GraphBuilder
Programmatic API to construct dataflow graphs of Transformer ops.
2. Pattern Matcher: 9 fusion patterns with priority-based selection.
Topological traversal, linear chain detection, no double-fusion.
3. Code Generator: Jinja2 templates -> CUDA source files.
WMMA API with vectorized float4 global loads,
register-resident accumulators, fused epilogue.
4. NVCC Wrapper: programmatic compilation to .so shared libraries.
5. Auto-Tuner: grid search over BM/BN/BK/threads with hardware constraints.
Shared memory budget validation (48KB SM75 limit).
Persistent JSON cache per (pattern, shape, arch).
Fusion Patterns
---------------
gemm_bias_gelu: GEMM + Bias + GeLU (fast approximation)
gemm_bias_silu: GEMM + Bias + SiLU
gemm_bias_relu: GEMM + Bias + ReLU
gemm_bias: GEMM + Bias
gemm_bias_residual: GEMM + Bias + Residual add
gemm_gelu: GEMM + GeLU
gemm_silu: GEMM + SiLU
layernorm_gemm: LayerNorm + GEMM
rmsnorm_gemm: RMSNorm + GEMM (two-kernel design)
Real Kernel Performance
-----------------------
Generated GEMM+Bias+GeLU fused kernel on RTX 2070:
Config: BM=64 BN=64 BK=32 threads=256
Shape Gen kernel TFLOPs vs Peak Max error
M=512 0.48ms 8.9 31.2% 0.002
M=1024 3.35ms 10.3 36.1% 0.004
M=2048 6.59ms 10.4 36.6% 0.004
M=4096 13.30ms 10.5 36.7% 0.004
Correctness: PASS at all shapes (max abs error < 0.004)
Optimization Journey
--------------------
V1 (scalar loads, smem accum): 3.0 TFLOPs (10.5%)
V2 (register accum): 3.1 TFLOPs (10.9%)
V3 (vectorized float4 loads): 10.5 TFLOPs (36.7%) <- 3.5x improvement
Vectorized global memory loads were the highest-impact optimization.
Performance Context
-------------------
cuBLAS (via PyTorch) at M=4096: 27 TFLOPs (reference)
Generated kernel at M=4096: 10.5 TFLOPs
Ratio: 0.39x cuBLAS
This is expected for a first-generation WMMA template without:
- Software pipelining (cp.async unavailable on SM75)
- Double buffering (48KB smem limit)
- Register-level tiling (CUTLASS-style)
The compiler infrastructure is the contribution.
Test Coverage
-------------
All tests pass:
- IR tests: 7/7
- Fusion tests: 11/11
- Codegen tests: 7/7
- Tuner tests: 7/7
- Correctness tests: 5/5 (24/24 sub-tests)
- Pipeline tests: 2/2
End-to-End Pipeline
-------------------
FFN block (SwiGLU style):
IR: 7 ops
Fusion: 4 ops (gemm_silu + gemm_bias_residual)
Codegen: 2 fused CUDA kernels generated and compiled
Tuner: 2 configs selected
Limitations
-----------
- Forward pass only
- FP16 input, FP32 accumulate only
- SM75 target (no cp.async, no persistent kernels)
- Programmatic IR only (no PyTorch/ONNX parser)
- Generated kernels at 37% peak (vs cuBLAS 95%)
- 48KB smem limit constrains tile sizes
Author
------
Joao Felipe De Souza
2026