제4부 데이터센터와 스케일

4.3 Hopper II — TMA·스레드블록 클러스터·DSMEM

4.3.1 이 장에서 배우는 것 / 선수 지식

4.2의 FP8이 정밀도 축의 진화였다면, 이 장은 Hopper의 실행·메모리 모델 진화다. 세 가지 신기능 — TMA(Tensor Memory Accelerator, 대량 비동기 복사), 스레드블록 클러스터(thread block cluster, 블록들을 묶음), DSMEM(distributed shared memory, 블록 간 공유메모리) — 이 3.4의 GEMM 파이프라인과 1.7의 온칩 메모리를 한 단계 끌어올린다. 대부분의 사용자는 이것을 직접 작성하지 않고 라이브러리(CUTLASS 3.x·cuDNN·TE)를 통해 누리지만, 왜 최신 커널이 빠른지·무엇이 바뀌었는지를 알아야 한다.

배우는 것:

• TMA 가 3.4의 cp.async를 어떻게 대체·발전시키는지 — 디스크립터 기반 다차원 대량 복사, 주소 계산을 하드웨어로.
• 스레드블록 클러스터 — grid/block 사이에 낀 새 계층. 여러 블록을 한 GPC에 공동 스케줄해 협력하게 한다.
• DSMEM — 클러스터 내 블록들이 서로의 공유메모리에 직접 접근, 공유메모리가 사실상 클러스터 전역 풀이 됨.
• warp specialization·WGMMA 와 함께 이들이 어떻게 최신 GEMM·어텐션 커널(CUTLASS 3.x·FA3)을 떠받치는지.

선수 지식: 3.4(CUTLASS 파이프라인·cp.async·multistage — 이 장이 그 진화), 1.7~1.8(공유메모리·온칩 — DSMEM은 그 확장), 1.3(GPC·SM 계층 — 클러스터는 GPC 단위), 1.5(워프 스케줄링 — warp specialization), 1.11~1.12(텐서코어 — WGMMA). 환경은 0.3. 이들은 Hopper+ 전용 기능이며 한도·세부는 세대·시점 의존(확인 필요). 코드는 개념 골격이며 직접 작성은 드물다(라이브러리가 사용).

4.3.2 동기 — 왜 블록 하나로 부족해졌나

3.4에서 GEMM 성능은 타일을 키워 재사용↑·HBM 트래픽↓ + 파이프라인으로 로드 지연 은닉에서 나왔다. 텐서코어가 세대마다 빨라지자(1.12) 두 가지가 한계에 부딪혔다:

Hopper의 답:

• 데이터 공급을 스레드에서 떼어 전용 하드웨어(TMA) 가 대량·다차원으로 비동기 운반.
• 블록 하나의 벽을 넘어 여러 블록을 클러스터로 묶고(DSMEM) 공유메모리를 합쳐 더 큰 협력 타일을 만든다.
• 워프들을 역할 분담(데이터 운반 전담 vs 연산 전담)해 파이프라인을 깊고 단순하게.

이 셋이 합쳐져 "텐서코어를 기아 상태로 두지 않고 더 큰 타일을 협력적으로" 처리한다 — 3.6의 FlashAttention이 FA3에서 더 빨라진 것도 이 기반이다.

4.3.3 TMA — 대량 비동기 복사 엔진

TMA(Tensor Memory Accelerator) 는 글로벌↔공유메모리 간 다차원 타일을 한 명령으로 대량 복사하는 전용 하드웨어 유닛이다.

3.4의 cp.async와 비교:

텐서 디스크립터. TMA는 "이 글로벌 텐서는 $M\times N$, 이 stride, 이 타일은 (i,j)" 같은 디스크립터를 받아 다차원 타일을 자동 운반한다. 경계 밖(out-of-bounds)은 0 채움 등 하드웨어가 처리 — 2.2의 손 경계검사가 사라진다.

비동기 완료 — mbarrier. TMA 복사는 비동기다. 완료를 mbarrier(메모리 배리어 객체)로 기다린다 — "이 타일 도착하면 신호". 소비 워프는 그 신호를 기다렸다 연산. 이것이 다음의 warp specialization과 맞물린다.

