warp 에 이은 lane specialization??

warp 는 독립적으로 스케줄링되는데, lane 은 같은 warp 에서 락스텝이라 

서로 다른 일을 시켜도 결국 그 warp 가 그 일들을 순차로 수행되기 쉬움 ( divergence / masking cost )

 

warp specialization 효과가 큰 이유

• warp 단위로 스케줄러가 따로 놀림 : warp 0 가 메모리 대기중일 때 warp 1 이 연산으로 SM 을 채울 수 있음
• 즉, memory warp 와 compte warp 를 분리하면 진짜 latency hiding 이 생김

현재 코드 구현도

• warp 0 : cp.async + wait
• warp 1 : wmma + softmax + PV

 

lane specialization 존재는 하지만, 목적이 다름

lane 단위로 역할을 나누는 건 주로 work partitioning 혹은 메모리 coalescing / transaction 정렬ㅇ르 위해 한다.

• 로드 분업 : lane 들이 각자 다른 chunk 를 옮겨서 한 타일을 완성
  • lane -> (k, seg) 매핑이 정확히 lane specialization 
• reduction 역할 분리 : lane 0 만 shared 에 write / lane 0 만 상태 업데이트
• 벡터 / 차원 분업 : lane 이 Dv 차원이나 head dim 을 나눠서 accumulate / store 

latency hiding 을 늘리는 것이 아니라 한 warp 안에서 일을 나눠서 한 번에 끝내기 / 메모리 효율 맞추기가 목적

 

그래도 의미 있게 쓰이는 lane-level specialization 

• producer / consumer 를 warp 내부에서 나누는 게 아니라 warp 간에 나눔
• lane 0 / 몇 lanes 만 특정 작업 : 포인터 계산, 상태 갱신, 큐 인덱스 관리, atomic 최소화
  • 스케줄 / 동기 줄여서 이득
• 메모리 transaction 맞춤: 16/32/128 B 정렬을 lane 에 매핑
  • coalescing 최대화
• Tensor Core operand layout : idmatrix / wmma feed 를 위해 lane 별 주소 패턴을 강제

 

결론

• warp specialization = latency hiding 을 위한 주력 축
• lane specialization = 분업/메모리 정렬/리덕션 wirter 같은 미세 역할