Warp-Specializaed Pipeline & cp.async Multi-Stage Overlap 개념

1. 기본 전제 : warp specialization 이란?

warp specialization 은 CTA 내부의 warp 들이 서로 다른 역할을 맡는 구조

• warp 0 글로벌 메모리 - shared memroy 로드
• warp 1 shared memroy - 레지스터 - 연산 

목표 : 메모리 지연 latency 를 연산 시간 뒤에 숨겨서 SM 이 놀지 않게 만드는 것

여기서 함정이 발생

 

2. 문제의 본질 : warp 간 속도는 절대 같지 않다.

두 warp 는 다른 종류의 작업을 한다.

그래서 타일 단위로 항상 특정 warp 에서의 속도가 다름 - 이것 자체가 문제가 아닌, 동기화 위치가 너무 촘촘할 경우 발생

 

3. 1-stage 파이프라인의 구조적 한계

[warp0] cp.async(t)
[warp0] wait
[warp0] publish ready(t)
---------------------- (동기화)
[warp1] compute(t)
---------------------- (다음 타일)

이 구조의 문제

• 타일 매번
  • wait
  • ready 확인
  • 사실상 CTA barrier 수준의 정렬 발생
• warp 간 속도 차이가 즉시 stall 로 변호나됨

즉, latency hiding 을 하려고 warp specialization 을 썼는데 동기화가 잦아서 오히려 latency 를 노출시킴

 

4. cp.async 2-stage 로 쌓는다 의 진짜 의미

핵심 개념 : 로드와 연산을 1:1 로 맞추지 말고, in-flight 상태로 여러 타일을 겹쳐 준다

이것이 흔히 말하는

• double buffering
• multi-stage pipeline
• cp.async 그룹 누적

개념적 변화

시간 →
warp0: cp.async(t) ─ cp.async(t+1) ─ cp.async(t+2) ─ …
warp1:           compute(t) ─ compute(t+1) ─ …

중요한 변화는 두 가지

• cp.async 를 먼저 여러 개 발사
• wait 를 바로 하지 않고, 소비 직전까지 미룸

 

5. wait 를 늦춘다는 말의 정확한 의미

버퍼를 재사용하거나 읽기 직전에만 wait 한다

즉,

• cp.async는 발사 순서가 중요
• wait 은 의존성 경계에서만 필요

올바른 패턴은

cp.async(t+1)
cp.async(t+2)
commit_group
...
(wait)  // t 버퍼를 실제로 읽기 직전
compute(t)

• warp 0 는 다음 주문서를 계속 쌓고
• warp 1 은 이미 도착한 재료로 일함

 

6. 이전에서 계속 발생한 문제

• 배치가 커지면 일부 batch 에서 무한 대기
• ready 값은 증가, 소비가 안 됨

 

7. 핵심 해결 방법 : handshake

ready 단방향 신호에서 ready / consumed 양방향 핸드셰이크