요지: TMA는 데이터 이동을 스레드의 일에서 하드웨어의 일로 옮겼다. 스레드(특히 텐서코어를 실행할 워프)는 주소 산수에서 해방돼 연산에 집중한다.

4.3.4 스레드블록 클러스터 — 새 계층

기존 CUDA 실행 계층은 grid → block → warp → thread(2.3)였다. Hopper는 grid와 block 사이에 클러스터(cluster) 를 넣는다:

핵심 보장: 한 클러스터의 블록들은 같은 GPC(1.3)에 공동 스케줄된다. 보통 블록들은 어느 SM에 언제 갈지 모르고 서로 독립이라 협력 불가였다(2.8의 grid 동기화 제약). 클러스터는 블록들이 동시에·인접 SM에 살아있음을 보장해, 블록 간 직접 협력(동기화·데이터 공유)을 가능케 한다.

클러스터 동기화는 cluster.sync() — 클러스터 내 全블록 배리어로 블록 경계를 넘는 협력을 제공한다.

왜 필요: 한 블록의 공유메모리·워프 수로는 부족한 더 큰 협력 단위가 필요해졌다. 클러스터는 "여러 SM을 하나의 큰 협력 블록처럼" 쓰게 한다.

4.3.5 DSMEM — 블록 간 공유메모리

클러스터가 블록을 묶으면, DSMEM(Distributed Shared Memory) 으로 한 블록이 같은 클러스터 내 다른 블록의 공유메모리에 직접 접근할 수 있다(SM-to-SM 네트워크 경유).

효과:

• 더 큰 유효 타일. 4블록 클러스터면 공유메모리 풀이 ~4배 — 3.4의 타일을 블록 하나 한도 너머로 키울 수 있다(데이터 재사용↑).
• 블록 간 데이터 교환을 HBM 없이. 한 블록이 로드한 데이터를 옆 블록이 DSMEM으로 바로 씀 → 중복 로드·HBM 트래픽 제거(3.6 융합 정신의 클러스터판).
• 협력 패턴. reduce·broadcast를 클러스터 범위에서 공유메모리 속도로.

요지: 공유메모리(1.7)의 "블록 사적 스크래치패드" 정의가 클러스터 안에서 "여러 블록이 나눠 쓰는 공동 풀" 로 확장됐다.

4.3.6 warp specialization·WGMMA — 파이프라인 재편

TMA·클러스터·DSMEM을 실제 커널에서 엮는 방식이 warp specialization(워프 역할 분담) 이다. 3.4의 파이프라인은 모든 워프가 로드·연산을 같이 했다. Hopper 스타일은 워프 그룹을 역할별로 나눈다:

둘은 mbarrier로 핸드셰이크한다 — producer가 '도착' 신호, consumer가 '소비 완료' 신호.

파이프라인이 역할로 분리돼 깊고 단순하다 — 텐서코어가 거의 안 쉰다.

• WGMMA(warpgroup MMA): Hopper 텐서코어 명령. 워프 하나가 아니라 워프그룹(4워프) 단위로 더 큰 MMA를 비동기로 발행 — 더 큰 타일을 한 번에, 연산과 다음 로드를 겹침(1.12 텐서코어 세대 진화).

이 구조가 CUTLASS 3.x(CuTe, 3.5)·FlashAttention 3·cuDNN의 최신 커널이 Hopper에서 빠른 이유다. 사용자는 이것을 직접 안 짠다 — 라이브러리가 TMA·클러스터·warp-spec·WGMMA를 조립해 제공하고, 우리는 그 위에서 GEMM·어텐션을 부른다.

4.3.7 워크스루 — 클러스터 런치 골격(개념)

직접 쓸 일은 드물지만, 무엇이 노출되는지 본다(개념 골격, 정확한 API는 CUDA 버전·CUTLASS 예제로):

