TMA 사용 시점 및 이유: NVIDIA Hopper GPU에서의 텐서 메모리 가속에 대한 워크로드 기반 특성화

NVIDIA Hopper에서 도입된 Tensor Memory Accelerator (TMA)는 프로덕션 GPU 커널(kernel) 전반에 걸쳐 광범위하게 사용되지만, 그 활용은 눈에 띄게 선택적입니다. CUTLASS, FlashAttention-3, 그리고 FlashInfer의 전문가급 커널 개발자들은 커널 내 일부 텐서 연산자에 TMA를 적용하는 반면, 나머지 부분에는 cp.async, 직접적인 글로벌 로드(direct global loads), 또는 분산 공유 메모리(distributed shared memory, DSM)에 의존합니다. 공급업체 문서는 API를 설명하고 있으며, 최근의 마이크로벤치마킹 작업은 TMA 처리량(throughput)을 독립적으로 특성화했지만, 이들 어느 것도 실무자들이 실제로 직면하는 질문, 즉 커널 내 메모리 연산자에 대해 TMA를 사용해야 하는지, 그리고 그렇다면 클러스터 멀티캐스트(cluster multicast)와 함께 사용할지 여부를 답해주지는 못합니다. 우리는 H100을 대상으로 워크로드 기반 특성화 연구를 통해 이러한 격차를 해소하고자 합니다. 우리는 FlashAttention-3 포워드 과정에서 사용된 TMA, 멀티캐스트, DSM, cp.async 선택 사항들을 목록화하고, 목표 마이크로벤치마킹(전송 크기 교차점, 멀티캐스트 스케일링)을 통해 판별적인 특징들(discriminating features)을 분리하며, Hopper에서 로드 메커니즘 선택에 대한 경험적 규칙(rules of thumb)을 제시합니다. 나아가 우리는 경쟁하는 클러스터 브로드캐스트 프리미티브로서 TMA 멀티캐스트와 DSM을 비교하고, 페이지 크기, 데이터 타입, 모델 형태, 실행 단계 전반에 걸쳐 H100에서 FlashAttention-3의 paged-key-value (KV) TMA 대 cp.async 디스패치를 분리합니다. 우리의 결과는 Hopper에서 언제 TMA가 효과를 발휘하는지, 그리고 멀티캐스트가 도움이 되는지 혹은 방해가 되는지에 대한 구체적이고 예측 가능한 지침을 제공하며, 이 분석을 B200의 CUTLASS GEMM 연구를 통해 서버 Blackwell (SM100)로 확장합니다. 이 연구는 정성적인 디스패치 규칙이 세대를 거쳐 유지되는 반면, 절대적인 교차점(absolute crossovers)은 텐서 코어 및 대역폭 스케일링에 따라 이동한다는 것을 보여줍니다.