본문 바로가기

FlashAttention Kernel Project

 (21)
FlashAttention 과 Compiler Optimization 의 한계 - Operator Graph 최적화에서 Algorithm Transformation 으로 1. 문제 제기딥러닝 컴파일러는 일반적으로 Operator Graph 최적화를 수행한다.모델을 연산 그래프로 표현하고, 각 노드를 최적화하거나 fusion 을 통해 실행 효율을 높이는 방식일반적인 파이프라인Model ↓Operator Graph ↓Graph Optimization (Fusion / Layout / Scheduling) ↓Kernel Execution대표적인 최적화는 다음과 같아Kernel FusionLayout TransformationTiling / VectorizationKernel Parameter Tuning하지만 이러한 방식은 연산 그래프의 구조를 유지하는 범위에서만 동작한다.즉연산 그래프는 고정실행 방식만 최적화라는 특징을 가진다. 일정 수준의 성능 향상을 제공하지만..
현재 성공 케이스 스펙 ( core 축소, ir 기반 실행 이후 기준 ) 이 스펙이 보장하는 것Warmup 은 커널 / 버퍼 예열만 하고 학습 상태를 절대 바꾸지 않는다Capture / Replay 는 실제 학습 업데이트를 포함한다IRExecutor 실행 결과가 Replay 와 동일한 state 변화를 만든다Restore 후 N 회 replay 시퀀스가 완전 결정적 Warmup 은 no-drift 강제변하면 안 되는 것stepparamsm/vwarmup 에서 해야 하는 것leaf param.grad buffer materializelazy init 방출필요한 임시 버퍼들 선할당현재 성공 원리warmup 동안만 acif_warmup=1capture 구간은 warmup 플래그가 켜져있으면 안 됨 Backward lowering 은 leaf param.grad 에 절대 write ..
warp 의 역할을 분리했을 때, hang 발생.. ㄷㄷ producer 가 consumer 를 추월해서 ping-pong 버퍼를 덮어씀t = 0 에서 warp1 이 ready[cur] == t 를 기다리는데출력은 ready0=6, ready1=7 처럼 이미 뒤쪽 타일 번호로 덮여 있음(A) 설계상 버퍼는 2개 뿐인데warp0(loader) 는 t+1 프리페치 계속 발사해서buf0 에는 0, 2, 4, 6, ...buf1 에는 1, 3, 5, 7, ...이렇게 최신 타일 번호로 계속 생신해버림(B) 그런데 warp1(compute)가 아직 tile0 을 처리하지 못한 상태면warp1 은 ready[cur]==0 을 기대하고 기다리는데, 이미 buf0 의 ready0 가 6으로 바뀌어버려서 영원히 0이 될 수가 없음그래서 t = 0 에서 무한 대기 발생즉, 문제..
NCU 리포트 해석 flashattn_streaming_16x16_kernel_mw_v6 실험 조건Grid/Block : 1024, 1, 1 x 64, 1, 1 = block 당 2 warps핵심 구조warp 0 : cp.async global - shared + waitwarp 1 : WMMA + streaming softmax + PV타일 경계마다 syncthreads 로 안정성 보장 1. GPU Speed Of Light Throughput관측값Memory Throughpupt : 89.10%Compute SM Throughput 89.10%L1/TEX Cache Throughput 90.22%DRAM Throughput 3.28%Duration 1.12 ms해석DRAM 이 3.28% 인데도 Memory Throughput 이 89% 로 뜨는 건 전형적으로실제 병목이 DRAM 대역폭이..

티스토리툴바