// 클러스터 차원을 런치에 명시 (블록 2×1×1 = 클러스터당 2블록)
cudaLaunchConfig_t config = {};
config.gridDim = grid; config.blockDim = block;
cudaLaunchAttribute attr[1];
attr[0].id = cudaLaunchAttributeClusterDimension;
attr[0].val.clusterDim = {2, 1, 1};       // 클러스터 형상
config.attrs = attr; config.numAttrs = 1;
cudaLaunchKernelEx(&config, my_kernel, args...);

// 커널 안 (cooperative_groups)
#include <cooperative_groups.h>
namespace cg = cooperative_groups;
__global__ void my_kernel(...) {
    cg::cluster_group cluster = cg::this_cluster();
    unsigned rank = cluster.block_rank();          // 클러스터 내 내 블록 번호
    // DSMEM: 다른 블록의 공유메모리 포인터를 map
    extern __shared__ float smem[];
    float* other = cluster.map_shared_rank(smem, (rank ^ 1)); // 짝 블록의 smem
    // ... TMA로 타일 적재, other로 옆 블록과 공유, cluster.sync()로 동기화 ...
    cluster.sync();                                 // 클러스터 전역 배리어
}

읽는 법:

• clusterDim — 런치 시 클러스터 형상 지정. 클러스터 크기 한도는 세대·점유율 의존(확인 필요).
• block_rank()·map_shared_rank() — DSMEM의 핵심: 같은 클러스터 다른 블록의 공유메모리를 가리킨다.
• cluster.sync() — 블록 경계를 넘는 배리어(2.8의 __syncthreads가 블록 안이었다면, 이것은 클러스터 안).
• 실무에선 이것을 직접 작성하지 않고 CUTLASS·cuDNN이 내부에서 사용. 우리는 "왜 빠른가"를 이해하는 선에서.

4.3.8 중간 점검 (체크포인트)

• [ ] TMA가 cp.async 대비 무엇을 하드웨어로 옮겼는지(다차원 주소·대량·경계) 설명한다.
• [ ] 스레드블록 클러스터가 grid/block 사이 새 계층이며 같은 GPC 공동 스케줄로 블록 간 협력을 가능케 함을 안다.
• [ ] DSMEM이 공유메모리를 "블록 사적"에서 "클러스터 공동 풀"로 확장해 HBM 없이 블록 간 공유함을 말한다.
• [ ] warp specialization(producer/consumer)·WGMMA가 이 셋을 엮어 텐서코어를 기아 상태로 두지 않음을, 그리고 사용자는 라이브러리로 누림을 안다.

증상→의심: Ampere용 커널을 Hopper에서 실행했는데 기대만큼 안 빠름 → TMA·클러스터·WGMMA 미사용(Hopper 최적 커널·CUTLASS 3.x·최신 라이브러리 필요). 클러스터 런치 실패 → 클러스터 크기·점유율 한도 초과 또는 비-Hopper. DSMEM 접근 오류 → 클러스터 밖 블록 접근(같은 클러스터 내에서만). 직접 작성한 Hopper 커널이 라이브러리보다 느림 → warp-spec·파이프라인 미성숙(라이브러리 사용 권장).

4.3.9 트레이드오프 — 직접 쓸 것인가, 라이브러리로 누릴 것인가

견해: 이 세 기능은 거의 항상 라이브러리를 통해 쓴다. 직접 TMA·클러스터·warp-spec 커널을 작성하는 것은 CUTLASS 기여자·극소수 연구 영역이고, 직접 라이브러리를 이기기 매우 어렵다. 실무 행동은 단순하다: Hopper에서 최신 CUTLASS·cuDNN·Transformer Engine·FlashAttention을 쓰면 이 기능들이 자동 적용된다. 우리 몫은 "왜 최신 커널·세대가 빠른지"를 이해해 올바른 라이브러리·세대를 고르고 폴백을 검출하는 것.

비NVIDIA 정직: "데이터 이동을 전용 엔진으로, 협력 단위를 키운다"는 발상은 보편적이다 — 다른 가속기도 DMA 엔진·온칩 협력 구조를 둔다. TMA·클러스터·DSMEM은 그 NVIDIA Hopper 구현이며, 이식성은 없다(Hopper+ 전용). 따라서 이 기능에 의존하는 손커널은 세대 락인 — 라이브러리가 세대 차이를 흡수해 주는 이유.

