《GPU Mode》
L083
2025
High priority
transcript · failed
Formalized Kernel Derivation — 커널 도출의 형식화
FlashAttention 같은 “손 깎은 커널” 을 수학적으로 한 단계씩 도출해서 같은 결과를 자동으로 만들 수 있는 프레임워크의 가능성. 명세(spec) 한 줄에서 시작해서 tiled / fused / online-softmax 의 형태까지 — 사람의 직관 대신 derivation rule 의 적용으로 — 내려가는 길. 자막이 실패한 강의를 도메인 지식 + 관련 논문으로 재구성한 학습 노트.
kernel derivation
FlashAttention
algebraic transformation
online softmax
tiling
fusion
verification
?
Speaker
미상 · 확인 필요
학계 · 컴파일러 / 형식 의미론 연구자로 추정
§ 01강의가 풀려는 문제· why this lecture exists
“손 깎은 커널” 을 사람 없이 도출할 수 있을까
FlashAttention 같은 커널은 발명품이다 — Tri Dao 가 online softmax 라는 이전부터 알려진 trick 과 tiling 을 결합해 직접 코드를 짰다. 이 강의는 그 “직접” 의 자리를 도전한다. 같은 결과를 형식적 변환의 연속으로 도출할 수 있는가 — 그게 가능하다면 새 hardware 마다, 새 dtype 마다 사람이 다시 깎는 비용이 사라진다.
강의 자막이 실패해 발표자의 정확한 입장은 직접 잡지 못 했다. 다만 “Formalized Kernel Derivation” 이라는 제목이 가리키는 학계의 흐름은 분명하다 — Halide, Exo, ATL, GEMMINI 등 일련의 연구가 spec → schedule → kernel 의 분리, 그리고 schedule 의 형식적 정당성을 다뤄왔다. 이 강의는 그 흐름의 한 응용으로 추정된다.
- 커널 = spec + schedule — 같은 수학적 명세에 대해 여러 schedule 이 가능하고, 각 schedule 은 같은 결과를 보장해야 한다.
- derivation rule — 한 schedule 에서 다른 schedule 로 가는 변환. 각 변환은 의미를 보존해야 (correctness-preserving) 한다.
- FlashAttention 은 derivation 의 결과로 표현 가능 — naive attention 명세에 reorder, tile, fuse, online-aggregate 같은 rule 을 적용한 결과.
자료 한계 · 확인 필요
발표자가 누구인지 — 학계의 어느 그룹(MIT CSAIL, CMU, Stanford 의 ML systems 그룹 등) 인지 — repo / 영상 메타데이터에서 확인되지 않는다. 노트의 입장은 강의가 다룰 법한 형식화의 일반적 형태 다. 실제 발표 내용과 차이가 있을 수 있다.
“손으로 깎은 커널은 발명이지만, 같은 알고리즘에 대한 서로 다른 schedule은 도출의 대상이다.”학습 노트 · 재구성
§ 02커널 derivation 의 동기· spec → kernel 의 거리
spec 한 줄에서 production kernel 까지의 거리는 너무 멀다
attention 의 수학 명세는 한 줄이다 — O = softmax(Q K^T / √d) V. 그런데 production attention 커널 한 개는 수백 줄의 Triton/CUDA 다. 그 사이의 “거리” 안에 무엇이 들어 있는가 — 형식적으로 깐다는 게 이 강의의 입장.
spec → kernel 의 거리에 들어가는 결정들은 대략 일곱 가지로 분류된다.
- tiling — 큰 텐서를 어떤 블록 단위로 나누는가
- loop reorder — m, n, k 의 어느 차원이 outer loop 인가
- fusion — 어떤 op 들을 한 커널 안에 묶는가
- memory hierarchy 배치 — register / SRAM / HBM 중 어디에 어떤 텐서를 둘지
- parallelization — 어느 차원을 thread / block / grid 로 펼칠지
- dtype / precision — accumulator 의 정밀도, 입력의 dtype
- online aggregation — softmax 같은 연산의 streaming 형태로의 변환
FIG · spec → kernel 의 의사결정 트리같은 명세 → 다른 schedule
spec: O = softmax(QK^T/√d) V
├── naive 3-pass
│ ├── HBM materialize S = QK^T
│ ├── softmax(S)
│ └── attn_out = P @ V
├── tiled / fused
│ ├── tile_m × tile_n
│ ├── inner: streaming K,V load
│ └── 3 op fused → 1 kernel
└── FlashAttention
├── online softmax (m,l)
├── tile-stride accumulator
└── single-pass over K,V
왼쪽이 같은 spec 의 세 schedule. 결과는 동일, 성능은 1× / 5× / 10× 수준 차. derivation 은 이 세 가지가 “같은 spec 의 다른 도출” 임을 형식적으로 보장.
학계의 기존 시도들 — Halide 의 schedule language, Exo 의 user-scheduling primitive, TVM 의 tensor IR — 가 모두 같은 문제를 다른 각도에서 푼다. 이 강의의 방향은 “알고리즘 자체의 변형 (online aggregation 같은) 까지 derivation rule 안에 넣는다” 는 점에서 더 야심차게 추정된다.
§ 03formal model 의 정의· algebra of kernels
커널을 어떤 대수적 객체로 다룰 것인가
형식화를 시작하려면 “커널” 이 무엇인지를 수학적으로 정의해야 한다. 학계의 기존 모델은 — Halide / Exo / einsum 의 흐름에서 — 대략 다음 형태로 모인다.
spec
입력 텐서 모양과 의미적 결과를 정의하는 수학적 함수. 예: O[m,d] = Σⱼ softmax(QKᵀ/√d)[m,j] · V[j,d]
schedule
spec 을 어떻게 “실행” 할지 결정하는 변환의 순서. 메모리 / 병렬 / dtype 의 선택을 모두 포함.
execution context
schedule 이 매핑되는 실제 hardware 의 추상 — register, SRAM, HBM, threadblock, warp.
derivation rule
한 schedule 에서 다른 schedule 로 가는 의미-보존 변환. 각 rule 은 precondition 과 correctness proof 를 가진다.
soundness
모든 derivation rule 이 spec 의 의미를 보존한다는 메타 속성. 이게 깨지면 도출 결과가 틀린다.
completeness
원하는 schedule 이 정의된 rule 집합으로 도달 가능한가. 보통은 strict completeness 보다 “FlashAttention 까지 도달 가능” 같은 부분 명제로 다룸.
이 정의가 깔리면 커널은 spec 과 schedule 의 쌍 이라는 형태로 다뤄진다. 같은 spec 에 다른 schedule 이 붙으면 다른 코드가 나오지만 결과는 같다. 같은 schedule 이 spec 에 따라 다른 코드를 만든다. 이 분리가 형식화의 기초다.
Halide 의 영향
Halide(Ragan-Kelley, 2013) 가 image processing 영역에서 “algorithm vs schedule 분리” 를 처음으로 본격화했다. ML 영역에서는 TVM, Exo, Triton 의 lowering 이 같은 흐름. 이 강의의 “formalized” 가 가리키는 건 — schedule 자체를 증명 가능한 변환의 sequence 로 다루겠다는 한 단계 위의 입장으로 추정.
§ 04derivation rules· tiling · fusion · reorder
어떤 변환이 “rule” 로 들어가는가
강의가 다룰 법한 derivation rule 의 카테고리를 학계의 기존 분류로 묶어둔다. 각 rule 은 — 형식적으로 — precondition (언제 적용 가능한가) 과 의미 보존 증명을 가진다.
TILE
for i in [0,N) ⇒ for ii for it
한 loop 를 outer × inner 로 분해. precondition — N 이 T 의 배수이거나 mask/peel 처리.
REORDER
for i for j ⇒ for j for i
두 loop 의 순서 교환. precondition — loop body 의 dependency 가 reorder 에 안 영향.
FUSE
k1: A→B; k2: B→C ⇒ k_fused
두 커널을 한 커널 안으로 묶음. precondition — B 의 producer/consumer 가 같은 tile 단위로 정렬.
SPLIT
k: A→C ⇒ k1 then k2
FUSE 의 역. memory pressure 를 줄이거나 kernel launch 차원을 분리하기 위해 쓴다.
PARALLEL
for i seq ⇒ for i par
loop 차원을 thread/block 으로 펼침. precondition — iteration 간 dependency 없음.
CACHE
load(HBM) ⇒ stage(SRAM)
반복 접근되는 데이터를 더 빠른 메모리 hierarchy 로 stage. precondition — reuse 횟수 ≥ 2.
ASSOC
(a + b) + c ⇒ a + (b + c)
결합법칙 적용. floating-point 에서는 의미 보존 ≠ 비트 일치 — 별도 토론 필요.
ONLINE-AGG
batch softmax ⇒ streaming m,l
FlashAttention 의 핵심. softmax 의 정의식을 streaming 형태로 변환. 의미 보존 증명이 가장 비싼 rule 중 하나.
이 rule 집합이 완전하지 않다 — 새 rule 을 추가하는 게 연구의 본업이다. 가령 BLOCK 단위 quantization 같은 새 변환은 새 rule 로 정의되고, 그 의미 보존이 따로 증명된다.
“변환이 어렵지 않다. 어려운 건 ‘이 변환이 의미를 보존한다’ 의 증명이다 — 그게 derivation 의 본업.”학습 노트 · 재구성
§ 05예시: GEMM 의 도출· naive → tiled GEMM
가장 단순한 예 — naive matmul 에서 출발해 tiled / blocked GEMM 까지
GEMM 이 가장 잘 풀린 예다. spec 한 줄 — C[i,j] = Σ_k A[i,k]·B[k,j] — 에서 시작해서 production GEMM 의 모양까지 도출하는 과정을 한 단계씩 깐다.
FIG · GEMM derivation 시퀀스spec → tiled / shared-mem / Tensor Core
D0
spec
C[i,j] = Σ_k A[i,k]·B[k,j] · 그냥 수학적 정의
D1
TILE on i, j
i = ii·BM + it · j = jj·BN + jt — outer loops over (ii, jj)
D2
TILE on k
k = kk·BK + kt — outer loop over kk, inner accumulator
D3
CACHE A_tile, B_tile to SRAM
한 (kk) iteration 에서 BM·BK 와 BK·BN 만 SRAM 에 stage. reuse 가 BM, BN 회
D4
PARALLEL on (ii, jj)
각 outer iteration 을 threadblock 으로 매핑. 자연스러운 grid 생성
D5
CACHE accum to register
inner accumulator C_reg 를 register 에 둠. 마지막에만 HBM write
D6
MAP inner mma → Tensor Core
register-level outer product 를 mma.sync (PTX) 로 매핑. WMMA / MMA layout 의 swizzle 자동 도출
D6 까지 가면 production CUTLASS 커널의 외형이 거의 잡힌다. 각 단계는 의미 보존 증명을 갖는 rule이며, 다른 순서로 적용해도 (대부분) 같은 도착점에 도달한다 — 이게 confluence 의 형식적 의미.
강의에서 발표자가 이 예시 자체를 라이브로 깠을 가능성이 높다 — GEMM 은 모든 schedule language 의 “hello world”. 차이는 — 이 강의가 derivation 의 형식화 에 초점을 두고 있다면 — 각 단계의 의미 보존이 증명 / 검증 가능하다는 점을 강조했을 것이다.
§ 06예시: FlashAttention 의 도출· online softmax 의 자리
“발명품” 의 도출 — online softmax 가 derivation 안으로 들어온다
제목의 키워드 “FlashAttention” 이 정확히 이 자리. naive attention 의 spec 에서 FlashAttention 의 형태까지 — derivation rule 로 표현 가능한가? 가능하다는 게 이 강의의 강한 명제로 추정된다.
FIG · FlashAttention 의 derivation 시퀀스naive → flash · 핵심 rule 만
F0
spec
O[m,d] = Σⱼ softmax(QKᵀ/√d)[m,j] · V[j,d]
F1
SPLIT softmax
softmax 정의를 (max, sub, exp, sum, div) 의 분해. m = max(s); l = Σ exp(s-m); p = exp(s-m)/l
F2
TILE on j (key/value 차원)
j = jj·BN + jt. 각 jj 에서 Kⱼⱼ, Vⱼⱼ 의 부분만 본다
F3
ONLINE-AGG on (m, l, acc)
key/value 의 tile 을 streaming 으로 처리하면서 m, l, acc 를 online 으로 갱신. 의미 보존이 가장 비싼 rule
F4
FUSE 5 ops into 1 kernel
QKᵀ, mask, softmax, weighted-sum 을 한 커널 안에 묶음. 중간 score matrix 가 HBM 에 안 떨어짐
F5
CACHE Qₘ, Kⱼⱼ, Vⱼⱼ in SRAM
tile 단위 데이터를 SRAM 에 stage. reuse 패턴은 outer loop 가 m 인지 j 인지에 따라 다름
F6
PARALLEL on m
m 차원을 threadblock 으로 매핑. 각 block 이 한 query tile 을 책임
F3 의 ONLINE-AGG 가 이 derivation 의 핵심 — softmax 의 정의가 max-sub-exp-sum-div 의 한 번-패스 형태에서 streaming 형태로 변환되려면 특수한 algebraic identity(running max 의 갱신, l 의 rescale)가 필요하다. 그게 의미 보존 증명의 본업.
이 derivation 의 출력이 — 적절한 hardware 매핑(F5, F6, 그리고 D6 의 mma 매핑) 까지 거치면 — Tri Dao 의 FlashAttention 코드와 구조적으로 같은 코드가 된다. “발명을 도출로 바꾸는” 의 한 사례.
현실적 한계
물론 자동 derivation 이 무한히 잘 작동하지는 않는다 — search space 가 폭발한다. 실용적으로는 사람이 “이 rule 들을 이 순서로 적용해라” 의 schedule 을 깔고, 시스템이 의미 보존을 검증하는 형태(Exo 가 가까운 예시)가 더 가능해 보인다.
“derivation 의 가치는 자동 발견이 아니다 — 사람의 schedule 을 검증해주는 형식적 잣대다.”학습 노트 · 재구성
§ 07검증· correctness preservation
“의미 보존” 을 어떻게 보장하는가
derivation 의 형식적 가치는 검증 가능성에서 온다. 각 rule 의 의미 보존을 어떻게 증명하고, 시스템 차원에서 어떻게 자동화할 수 있는가.
학계의 기존 검증 접근은 대략 네 갈래.
- theorem proving — Coq / Lean 같은 증명 보조기 위에서 각 rule 의 정확성을 사람이 직접 증명 (e.g., CompCert).
- refinement type — 변환 전후의 semantic invariant 를 type 으로 표현. Exo 가 부분적으로 사용.
- SMT 기반 검증 — Z3 등에 변환 전후의 동등성을 풀게 함. 정수 / 비트 차원에서 작동, 실수에서는 한계.
- 경계 검증 + property test — 모든 (가능한) tile 크기와 입력 shape 에 대해 numerical 동등성을 random 검증. 실용적 baseline.
각각의 trade-off — proving 은 강력하지만 비싸고, property test 는 싸지만 약하다.
FIG · 검증 도구의 비용 vs 보장네 가지 위치
theorem proving
강한 보장 / 매우 비쌈 / 사람이 직접 증명
refinement type
중간 보장 / 중간 비용 / 자동화 가능
property test
약한 보장 / 싸다 / 실용 baseline
실제 시스템은 이 네 가지의 조합. 핵심 rule 은 proving, 일반 rule 은 refinement type, 가장자리는 property test.
floating-point 의 비결합성은 이 영역의 가장 큰 함정이다. (a+b)+c ≠ a+(b+c) bit-exact. 그래서 의미 보존을 “bit-exact” 가 아니라 “up to ULP / relative error” 로 정의하는 게 일반적. 실용적 검증은 — 어디까지의 numerical 오차를 의미 보존으로 인정할지 — 의 정책 문제로 이어진다.
§ 08채택과 한계· where this scales
이 접근이 잘 도는 자리와 안 도는 자리
형식화의 가치는 자명하지만, 실용 컴파일러 영역에서 항상 잘 도는 건 아니다. 어디서 잘 도는지, 어디서 안 도는지를 정리.
잘 도는 자리
정형화된 dense linear algebra (GEMM, GEMV, conv). spec 이 깨끗하고 hardware 매핑이 정형화돼 있어 derivation 이 닫힌 형태로 가능.
FlashAttention 류 fusion
algebraic identity 가 잘 알려진 경우 — softmax 의 streaming 형태, layer norm 의 streaming 통계 등.
새 hardware 적응
CACHE / MAP rule 의 정의를 새 메모리 hierarchy / 명령에 맞춰 추가하면 같은 spec 이 새 hardware 에 자동 lowering.
안 도는 자리 — sparse
sparse matrix, irregular access pattern. tile 단위 derivation 이 깨끗하지 않다.
안 도는 자리 — dynamic
runtime 에 결정되는 모양 (variable-length attention 일부). schedule space 가 예측 불가.
안 도는 자리 — communication
multi-GPU collectives. derivation rule 이 single-GPU 추상 위에 있는 게 일반적.
L087 NVSHMEM 의 영역.
실용적 입장
현재 산업에서 “formal kernel derivation” 이 직접 production 에 들어간 사례는 드물다. 더 흔한 패턴 — derivation 의 일부를 컴파일러 내부에 넣고, 사용자에게는 schedule primitive 로 노출(Exo, Triton 의 일부 lowering, ATL). 이 강의가 가리키는 다음 단계는 — 그 schedule primitive 자체를 더 깊은 검증 위에서 다루자는 흐름.
§ 09다른 DSL 과 비교· vs Halide · Exo · TVM
같은 문제에 대한 다른 입장
Halide (2013)
image processing 의 algorithm vs schedule 분리를 도입한 첫 번째 DSL. derivation 의 형식적 검증보다는 schedule 표현력 자체에 초점.
TVM (2018)
ML 영역의 Halide 후속. tensor IR + schedule. 자동 스케줄링(AutoTVM, Ansor) 으로 search 강조. derivation rule 의 형식적 검증은 약함.
Exo (2022)
user-scheduling 을 명시적으로 만든 시스템. schedule rule 이 “primitives” 로 노출됨. refinement type 기반 검증 시도.
Triton (2019–)
user-facing DSL 로는 schedule 의 일부(launch shape) 만 노출. lowering 안에 derivation rule 이 숨어 있음. 형식적 검증 약함.
ATL / 학계 시도
linear algebra 영역의 형식적 schedule. 학계의 가장 가까운 방향성. 자동 정리 증명과 결합된 시도.
이 강의의 위치
위의 흐름들의 한 응용으로, FlashAttention 같은 알고리즘적 변환까지 derivation 안으로 끌어오려는 시도로 추정. 확인 필요.
“DSL 의 미래는 사용자 표면이 아니라 — 그 표면이 어떤 검증 위에서 도는가로 갈린다.”학습 노트 · 재구성
§ 10기억할 메모· key takeaways
다시 열었을 때 5분 안에 손에 잡혀야 할 것
spec vs schedule 분리
커널 = spec(수학적 의미) + schedule(실행 결정). 같은 spec 의 다른 schedule 은 같은 결과를 보장.
derivation rule
한 schedule 에서 다른 schedule 로 가는 의미-보존 변환. 각 rule 은 precondition + 정당성 증명을 가진다.
8 가지 rule 카테고리
TILE · REORDER · FUSE · SPLIT · PARALLEL · CACHE · ASSOC · ONLINE-AGG. § 04 참조.
FlashAttention 의 자리
naive attention 에서 ONLINE-AGG + FUSE + TILE + CACHE 의 적용으로 도달 가능. “발명을 도출로” 의 표본.
검증 4 가지
theorem proving · refinement type · SMT · property test. 비용/보장의 trade-off.
FP 비결합성
의미 보존을 bit-exact 가 아니라 ULP/relative error 단위로 정의해야 함. floating-point 의 비결합성이 함정.
잘 도는 자리
dense linear algebra · streaming aggregation · 새 hardware 적응. dense / 정형화된 영역.
안 도는 자리
sparse · dynamic shape · multi-GPU collective. derivation 의 표현력 밖.
Slides
official repo 에 미공개
손에 새기기 — 학습 시퀀스
- GEMM 의 derivation 손으로 따라가기 — § 05 의 D0 → D6 시퀀스를 종이 위에서 직접 한 단계씩 적는다. 각 단계의 precondition 을 명시.
- FlashAttention 의 ONLINE-AGG 의미 보존 증명 — running max 의 갱신과 l 의 rescale 이 어떻게 동등한 결과를 주는지 — 종이 위에서 증명.
- Halide schedule 한 번 짜보기 — schedule language 를 직접 만져 “algorithm vs schedule 분리” 의 감을 잡는다.
- Exo repo 의 user-scheduling 예시 읽기 — derivation rule 이 사용자 표면에 노출된 가장 가까운 예.
- Triton lowering 안의 schedule — TritonGPU dialect 의 transform passes 를 한 번 펼친다. derivation rule 이 어디에 숨어 있는지 찾기.
§ 11다른 강의로 이어지는 길· connections
derivation 의 자리에 등장하는 다른 강의들
§ 12열린 질문· open questions
다음에 다시 들었을 때 직접 검증해야 할 것들
- 발표자의 정확한 신원과 그룹 — 자막이 실패해 발표자 정보가 잡히지 않았다. 영상 첫 부분의 자기 소개 직접 확인 필요.
- 이 강의가 기반한 정확한 논문 — “Formalized Kernel Derivation” 이라는 워딩이 어느 논문의 인용인지. ASPLOS / OSDI / PLDI 같은 학회의 최근 발표 추적 필요.
- 구현 가능성 — derivation 의 실제 구현이 어디까지 되어 있는가. github 에 공개된 코드가 있는가.
- FlashAttention-2/3 의 ONLINE-AGG — FA-2 의 sequence parallelism, FA-3 의 producer-consumer 분리는 derivation rule 로 어떻게 표현되는가.
- quantization 과 derivation — FP8/MXFP4 같은 dtype 변환을 rule 로 정의할 때의 의미 보존은 어떤 의미인가. L084 Numerics 와의 교차점.
- multi-GPU 확장 — single-GPU schedule 을 multi-GPU schedule 로 derivation 하는 rule 이 정의 가능한가. communication overlap 까지 형식화하는 길.
검증 메모
이 노트는 자막 실패 상태에서 도메인 지식 + 학계의 인접 연구로 재구성한 형태다. 발표자의 실제 입장 / 정확한 형식 모델 / 구현 디테일은 영상을 직접 보면서 보강 필요. § 04 의 rule 8개와 § 05/06 의 derivation 시퀀스는 학계의 일반적 표현이지, 이 강의의 정확한 표기는 아니다.