1. 흔한 오해
- warp 는 병렬 실행 단위다
- warp 수를 늘리면 병렬성이 늘어난다
이건 하드웨어 실행 모델 설명일 뿐, 커널 설계 관점에서는 반쪽짜리 이해
2. 커널 설계 관점의 진짜 의미
warp 는 동시에 실행되는 계산 단위면서 동시에 서로 다른 책임을 맡기는 최소 스케줄링 단위다.
그래서 warp specialization 이 성립한다
3. FlashAttention 에서의 구체적 의미
CTA
├─ warp 0 → 역할 A
│ - QK score 계산
│ - online softmax stats (m, l)
│ - exp(score) 생성
│
└─ warp 1 → 역할 B
- O 누적 (P * V)
- scale_old / scale_tile 적용
- 최종 normalize + store
2-warp CTA 예시
- 두 warp 는 같은 데이터를 동시에 계산하지 않는다.
- warp 0 결과 없이는 warp1 이 할 일이 없다.
즉,
- 병렬성은 lane 내부 SIMD
- 구조는 warp 간 파이프라인
4. warp 를 역할 단위로 봐야 하는 이유
- control flow 를 깨끗하게 분리 가능
- 레지스터와 smem 생존 범위르 ㄹ줄인다
- 동기화 비용을 의식적으로 통제할 수 있음
5. 병렬은 어디서?
- lane-level SIMD
- CTA-level parallelism
- warp-level specialization
- warp 수를 늘린다고 병ㄹ려 계산량이 늘어나느 것이 아닌
- warp 수는 역할 분해의 granularity
'GPU-KERNEL' 카테고리의 다른 글
| CUDA Kernel Analysis System - Idea Exploration Draft (0) | 2026.03.05 |
|---|---|
| AICF Kernel Engineering Report ( GEMM & BiasAdd ) (0) | 2026.02.16 |
| SMEM 에 대한 접근 - 저장소가 아닌 연산 스케줄의 일부, layout 이 알고리즘 그 자체 (0) | 2025.12.16 |
| Warp-Specializaed Pipeline & cp.async Multi-Stage Overlap 개념 (0) | 2025.12.15 |
| GPU memory transaction, Byte 단위 사고 정리 (0) | 2025.12.15 |