4.3.10 흔한 함정

• 구세대 커널로 Hopper 성능 기대. Ampere 최적 커널은 TMA·클러스터·WGMMA를 안 쓴다. Hopper 칩만 새것이라고 자동으로 안 빨라진다 — Hopper 빌드 타깃·최신 라이브러리·CUTLASS 3.x가 필요(2.12 아키텍처 타깃).
• 클러스터를 맹목적으로 키움. 클러스터가 클수록 협력 풀은 크지만 공동 스케줄 제약·점유율 제약이 커진다. 크기 한도와 점유율을 함께 고려(세대 의존).
• DSMEM을 클러스터 밖에서 기대. 다른 블록 공유메모리 접근은 같은 클러스터 내에서만. 클러스터 경계를 넘으면 불가 — 협력 범위를 혼동하지 말 것.
• 이식성 무시. TMA·클러스터·DSMEM에 의존한 손커널은 Hopper+ 전용이라 구세대·타 벤더에서 깨진다. 이식이 필요하면 라이브러리에 맡겨 세대 분기를 흡수.
• 직접 구현 과신. warp specialization·mbarrier 핸드셰이크·WGMMA 파이프라인은 미묘하고 버그가 잦다. 라이브러리가 수개월 튜닝한 것을 직접 작성해 따라잡기 어렵다 — 연구 목적이 아니면 라이브러리.
• mbarrier·비동기 완료 오해. TMA·WGMMA는 비동기라 완료 신호(mbarrier)를 기다려야 데이터가 유효. 동기 코드처럼 다루면 경쟁·오류. (직접 작성할 때.)

4.3.11 연습문제

연습 1 (TMA vs cp.async). TMA가 cp.async(3.4) 대비 무엇을 개선했는지 세 가지를 들고, 그로 인해 연산 워프가 얻는 이득을 설명하라.

연습 2 (클러스터의 보장). "보통 블록들은 협력 못 한다"의 이유와, 스레드블록 클러스터가 무엇을 보장해 협력을 가능케 하는지 설명하라.

연습 3 (DSMEM 타일 산수). 블록당 공유메모리 228KB(개략, 세대 의존) 한도에서 GEMM 타일이 더 커야 하는데 한 블록으론 부족하다. 4블록 클러스터 + DSMEM이 어떻게 해결하는지, 유효 공유메모리 풀과 이득을 설명하라(정확한 한도는 확인 필요 전제).

연습 4 (실무 행동). 팀이 H100으로 옮겼는데 학습이 A100 대비 기대만큼 안 빨라졌다. TMA·클러스터·WGMMA 관점에서 점검할 것과, 직접 커널을 작성하는 대신 권장되는 조치.

4.3.12 요약 · 다음 장

• Hopper는 실행·메모리 모델을 세 축으로 진화시켰다: TMA(다차원 대량 비동기 복사를 하드웨어로 — cp.async 발전), 스레드블록 클러스터(grid/block 사이 새 계층, 같은 GPC 공동 스케줄로 블록 협력), DSMEM(클러스터 내 블록 간 공유메모리 직접 접근).
• 이들은 warp specialization(producer/consumer)·WGMMA(워프그룹 MMA)와 엮여 텐서코어를 기아 상태로 두지 않는 깊은 파이프라인을 만든다 — CUTLASS 3.x·FA3·cuDNN 최신 커널의 기반.
• 사용자는 직접 안 짠다 — 라이브러리가 제공하고, 우리 몫은 "왜 최신 세대·커널이 빠른지" 이해해 올바른 세대·라이브러리·빌드 타깃을 고르고 폴백을 검출하는 것. 이식성 없는 Hopper+ 전용.

다음 장: 4.4 Blackwell I — 듀얼다이·FP4·2세대 트랜스포머 엔진. Hopper 다음 세대다. 두 다이를 하나처럼 잇는 패키징, FP4로 더 내려간 정밀도, 2세대 트랜스포머 엔진을 다룬다 — 발표 기준·변동 가능으로 신중히 표기하며, 무엇이 검증된 방향이고 무엇이 로드맵인지 가른다.