GPU의 메모리 모델과 인터커넥트(Network) 설계

사실 많은 사람들이 Scale‑UP과 Scale‑Out 버스를 이야기할 때, 거의 항상 네트워크만 이야기하고 GPU Memory Model 관점을 빼먹는다. 한편으로는, 황 회장이 말한 “먼저 Scale‑UP, 그다음에 Scale‑Out”이라는 것도, 솔직히 말하면 일종의 세일즈 화법이라고 이해해도 된다. 되묻자. NV가 Scale‑Out 쪽에서 올해 GTC에서 내놓을 만한 게 뭐가 있나? IB 스위치는 조용하고, 이더넷 스위치와 NIC는 RoCE에서 여전히 문제가 산더미다… 이게 본질적인 문제다. 이런 것들을 제쳐 두더라도, 결국 본질은 여전히 메모리에 있다.

그래서 오늘은 메모 관련해서 Nvidia GPU를 처음부터 끝까지 한번 제대로 정리해 보려고 한다. 마침 GTC25에 《CUDA Techniques to Maximize Memory Bandwidth and Hide Latency》[1], GTC24에 《Advanced Performance Optimization in CUDA》[2] 세션이 있어 같이 참고할 수 있다.

여기에 Blackwell의 Tensor Memory 도입으로 바뀌는 메모리 모델, Tile 기반 IR 등 여러 요소를 합쳐서 정리해 보고, 마침 화웨이 UB와 UALink 1.0 표준도 공개됐으니 마지막에는 Scale‑Out과 Scale‑UP에서 요구되는 메모리 모델을 이야기해 보겠다. 덤으로 eRDMA가 다중 경로(multipath)에서 메모리 모델을 어떻게 구현하는지, 그리고 표준 RC 및 AWS SRD와 어떻게 다른지도 비교할 것이다.

또한 GTC25 세션에는 "저지연 Cluster 동기화"와 "메모리 대역폭 최대화"라는 두 가지 중요한 주제도 있는데, 이것도 함께 살펴볼 예정이다. 글의 구조는 다음과 같다.

1. 메모리 일관성 모델
   1.1 기초 예제에서 출발
   1.2 순차 일관성(Sequential Consistency)
   1.3 Total Store Order
   1.4 Relaxed Consistency
   1.5 Cache 일관성과 Memory Model의 차이
2. Nvidia GPU Memory Model
   2.1 Single Thread
   2.2 Multi‑Thread
   2.2.1 순차 일관성(SC)
   2.2.2 Acquire
   2.2.3 Release
   2.2.4 Relaxed
   2.2.5 Scope
   2.3 Async Thread
   2.4 Async Proxy
3. memory order를 이용한 저지연 Cluster 동기화
   3.1 Thread Block Cluster 프로그래밍
   3.2 Cluster 저지연 동기화
   3.3 DSMEM 기반 멀티캐스트
4. 메모리 대역폭 최대화
   4.1 메모리 계층 구조 복습
   4.2 Little's Law
   4.3 병렬 최적화와 비동기 접근
5. LD/ST 명령으로 Cache 제어
   5.1 LD 명령
   5.2 ST 명령
6. Scale‑UP과 Scale‑Out 네트워크 설계 논의
   6.1 메모리 접근 Size 관점
   6.2 메모리 접근 지연
   6.3 Memory Model

---

1. 메모리 일관성 모델

일관성(consistency)의 출발점은 폰 노이만 구조에서의 다음과 같은 가정이다.

“어떤 read 연산도 최근에 write된 결과를 반환해야 한다.”

하지만 분산 시스템이나 멀티코어 CPU 시스템에서는, 연산 지연 등의 요인 때문에 결과가 예측 불가능해진다.

먼저 코어가 둘 있는 프로세서를 보자. 하나는 Producer, 하나는 Consumer라고 하자.

r2가 Core1이 만든 새로운 데이터를 제대로 읽을 수 있을까? 실제 실행 시 Core1의 S1과 S2가 ReOrder(재정렬)될 수 있고, 이 경우 Core2 쪽의 L1이 L2보다 먼저 실행되어 r2가 옛날 데이터를 읽게 된다.

Reorder가 발생하는 경우는 몇 가지 유형이 있다.

가장 직관적인 메모리 일관성 모델은 순차 일관성(Sequential Consistency, SC)이다. 가장 이른 형식적 정의는 Lamport의 논문 《How to Make a Multiprocessor Computer that Correctly Executes Multiprocess Programs. IEEE Transactions on Computers, C-28(9):690–91, Sept. 1979》에 나온다. 멀티코어 프로세서에서는, 멀티코어 실행 시에도 프로그램 실행 순서(Program Order)가 단일 코어에서의 순서와 동일해야 한다는 의미다.

문제의 본질로 돌아가면, Program Order와 Memory Order가 서로 다른 Load/Store 조합에서 어떤 제약을 받는가 하는 문제다. 형식적 정의는 아래와 같다.

SC를 유지하는 실질적인 방법은 두 가지다.

1. 단일 코어만 사용해 실행한다.
2. 메모리 접근 시 순차적 접근을 강제한다.

즉, 매 시간 스텝마다 switch가 실행할 스레드를 하나 선택하고, 그 스레드의 “다음 이벤트” 하나를 완전히 수행하는 모델이다. 이 모델은 순차 일관성 규칙을 충실히 따른다. 하지만 치명적인 단점이 있다. 너무도 느리다는 점이다. 매번 한 개의 명령어만 실행할 수 있으니, 멀티스레드 병렬 실행의 장점을 대부분 잃게 된다.

더 나쁘게는, 각 명령어가 끝나기 전까지 다음 명령어를 실행할 수 없다. 즉, 현재 명령어의 효과가 다른 모든 스레드에 관측 가능해지기 전에는 더 이상 진행할 수 없다.

프로세서 관점에서 메모리 write를 전부 기다리면 Store가 너무 느려진다. 그래서 보통 Store Buffer를 두어 지연을 숨기고 stall을 피한다. 멀티코어 프로세서에서는 각 코어가 독립적인 Store Buffer를 가진다.

이런 구조에서는, 그림과 같이 두 코어 모두 옛날 값을 읽을 수 있다.

하지만 이런 trade‑off로 얻는 성능 이득은 매우 크다. 그래서 등장한 것이 Total Store Order(TSO)다. 형식적으로는 Store→Load 순서 보장을 포기하고, Store Buffer를 도입하는 모델이다.

Store→Load의 문제는 FENCE로 해결할 수 있다. Fence 구현 자체도 간단하다. 예를 들어 store buffer를 비우는 식으로 주 메모리에 대해 Read‑Write coherence를 맞추면 된다.

더 나아가, 더 많은 Reorder를 허용하면 프로그램의 실행 병렬성을 높일 수 있을까? 그리고 Fence(Memory Barrier)를 적절히 사용해 프로그램 실행 순서의 “논리적” 정합성만 보장할 수 있을까?

산업계에서 Relaxed Consistency를 분류할 때는, TSO(앞서 본 대로 Store→Load 제약을 포기), Partial Store Order(PSO, 또 다른 제약 제거), 그리고 네 가지 제약을 전부 포기하는 Relaxed Memory Order(RMO) 등으로 구분하기도 한다. 실제로 많은 프로세서가 이런 relax를 지원하며, GPU도 근본적으로는 Relaxed Consistency 시스템이다.

이 지점이 많은 사람이 헷갈려 하는 부분이다. Cache 일관성은 Store를 어떻게 "필요할 때" 다른 프로세서에게 전파해, 쓰기가 다른 프로세서에서도 관측 가능하게 할지를 정의하는 메커니즘이다. 반면 Memory Model은 연산이 다른 프로세서에 전파될 때의 "순서" 경계를 정의한다.

Nvidia GPU의 메모리 계층 구조는 아래와 같다. 수천 개의 CUDA 코어에 대해 TSO(Total Store Order)를 유지하려면 비용이 엄청나게 크다. 그래서 Nvidia GPU는 Partial Store Order 메모리 모델을 채택했다.

아키텍처마다 미묘한 차이가 있는데, 예를 들면 인텔이 스스로 요약한 네 가지 메모리 모델 버전 같은 것들이 있다.

단일 스레드 관점에서 보면, 동일 주소에 대한 LD/ST는 순서를 유지한다.

다만 예외가 하나 있다. 아래 프로그램의 출력 결과를 추측해 보자. 실제로는 정의되지 않은 동작(undefined behavior)이다.

cpp
#include<iostream>
#include<cuda.h>

__constant__ int val = 1;

__global__ void kernel_constant_sc() {
  int tid = threadIdx.x + blockDim.x * threadIdx.y;
  if (tid != 0)  {
    printf("Thread %d, val %d\n", tid, val); // val을 Const$로 load
  } else {
    // constant 제거
    int* mut_val = const_cast<int*>(&val);
    asm volatile("" : "+l"(mut_val)); // store new value
    *mut_val = 42;
  }
}

int main(int argc, const char* argv[]) {
  int n = 2;
  if (argc == 2) {
    n = strtol(argv[1], NULL, 10);
  }
  kernel_constant_sc<<<1, n>>>();
  cudaDeviceSynchronize();
  return 0;
}

이는 SM 내에 Read‑Only Cache가 따로 있기 때문이다. 상수는 이 공간에 저장되며, 이 캐시는 L2 Cache와 독립된 데이터 경로를 사용한다. 그래서 문제가 발생한다.

이런 상수 값을 수정하면 정의되지 않은 동작이 발생할 수 있다.

---

2. Nvidia GPU Memory Model

앞서 봤듯이, 순차 일관성(SC)을 유지하려면 다음 네 가지 규칙을 충족해야 한다.

(그림 생략, 일반적인 Load‑Load, Load‑Store, Store‑Store, Store‑Load 순서 보장 규칙을 의미)

Total Store Order(TSO)는 Store Buffer를 도입하기 위해 네 번째 규칙(Store→Load)을 포기한다. 그리고 단일 코어 내부에서는 Load가 Write Buffer를 bypass해서 읽을 수 있게 하고, 코어 간에는 Fence를 사용해 순서를 맞춘다.

그러나 GPU에는 수천 개의 CUDA 코어가 있다. 이 모든 코어의 메모리 연산을 TSO로 유지하려면, 명령어 수준 병렬성과 데이터 수준 병렬성에 막대한 성능 손실을 야기한다. 즉, GPU에서 TSO를 유지하는 비용은 매우 크다. 따라서 Nvidia는 GPU 내에서 Relaxed Order를 지원하고, ATOMIC과 FENCE를 활용하는 쪽을 택했다. GPU에서 Nvidia는 크게 4가지 메모리 순서를 지원한다.

2.2.1 Sequential Consistency (SC)

아래 그림처럼, SC에서는 LD/ST가 특정 연산 전·후 어느 쪽으로도 이동할 수 없다. 프로그램은 쓰기 쉽지만, 성능은 매우 느리다.

다음 코드를 PTX로 컴파일해 보며 분석해 보자.

cpp
__global__ void kernel_seq_constant(int* array) {
  cuda::atomic<int> a;
  int val;
  // prior load/store
  int before = array[0];
  array[0] = 3;

  // atomic load
  val = a.load(cuda::std::memory_order_seq_cst);

  // later load
  int after = array[0];
  printf("before %d, after %d, val %d", before, after, val);
}

int main(int argc, const char* argv[]) {
  int* array;
  cudaMalloc(&array, sizeof(int) * 4);
  kernel_seq_constant<<<1, 2>>>(array);
  cudaDeviceSynchronize();
  return 0;
}

PTX를 보면:

ptx
fence.sc.sys

가 있어 Prior load/store가 atomic 이후로 이동하는 것을 막는다. 동시에 atomic load는

ptx
ld.acquire

를 사용해 후속 LD/ST가 이 명령어 앞에 실행되지 못하게 한다.

ptx
ld.global.u32 %r3, [%rd3];     // before = array[0]
st.global.u32 [%rd3], %r2;    // array[0] = 3

// begin inline asm
fence.sc.sys;                 // 후속 LD/ST가 앞으로 끌려오지 못하게 막음
// end inline asm

add.u64 %rd1, %SP, 0;
// begin inline asm
ld.acquire.sys.b32 %r1, [%rd1]; // acquire: 뒤의 LD/ST가 앞으로 끌려오는 것 방지
// end inline asm

ld.global.u32 %r4, [%rd3];    // after = array[0]

이제

cpp
val = a.load(cuda::std::memory_order_seq_cst);

를

cpp
val = a.load(cuda::std::memory_order_acquire);

로 바꿔 보면, PTX에서 fence.sc.sys가 제거된다.

ptx
ld.global.u32 %r3, [%rd3];     // before = array[0]
st.global.u32 [%rd3], %r2;    // array[0] = 3

add.u64 %rd1, %SP, 0;
// begin inline asm
ld.acquire.sys.b32 %r1, [%rd1];   // acquire: 뒤의 LD/ST가 앞으로 끌려오지 못하게
// end inline asm

ld.global.u32 %r4, [%rd3];    // after = array[0]

이제 atomic load 이전의 LD/ST는 atomic 이후로 이동할 수 있지만, atomic 이후의 Later Load는 여전히 막힌다.

2.2.2 Acquire와 2.2.3 Release

앞서 "뒤"를 막는 acquire가 있다면, "앞"의 LD/ST만 막고 뒤는 막지 않는 모델도 있어야 한다. 그것이 release 모드다.

cpp
__global__ void kernel_release(int* array) {
  cuda::atomic<int> a;
  // prior LD/ST
  int before = array[0];
  array[0] = 3;

  // atomic store.release
  a.store(1, cuda::std::memory_order_release);

  // later load
  int after = array[0];
  printf("before %d, after %d", before, after);
}

PTX를 보면 st.release 명령이 사용된 것을 볼 수 있다. 이 명령은 앞선 LD/ST를 막지만, Later Load는 앞당겨 실행될 수 있다.

ptx
ld.global.u32 %r3, [%rd3];      // before = array[0]
st.global.u32 [%rd3], %r2;     // array[0] = 3
mov.u32 %r1, 1;
add.u64 %rd1, %SP, 0;
// begin inline asm
st.release.sys.b32 [%rd1], %r1; // store.release
// end inline asm

ld.global.u32 %r4, [%rd3];      // later load

2.2.4 Relaxed

마지막은 가장 느슨한 Relaxed 메모리 모델이다. 앞·뒤 LD/ST 모두 자유롭게 재정렬될 수 있다.

cpp
__global__ void kernel_relaxed(int* array) {
  cuda::atomic<int> a;
  // prior LD/ST
  int before = array[0];
  array[0] = 3;

  // atomic store.relaxed
  a.store(1, cuda::std::memory_order_relaxed);

  // later load
  int after = array[0];
  printf("before %d, after %d", before, after);
}

PTX:

ptx
ld.global.u32 %r3, [%rd3];      // before = array[0]
st.global.u32 [%rd3], %r2;     // array[0] = 3
mov.u32 %r1, 1;
add.u64 %rd1, %SP, 0;
// begin inline asm
st.relaxed.sys.b32 [%rd1], %r1; // store.relaxed
// end inline asm

ld.global.u32 %r4, [%rd3];      // later load

2.2.5 Scope

앞선 예제의 PTX 명령들에는 모두 .sys 속성이 붙어 있다. 실제로는 필요에 따라 다른 "범위(scope)"를 선택할 수 있다. CUDA C++ API에서 정의하는 scope는 다음과 같다.

PTX에서 scope 정의는 다음과 같다.

Nvidia GPU의 메모리 계층 구조를 다시 보면, Block Scope는 SM 내부에서 L1 Cache를 기반으로 일관성을 유지하는 범위다.

Cluster Scope는 Thread Block Cluster 수준에서의 범위이며, 하드웨어 관점에서는 GPC 내에서 L2 Cache를 기반으로 일관성을 유지한다.

Device Scope는 GPU 칩 전체, 즉 모든 SM에 대해 L2를 기반으로 일관성을 유지하는 범위다.

Sys Scope는 전체 시스템을 포함하는 범위다.

Block Scope 예제

간단한 block_scope 예제를 보자.

cpp
#include<iostream>
#include<cuda.h>
#include<cuda/atomic>

#define CUDAASSERT(condition)  \
 if(!(condition))  \
 {  \
   printf("Assertion %s failed!\n", #condition); \
 }

__device__ void producer(
  cuda::atomic_ref<int, cuda::thread_scope_block> val)
{
  val.store(42, cuda::memory_order_relaxed);
}

__device__ void consumer(
  cuda::atomic_ref<int, cuda::thread_scope_block> val)
{
  volatile int tmp = -1;
  while(tmp == -1) {
    tmp = val.load(cuda::memory_order_relaxed);
  }
  CUDAASSERT(tmp == 42);
}

__global__ void kernel_scope_test(int* array)
{
  if(blockIdx.x == 0) {
    producer(array[0]);
  } else {
    consumer(array[0]);
  }
}

int main(int argc,const char*argv[]) {
  int* array;
  cudaMalloc(&array, sizeof(int) * 4);
  dim3 grid(2,1);
  kernel_scope_test<<<grid,1>>>(array);
  cudaDeviceSynchronize();
  return 0;
}

cuda::atomic_ref<int, cuda::thread_scope_block>로 정의했기 때문에, PTX에서 LD/ST relaxed의 scope는 .cta로 나온다.

ptx
// consumer
mov.u64 %rd6, %rd15;
// begin inline asm
ld.relaxed.cta.b32 %r3, [%rd6];
// end inline asm

// producer
mov.u64 %rd12, %rd15;
mov.u32 %r5, 42;
// begin inline asm
st.relaxed.cta.b32 [%rd12], %r5;
// end inline asm

실행 결과는 다음과 같다. 다른 block의 load는 다른 데이터 경로를 타기 때문에 막히지 않는다.

Scope를 device level로 확대해 보자.

cpp
__device__ void producer(
  cuda::atomic_ref<int, cuda::thread_scope_device> val)
{
  val.store(42, cuda::memory_order_relaxed);
}

__device__ void consumer(
  cuda::atomic_ref<int, cuda::thread_scope_device> val)
{
  int tmp = -1;
  while(tmp == -1) {
    tmp = val.load(cuda::memory_order_relaxed);
  }
  CUDAASSERT(tmp == 42);
}

PTX scope는 .gpu로 바뀐다.

ptx
// consumer
mov.u64 %rd6, %rd15;
// begin inline asm
ld.relaxed.gpu.b32 %r3, [%rd6];
// end inline asm
setp.eq.s32 %p2, %r3, -1;
@%p2 bra $L__BB0_2;

// producer
mov.u64 %rd12, %rd15;
mov.u32 %r5, 42;
// begin inline asm
st.relaxed.gpu.b32 [%rd12], %r5;
// end inline asm

그러나 프로그램은 Nvidia GTC25 PPT에 나오는 것처럼 "잘" 작동하지 않는다. 왜일까?

Scope를 GPU 전체로 설정했더라도, GTC25 발표자가 아마도 relaxed order에 대해 약간의 typo를 한 듯하다. 물론 이건 본질은 아니다. 보다 올바른 접근은 flag를 사용하고 release/acquire로 처리하는 것이다.

cpp
#include<iostream>
#include<cuda.h>
#include<cuda/atomic>

#define CUDAASSERT(condition)  \
 if(!(condition))  \
 {  \
   printf("Assertion %s failed!\n", #condition); \
 }

__device__ void producer(
 int& val,
 cuda::atomic_ref<int, cuda::thread_scope_device> flag)
{
  val = 42;
  flag.store(42, cuda::memory_order_relaxed);
}

__device__ void consumer(
 int& val,
 cuda::atomic_ref<int, cuda::thread_scope_device> flag)
{
  while(flag.load(cuda::memory_order_acquire) != -1) {
  }
  int tmp = val;
  CUDAASSERT(tmp == 42);
}

__global__ void kernel_scope_test(int* array)
{
  array[0] = 0;
  int flag = -1;
  __syncthreads();
  if(blockIdx.x == 0) {
    producer(array[0], flag);
  } else {
    consumer(array[0], flag);
  }
}

int main(int argc,const char*argv[]) {
  int* array;
  cudaMalloc(&array, sizeof(int) * 4);
  dim3 grid(2,1);
  kernel_scope_test<<<grid,1>>>(array);
  cudaDeviceSynchronize();
  cudaFree(&array);
  return 0;
}

Relaxed와 Acquire‑Release의 비교는 아래와 같다. Relaxed는 더 빠르며, 두 스레드가 단일 값을 교환하기만 하면 되는 경우에 유용하다. Release‑Acquire는 cache flush가 들어가기 때문에 더 느리고, 여러 스레드가 여러 값을 교환해야 하는 상황에 유리하다.

---

2.3 Async Thread

Ampere 세대부터 Async Thread 프로그래밍 능력이 도입되었다. 주로 GMEM→SMEM 비동기 복사(Load), SMEM→GMEM 비동기 Store를 제공해, 레지스터와 L1 점유를 피하면서 대용량 데이터 처리 시 처리량을 높이는 기능이다.

먼저 비동기 prefetch를 한다.

그다음 계산을 수행한다.

대략적인 계산 흐름은 다음과 같다. 자세한 코드는 아래에서 볼 수 있다.

https://github.com/zartbot/tensorcore_gemm/blob/main/05_pipeline_gmem_to_smem.cu

text
Async Copy A-Chunk from GMEM --> SMEM(Buffer_1)
Async Copy B-Chunk from GMEM --> SMEM(Buffer_1)
Wait for Async Copy Completion

for (size_t tile_k = CHUNK_K; tile_k < K_tiles; tile_k += CHUNK_K) {
  Swap Buffer_1/Buffer_2 Offset
  // Buffer_2를 비동기 로딩하면서 Buffer_1을 계산해 overlap
  Async Copy A-Chunk from GMEM --> SMEM(Buffer_2)
  Async Copy B-Chunk from GMEM --> SMEM(Buffer_2)

  for (size_t k_step = 0; k_step < CHUNK_K; ++k_step) {
    for (size_t i = 0; i < WT_COL_MMA_NUM; ++i) {
      Load-SMEM(Buffer_1)-to-A_fragment
      for (size_t j = 0; j < WT_ROW_MMA_NUM; ++j) {
        Load-SMEM(Buffer_1)-to-B_fragment
        wmma::mma_sync; // Tensor Core 계산
      }
    }
  }

  Wait for Async Copy Completion
}

Calculate Last Buffer
WarpTile WMMA-Store-to-SMEM
Store-SMEM -> GMEM

Hopper에서는 st.async 명령어를 도입해 DSMEM으로의 비동기 저장을 지원한다. 자세한 테스트 코드는 3장에서 다룬다.

하지만 이런 비동기 연산들은 새로운 문제를 낳는다. 비동기 데이터 경로가 하나라도 등장하면 Data Race 가능성이 생기므로, 각별히 주의해야 한다.

메모리 계층 구조에 데이터 경로가 여러 개 있다는 점, 특히 Hopper부터는 TMA가, Blackwell에서는 Tensor Memory가 도입된 점을 감안하면, 서로 다른 데이터 경로 사이의 Data Race를 더 잘 관리·추상화해야 한다. Hopper에서는 이를 위해 Async Proxy 개념을 도입하여 General Proxy와 Async Proxy로 메모리 접근 경로를 구분한다.

이렇게 구분하면 async proxy의 메모리 연산에 대해 fence를 설정할 수 있다.

반대로, Async Proxy의 연산 쪽에서는 일반적으로 memory barrier를 두고, General Proxy의 LD/ST가 이 barrier 완료를 기다리는 식으로 동작한다.

구체적인 절차는 다음과 같다.

1. SMEM에 mbarrier를 하나 할당한다.
2. 한 스레드가 TMA 명령(UBLKCP)을 issue한다.
3. SMEM에 복사되는 데이터는 반드시 alignas(16) bytes 이어야 한다.

완료 여부는 completion_tx 카운터 방식으로 관리된다.

TMA‑2D 예제

아래는 TMA‑2D 기반 예제다. (코드는 길지만 핵심은 cp_async_bulk_tensor_2d_global_to_shared와 mbarrier 사용 구조다.)

cpp
#include<cuda.h>
#include<cudaTypedefs.h>
#include<cuda/barrier>
#include<iostream>

#pragma nv_diag_suppress static_var_with_dynamic_init

using barrier_t = cuda::barrier<cuda::thread_scope_block>;
namespace cde = cuda::device::experimental;

constexpr size_t GLOBAL_M = 64;
constexpr size_t GLOBAL_K = 32;
constexpr size_t TILE_M = 8;
constexpr size_t TILE_K = 16;

inline PFN_cuTensorMapEncodeTiled get_cuTensorMapEncodeTiled() {
  cudaDriverEntryPointQueryResult driver_status;
  void* cuTensorMapEncodeTiled_ptr = nullptr;
  cudaGetDriverEntryPointByVersion("cuTensorMapEncodeTiled",
                                   &cuTensorMapEncodeTiled_ptr,
                                   12000, cudaEnableDefault, &driver_status);
  if (driver_status != cudaDriverEntryPointSuccess)
    throw std::runtime_error("driver_status != cudaDriverEntryPointSuccess");
  return reinterpret_cast<PFN_cuTensorMapEncodeTiled>(cuTensorMapEncodeTiled_ptr);
}

CUtensorMap make_2d_tma_desc(int32_t* global_address,
                              uint64_t global_dim[2], uint64_t stride,
                              uint32_t smem_dim[2], CUtensorMapSwizzle swizzle)
{
  CUtensorMap tensor_map = {};
  uint64_t global_stride[1] = { stride };
  uint32_t elem_stride[2] = { 1, 1 };
  auto encode = get_cuTensorMapEncodeTiled();
  auto res = encode(&tensor_map,
                    CUtensorMapDataType::CU_TENSOR_MAP_DATA_TYPE_INT32,
                    2, // rank = 2
                    global_address,
                    global_dim,
                    global_stride,
                    smem_dim,
                    elem_stride,
                    CUtensorMapInterleave::CU_TENSOR_MAP_INTERLEAVE_NONE,
                    swizzle,
                    CUtensorMapL2promotion::CU_TENSOR_MAP_L2_PROMOTION_L2_256B,
                    CUtensorMapFloatOOBfill::CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);
  assert(res == CUDA_SUCCESS && "make tma descriptor failed.");
  return tensor_map;
}

__global__ void tma_kernel(const __grid_constant__ CUtensorMap tensor_map,
                           uint32_t x, uint32_t y)
{
  __shared__ alignas(128) int tile_smem[TILE_M * TILE_K];
  __shared__ barrier_t bar;

  // Barrier 초기화
  if (threadIdx.x == 0) {
    init(&bar, blockDim.x);
    // TMA 호출 경로는 async proxy이므로 fence로 가시성 보장 필요
    cde::fence_proxy_async_shared_cta();
  }
  __syncthreads();

  barrier_t::arrival_token token;

  if (threadIdx.x == 0) {
    // TMA 복사 수행
    cde::cp_async_bulk_tensor_2d_global_to_shared(tile_smem,
                                                  &tensor_map, x, y, bar);
    token = cuda::device::barrier_arrive_tx(bar, 1, sizeof(tile_smem));
  } else {
    token = bar.arrive();
  }

  // 다른 일 수행...
  int value = threadIdx.x * 100 + threadIdx.x;

  // 모든 데이터 도착 대기
  bar.wait(std::move(token));
  printf("[tma_kernel] threadIdx.x %d arrived\n", threadIdx.x);

  for (int i = 0; i < TILE_M * TILE_K; i += blockDim.x) {
    tile_smem[i + threadIdx.x] += value;
  }

  // async proxy가 SMEM에 기록 완료되기를 fence
  cuda::ptx::fence_proxy_async(cuda::ptx::space_shared);
  __syncthreads();

  if (threadIdx.x == 0) {
    // TMA로 SMEM -> GMEM 복사
    cde::cp_async_bulk_tensor_2d_shared_to_global(&tensor_map, x, y, tile_smem);
    cuda::ptx::cp_async_bulk_commit_group();
    cuda::ptx::cp_async_bulk_wait_group_read(cuda::ptx::n32_t<0>());
  }

  printf("thread %d done\n", threadIdx.x);
}

int main(int argc, char** argv)
{
  uint64_t global_dim[2] = { GLOBAL_M, GLOBAL_K };
  size_t GLOBAL_SIZE = GLOBAL_K * GLOBAL_M;
  uint32_t tile_dim[2] = { TILE_M, TILE_K };

  int h_data[GLOBAL_SIZE];
  for (size_t i = 0; i < GLOBAL_SIZE; ++i) {
    h_data[i] = 1;
  }

  // GPU6 메모리 할당 및 P2P 허용
  cudaSetDevice(6);
  cudaDeviceEnablePeerAccess(7, 0);

  int* d_data;
  cudaMalloc(&d_data, GLOBAL_SIZE * sizeof(int));
  cudaMemcpy(d_data, h_data, GLOBAL_SIZE * sizeof(int), cudaMemcpyHostToDevice);

  // GPU7에서 NVLINK를 통해 TMA 테스트
  cudaSetDevice(7);
  cudaDeviceEnablePeerAccess(6, 0);

  CUtensorMap tensor_map = make_2d_tma_desc(
    d_data, global_dim, GLOBAL_K * sizeof(int), tile_dim,
    CUtensorMapSwizzle::CU_TENSOR_MAP_SWIZZLE_NONE);

  uint32_t coord_x = 16;
  uint32_t coord_y = 16;

  tma_kernel<<<1, TILE_M * TILE_K>>>(tensor_map, coord_x, coord_y);
  cudaDeviceSynchronize();

  cudaError_t err = cudaGetLastError();
  std::cout << cudaGetErrorString(err) << std::endl;

  cudaSetDevice(6);
  cudaMemcpy(h_data, d_data, GLOBAL_SIZE * sizeof(int), cudaMemcpyDeviceToHost);

  for (size_t i = 0; i < GLOBAL_M; ++i) {
    for (size_t j = 0; j < GLOBAL_K; ++j) {
      printf("%5d ", h_data[i * GLOBAL_K + j]);
    }
    printf("\n");
  }

  cudaFree(d_data);
  return 0;
}

메모리 비동기 복사 종류에 따라 완료 메커니즘이 서로 다르다는 점에 유의해야 한다.

또한 mbarrier의 경우, 어떤 것은 completion.tx 카운트 방식을 사용하고, 또 어떤 것은 waitgroup 방식을 사용한다.

마지막으로 저자가 정리한 내용을 보면, st.async / red.async / cp.async는 구현 측면에서 더 이르고, 데이터 경로 측면에서는 async.proxy를 지원하지 않는다. 반면 TMA/TMEM/WGMMA는 async proxy를 지원한다.

---

3. Thread Block Cluster와 저지연 동기화

Hopper의 Cluster 프로그래밍에 대한 공개 자료는 많지 않다. 《CUDA C Programming Guide》[5]에 일부 소개가 있다. Thread Block Cluster는 Hopper에서 도입된 새로운 계층이다. Hopper 내부에는 국소적인 SM‑to‑SM 데이터 경로가 구축되어 있고, Distribute Shared Memory(DSMEM) 개념을 제공한다.

소프트웨어 관점에서 보면, Grid와 Block 사이에 Cluster라는 계층을 새로 도입한 것이다.

간단한 예제를 보자. Kernel 함수 정의에서

cpp
__cluster_dims__(x, y, z)

를 사용해 cluster의 형태를 결정하고, cg::this_cluster()로 현재 cluster의 descriptor를 얻을 수 있다. 이식성을 고려해 단일 cluster에는 최대 8개의 Thread Block만 지원한다고 CUDA 문서에 적혀 있지만, Hopper의 DataSheet를 보면 H100은 8개의 GPC와 132개의 SM이 있어, 단일 Cluster에 최대 16개 Block도 가능하다. 다만 H20에서는 SM이 줄어서 실제 테스트 상 최대 8개까지만 지원되었다.

cpp
#include<iostream>
#include<cuda.h>
#include<cuda/atomic>
#include<cooperative_groups.h>

namespace cg = cooperative_groups;

__global__ void __cluster_dims__(4,2,1) kernel_cluster_test()
{
  cg::cluster_group cluster = cg::this_cluster();
  unsigned int cluster_block_rank = cluster.block_rank();

  printf("ThreadIdx [%d,%d,%d], BlockDIM [%d,%d,%d], BlockIdx [%d,%d,%d] "
         "Cluster rank %d dim [%d,%d,%d] idx [%d,%d,%d] GridDim [%d,%d,%d]\n",
         threadIdx.x, threadIdx.y, threadIdx.z,
         blockDim.x, blockDim.y, blockDim.z,
         blockIdx.x, blockIdx.y, blockIdx.z,
         cluster.block_rank(),
         cluster.dim_blocks().x, cluster.dim_blocks().y, cluster.dim_blocks().z,
         cluster.block_index().x, cluster.block_index().y, cluster.block_index().z,
         gridDim.x, gridDim.y, gridDim.z);
}

int main(int argc, const char* argv[])
{
  dim3 grid(4,8,1);
  dim3 block(4,4,4);
  kernel_cluster_test<<<grid, block>>>();
  cudaError_t err = cudaGetLastError();
  std::cout << cudaGetErrorString(err) << std::endl;
  cudaDeviceSynchronize();
  return 0;
}

__cluster_dims__ 대신 cudaLaunchKernelEx로 런타임에 cluster dim을 설정할 수도 있다.

cpp
__global__ void kernel_cluster_test(int var1, int var2) {}

int main()
{
  dim3 grid(4,8,1);
  dim3 block(4,4,4);

  cudaLaunchConfig_t config = {0};
  config.gridDim = grid;
  config.blockDim = block;

  cudaLaunchAttribute attribute[1];
  attribute[0].id = cudaLaunchAttributeClusterDimension;
  attribute[0].val.clusterDim.x = 2;
  attribute[0].val.clusterDim.y = 1;
  attribute[0].val.clusterDim.z = 1;

  config.attrs = attribute;
  config.numAttrs = 1;

  cudaLaunchKernelEx(&config, kernel_cluster_test, var1, var2);
}

Cluster 관련 기타 함수는 CUDA C Programming Guide의 Cluster group[6] 절에 정리되어 있다.

3.1 Distributed Shared Memory (DSMEM)

Cluster의 핵심 사용 사례는 Distributed Shared Memory다. Cluster 내부에서는 L2를 우회(bypass)하는 저지연 SMEM 상호 접근이 가능하며, LD/ST, ATOMIC, async DMA 등의 연산을 지원한다.

예를 들어 아래 예제를 보자.

cpp
#include<cstdio>
#include<iostream>
#include<cuda/ptx>
#include<cuda/barrier>
#include<cooperative_groups.h>

namespace cg = cooperative_groups;

__global__ void __cluster_dims__(8,1,1) kernel()
{
  cg::cluster_group cluster = cg::this_cluster();

  // SMEM 선언/초기화
  __shared__ int smem_x[32];
  smem_x[threadIdx.x] = blockIdx.x * 10000 + threadIdx.x;

  // cluster 내 동기화: 모든 스레드가 공유 메모리 초기화를 끝냈는지 보장
  cluster.sync();

  int peer_rank = cluster.block_rank() ^ 1;
  int* dst_mem = cluster.map_shared_rank(smem_x, peer_rank);

  dst_mem[threadIdx.x] += cluster.block_rank() * 100;
  cluster.sync();

  printf("threadIdx %d blockIdx %d clusterRank %d smem: %d\n",
         threadIdx.x, blockIdx.x, cluster.block_rank(), smem_x[threadIdx.x]);
}

int main()
{
  kernel<<<8,4>>>();
  cudaDeviceSynchronize();
  return 0;
}

실행 결과 예:

text
threadIdx 0 blockIdx 6 clusterRank 6 smem:60700
threadIdx 1 blockIdx 6 clusterRank 6 smem:60701
...
threadIdx 0 blockIdx 0 clusterRank 0 smem:100
...

Thread Block Cluster의 장점은 SM‑to‑SM 네트워크를 통해 데이터를 주고받으므로, 데이터를 L2/GMEM에 쓰지 않고도 통신할 수 있다는 점이다.

그러나 앞의 cluster.sync() 구현은 성능 병목이 있다. Cluster 전체를 동기화하면서, cluster 내 다른 thread에서도 LD/ST 결과가 보이도록 만들기 때문에 데이터가 L2를 가로지르게 된다. 실제로 PTX에서는 cluster::sync()가 다음 두 명령으로 확장된다.

ptx
barrier.cluster.arrive;
barrier.cluster.wait;

하지만 PTX 명령을 직접 사용해 arrive와 wait를 분리하고, 필요한 경우에만 release/relaxed를 선택적으로 써서 LD/ST 가시성을 제어할 수 있다.

예를 들어 barrier 초기화는 cluster::sync()를 써도 되지만, 이 경우 L2를 통과하므로 느리다.

이를 release‑acquire 방식으로 바꾸어 최적화할 수 있다.

그런 다음 Cluster 내 SM‑SM 통신은 비동기 store와 local mbarrier 대기 방식을 사용한다.

전체 코드는 다음과 같다. (핵심은 PTX의 mbarrier와 st_async, barrier.cluster.*를 직접 사용하는 부분이다.)

cpp
#include<cstdio>
#include<cuda/ptx>
#include<cuda/barrier>
#include<cooperative_groups.h>

namespace cg = cooperative_groups;

using cuda::ptx::scope_cluster;
using cuda::ptx::sem_acquire;
using cuda::ptx::sem_relaxed;
using cuda::ptx::sem_release;
using cuda::ptx::space_cluster;
using cuda::ptx::space_shared;

namespace ptx {

__device__ __forceinline__ uint32_t __as_ptr_smem(const void* __ptr)
{
  return static_cast<uint32_t>(__cvta_generic_to_shared(__ptr));
}

__device__ __forceinline__ void mbarrier_init(uint64_t* mbar, const uint32_t count)
{
  uint32_t mbar_ptr = __cvta_generic_to_shared(mbar);
  asm volatile("mbarrier.init.shared.b64 [%0], %1;" ::
               "r"(mbar_ptr), "r"(count) : "memory");
}

__device__ __forceinline__ void fence_mbarrier_init(cuda::ptx::sem_release_t,
                                                    cuda::ptx::scope_cluster_t)
{
  asm volatile("fence.mbarrier_init.release.cluster; // 3." ::: "memory");
}

__device__ __forceinline__ void barrier_cluster_arrive(cuda::ptx::sem_relaxed_t)
{
  asm volatile("barrier.cluster.arrive.relaxed;" :::);
}

__device__ __forceinline__ void barrier_cluster_wait(cuda::ptx::sem_acquire_t)
{
  asm volatile("barrier.cluster.wait.acquire;" ::: "memory");
}

__device__ __forceinline__ void barrier_cluster_wait()
{
  asm volatile("barrier.cluster.wait;" ::: "memory");
}

template<cuda::ptx::dot_scope Scope>
__device__ __forceinline__ uint64_t mbarrier_arrive_expect_tx(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> __scope,
  cuda::ptx::space_shared_t,
  uint64_t* __addr,
  const uint32_t& __txCount)
{
  uint64_t __state;
  if constexpr (__scope == cuda::ptx::scope_cta) {
    asm("mbarrier.arrive.expect_tx.relaxed.cta.shared::cta.b64 %0, [%1], %2;"
        : "=l"(__state)
        : "r"(__as_ptr_smem(__addr)), "r"(__txCount)
        : "memory");
  } else if constexpr (__scope == cuda::ptx::scope_cluster) {
    asm("mbarrier.arrive.expect_tx.relaxed.cluster.shared::cta.b64 %0, [%1], %2;"
        : "=l"(__state)
        : "r"(__as_ptr_smem(__addr)), "r"(__txCount)
        : "memory");
  }
  return __state;
}

template<cuda::ptx::dot_scope Scope>
__device__ __forceinline__ bool mbarrier_try_wait(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> __scope,
  uint64_t* __addr,
  const uint64_t& __state)
{
  uint32_t __waitComplete;
  if constexpr (__scope == cuda::ptx::scope_cta) {
    asm("{\n\t .reg .pred P_OUT; \n\t"
        "mbarrier.try_wait.acquire.cta.shared::cta.b64 P_OUT, [%1], %2;\n\t"
        "selp.b32 %0, 1, 0, P_OUT; \n"
        "}"
        : "=r"(__waitComplete)
        : "r"(__as_ptr_smem(__addr)), "l"(__state)
        : "memory");
  } else if constexpr (__scope == cuda::ptx::scope_cluster) {
    asm("{\n\t .reg .pred P_OUT; \n\t"
        "mbarrier.try_wait.acquire.cluster.shared::cta.b64 P_OUT, [%1], %2;\n\t"
        "selp.b32 %0, 1, 0, P_OUT; \n"
        "}"
        : "=r"(__waitComplete)
        : "r"(__as_ptr_smem(__addr)), "l"(__state)
        : "memory");
  }
  return static_cast<bool>(__waitComplete);
}

} // namespace ptx

__global__ void __cluster_dims__(8,1,1) low_latency_kernel(int iter_num)
{
  cg::cluster_group cluster = cg::this_cluster();

  __shared__ int      receive_buffer[4];
  __shared__ uint64_t bar;

  // barrier 초기화
  if (threadIdx.x == 0) {
    ptx::mbarrier_init(&bar, blockDim.x);
  }

  // barrier 가시화
  ptx::fence_mbarrier_init(sem_release, scope_cluster);
  ptx::barrier_cluster_arrive(sem_relaxed);
  ptx::barrier_cluster_wait(sem_acquire);

  // 원격 buffer/barrier 주소 획득
  unsigned int peer_rank = cluster.block_rank() ^ 1;
  uint64_t*    remote_bar = cluster.map_shared_rank(&bar, peer_rank);
  int*         remote_buffer = cluster.map_shared_rank(&receive_buffer[0], peer_rank);

  for (int iter = 0; iter < iter_num; ++iter) {
    cuda::ptx::st_async(remote_buffer, {iter, iter, iter, iter}, remote_bar);

    uint64_t token = ptx::mbarrier_arrive_expect_tx(
      sem_relaxed, scope_cluster, space_shared, &bar, sizeof(receive_buffer));

    bool ready = false;
    while (!ready) {
      ready = ptx::mbarrier_try_wait(sem_acquire, scope_cluster, &bar, token);
    }

    ptx::barrier_cluster_arrive(sem_relaxed);
    ptx::barrier_cluster_wait();
  }
}

__global__ void __cluster_dims__(8,1,1) standard_async_kernel(int iter_num)
{
  cg::cluster_group cluster = cg::this_cluster();

  using barrier_t = cuda::barrier<cuda::thread_scope_block>;

  __shared__ int       receive_buffer[4];
  __shared__ barrier_t bar;

  init(&bar, blockDim.x);
  cluster.sync();

  unsigned int other_block_rank = cluster.block_rank() ^ 1;
  uint64_t*    remote_bar = cluster.map_shared_rank(
    cuda::device::barrier_native_handle(bar), other_block_rank);
  int* remote_buffer = cluster.map_shared_rank(&receive_buffer[0], other_block_rank);

  for (int iter = 0; iter < iter_num; ++iter) {
    uint64_t arrival_token = cuda::ptx::mbarrier_arrive_expect_tx(
      sem_release, scope_cluster, space_shared,
      cuda::device::barrier_native_handle(bar), sizeof(receive_buffer));

    cuda::ptx::st_async(remote_buffer, {iter, iter, iter, iter}, remote_bar);

    while (!cuda::ptx::mbarrier_try_wait(
              sem_acquire, scope_cluster,
              cuda::device::barrier_native_handle(bar), arrival_token)) {
    }
  }
}

int main()
{
  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);

  int   num_iter = 10000;
  float time;

  cudaEventRecord(start);
  low_latency_kernel<<<128,32>>>(num_iter);
  cudaEventRecord(stop);
  cudaDeviceSynchronize();
  cudaEventElapsedTime(&time, start, stop);
  printf("low latency kernel elapsed %f\n", time);

  cudaEventRecord(start);
  standard_async_kernel<<<128,32>>>(num_iter);
  cudaEventRecord(stop);
  cudaDeviceSynchronize();
  cudaEventElapsedTime(&time, start, stop);
  printf("async kernel elapsed %f\n", time);
}

st_async에서 cluster‑scope relaxed를 사용하고, local barrier wait에서 acquire를 사용한다.

H20에서 실제 테스트하면 mbarrier_arrive_expect_tx(sem_release)를 사용하는 표준 방식 대비 약 46% 빠르다.

text
low latency kernel elapsed 2.068736
async kernel elapsed 3.714880

3.2 TMA Multicast

TMA에는 Multicast 기능도 추가되어, 하나의 block에서 로드한 데이터를 여러 block에 동시에 배포할 수 있다. 예제 코드는 다음과 같다. 컴파일 시 다음과 같은 경고를 볼 수 있는데, multicast::cluster는 sm_90a/sm_100a/sm_101a에서 사용하는 것이 더 좋다고 알려 주는 것이다.

text
ptxas ... warning : Advisory: '.multicast::cluster' modifier ...

cpp
#include<cuda.h>
#include<cudaTypedefs.h>
#include<cooperative_groups.h>
#include<cuda/barrier>
#include<iostream>

#pragma nv_diag_suppress static_var_with_dynamic_init

using barrier_t = cuda::barrier<cuda::thread_scope_block>;
namespace cde = cuda::device::experimental;
namespace cg  = cooperative_groups;

const int ARRAY_SIZE  = 512;
const int TILE_SIZE   = 16;
const int CLUSTER_DIM = 8;

inline PFN_cuTensorMapEncodeTiled get_cuTensorMapEncodeTiled() { /* 생략 */ }

CUtensorMap make_1d_tma_desc(int32_t* global_address,
                              uint64_t global_dim,
                              uint32_t smem_dim)
{
  CUtensorMap tensor_map = {};
  uint64_t global_size[1] = { global_dim };
  uint64_t global_stride[1] = { global_dim * sizeof(int) };
  uint32_t tile_size[1] = { smem_dim };
  uint32_t elem_stride[1] = { 1 };
  auto encode = get_cuTensorMapEncodeTiled();
  auto res = encode(&tensor_map,
                    CUtensorMapDataType::CU_TENSOR_MAP_DATA_TYPE_INT32,
                    1,
                    global_address,
                    global_size,
                    global_stride,
                    tile_size,
                    elem_stride,
                    CUtensorMapInterleave::CU_TENSOR_MAP_INTERLEAVE_NONE,
                    CUtensorMapSwizzle::CU_TENSOR_MAP_SWIZZLE_NONE,
                    CUtensorMapL2promotion::CU_TENSOR_MAP_L2_PROMOTION_L2_256B,
                    CUtensorMapFloatOOBfill::CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);
  assert(res == CUDA_SUCCESS && "make tma descriptor failed.");
  return tensor_map;
}

__global__ void __cluster_dims__(CLUSTER_DIM,1,1)
tma_kernel(const __grid_constant__ CUtensorMap tensor_map,
           uint32_t coord)
{
  __shared__ alignas(16) int tile_smem[TILE_SIZE];
  __shared__ barrier_t   bar;

  cg::cluster_group cluster = cg::this_cluster();
  unsigned int       cluster_rank = cluster.block_rank();

  if (threadIdx.x == 0) {
    init(&bar, blockDim.x);
    cde::fence_proxy_async_shared_cta();
  }
  __syncthreads();

  barrier_t::arrival_token token;
  if ((cluster_rank == 0) && (threadIdx.x == 0)) {
    uint16_t ctaMask = 0b10111011;
    asm volatile (
      "cp.async.bulk.tensor.1d.shared::cluster.global.tile.mbarrier::"
      "complete_tx::bytes.multicast::cluster "
      "[%0], [%1, {%2}], [%3], %4;\n"
      :
      : "r"(static_cast<_CUDA_VSTD::uint32_t>(
                __cvta_generic_to_shared(tile_smem))),
        "l"(&tensor_map),
        "r"(coord),
        "r"(static_cast<_CUDA_VSTD::uint32_t>(
                __cvta_generic_to_shared(
                  cuda::device::barrier_native_handle(bar)))),
        "h"(ctaMask)
      : "memory");

    token = cuda::device::barrier_arrive_tx(bar, 1, sizeof(tile_smem));
  } else {
    token = bar.arrive();
  }

  bar.wait(std::move(token));
  cluster.sync();

  if (threadIdx.x == 0) {
    printf("cluster %d smem[%d %d %d %d]\n",
           cluster_rank,
           tile_smem[0], tile_smem[1], tile_smem[2], tile_smem[3]);
  }
}

int main(int argc, char** argv)
{
  int* h_data = nullptr;
  cudaHostAlloc(&h_data, ARRAY_SIZE * sizeof(int), cudaHostAllocMapped);

  for (size_t i = 0; i < ARRAY_SIZE; ++i) {
    h_data[i] = i;
  }

  int* d_data;
  cudaHostGetDevicePointer(&d_data, h_data, 0);

  CUtensorMap tensor_map = make_1d_tma_desc(d_data, ARRAY_SIZE, TILE_SIZE);
  uint32_t    coord = 3 * TILE_SIZE;

  tma_kernel<<<CLUSTER_DIM, 32>>>(tensor_map, coord);
  cudaDeviceSynchronize();

  cudaError_t err = cudaGetLastError();
  cudaFree(d_data);
  return 0;
}

출력 예:

text
cluster 6 smem[0 0 0 0]
cluster 7 smem[48 49 50 51]
cluster 0 smem[48 49 50 51]
cluster 1 smem[48 49 50 51]
...

---

4. 메모리 대역폭 최대화

이제 GTC25 세션의 또 다른 주제인 "Maximizing Memory Bandwidth"로 넘어가자.

Nvidia GPU에는 많은 CUDA Core가 있고, 이에 맞추어 메모리도 매우 깊은 계층 구조를 가진다.

연산 코어에 가까운 캐시일수록 접근 지연을 크게 줄여 준다. Shared Memory는 세대마다 용량이 늘고 있지만 거의 한계에 근접했다. 예를 들어 Hopper와 Blackwell의 SMEM 용량은 더 이상 증가하지 않았다. 대신 더 국소적인 온칩(on‑chip) 네트워크를 구축해, SMEM을 이 작은 네트워크를 통해 Distributed SMEM처럼 묶어 Data Locality를 높이고 GMEM 접근 지연을 피하는 방향으로 진화하고 있다.

그래서 3장에서 언급했듯이, Hopper와 Blackwell에서는 DSMEM 기반 프로그래밍과 Cluster 추상이 더욱 중요해졌다.

Blackwell은 또한 Cluster의 혼합 배치 스케줄링을 개선해 자원 활용률을 높였다.

하지만 칩 면적 제약 때문에 SM 증가 속도는 둔화되고 있다. 더 큰 SMEM, 더 큰 연산 성능, 더 많은 SM을 동시에 갖고 싶다면, 결국 무언가를 포기해야 한다. 그래서 FP64가 잘려 나갔다…

한편으로 HBM3e/HBM4 덕분에 메모리 대역폭은 계속 증가하지만, 연산과 메모리 접근 사이의 튜닝 난도도 점점 높아지고 있다.

분산 시스템의 효율을 최적화할 때, 미시적으로는 메모리 접근 지연, 거시적으로는 부하 균형, 네트워크 상 혼잡 제어 등 여러 요소를 고려해야 한다. 이를 분석할 때는 대기행렬이론(Queueing Theory)의 시각에서 접근하는 것이 효과적이다. 통계·확률·비동기성을 이용해 지연을 다루어야 한다. 앞에서 Kingman 공식 등을 언급했는데, 여기서는 조금 더 단순한 모델인 Little's Law를 사용해 보자.

안정 상태의 시스템에서는, 시스템 내 평균 고객 수 L은 평균 도착률 λ와 고객이 시스템 내에서 보내는 평균 시간 W의 곱과 같다. 직관적인 공식이다.

예를 들어 에스컬레이터를 타는 시스템을 가정하자. 평균 2초마다 1명씩 도착(= 평균 1초에 1/2명), 에스컬레이터를 타는 데 40초가 걸린다면, 시스템이 동시에 수용할 수 있는 인원은 20명이다.

메모리 접근도 비슷하다. 메모리 대역폭과 평균 메모리 접근 지연을 알면 inflight‑bytes(동시에 날아가고 있는 데이터 양)를 계산할 수 있다. Hopper에서는 메모리 대역폭을 가득 쓰려면 약 32KB inflight가 필요하고, Blackwell에서는 거의 두 배인 64KB 정도가 필요하다.

간단한 커널의 경우, 한 스레드가 메모리를 몇 번 접근하는지, 한 번에 몇 byte를 접근하는지, 하나의 block에 몇 개 스레드가 있는지, 한 SM에 몇 개 block이 올라가는지를 곱해, 단일 SM 수준 inflight bytes를 계산할 수 있다.

Inflight bytes를 늘리는 방법은 크게 세 가지다.

1. 명령어 수준 병렬성(ILP)
2. 데이터 수준 병렬성(DLP)
3. 비동기 메모리 접근

예를 들어 UNROLL로 루프를 펼치면 한 번에 더 많은 명령을 병렬로 실행할 수 있다.

또한 Vector Load로 데이터 병렬성을 늘릴 수 있다.

하지만 데이터 병렬성이나 명령어 병렬성을 늘리면 레지스터 압력이 급격히 증가한다.

그래서 비동기 복사가 등장했다. 레지스터를 점유하지 않고 inflight bytes를 늘리는 방식이다.

동시에 비동기 로드는 데이터 복사와 계산을 겹쳐(overlap) 실행할 수 있게 한다.

또한 Producer‑Consumer 패턴을 warp specialization과 결합해 사용할 수 있다.

예를 들어 일부 thread를 메모리 복사 전담 Producer로 사용하고,

Consumer가 계산을 수행하면서, 동시에 또 다른 일부 thread가 데이터를 미리 prefetch하는 구조를 사용할 수 있다.

GTC25 세션에서는 로딩 최적화에 대한 권장 패턴도 보여 준다.

---

5. LD/ST 명령과 Cache 제어

명령어 병렬성과 데이터 병렬성을 늘리고, 대량의 비동기 복사를 사용하는 상황에서는 레지스터와 캐시 압력을 줄이는 것이 매우 중요하다. Ampere에서 도입된 Async Copy는 레지스터와 L1 Cache를 우회(bypass)해 GMEM에서 SMEM으로 직접 로드할 수 있게 한다. Hopper에서는 TMA를 도입해 n‑D 행렬 같은 패턴의 issue 수를 줄였고, Blackwell에서는 Tensor Memory를 도입해 MMA 시 레지스터 점유를 줄였다.

이번 장에서는 일반적인 LD/ST 명령이 캐시를 어떻게 제어하는지 살펴보자.

DeepEP에서 사용하는 "ld.global.nc.L1::no_allocate.L2::256B", "st.global.L1::no_allocate" 같은 패턴이 대표적이다. DeepEP의 아래 파일은 참고할 가치가 있다.

https://github.com/deepseek-ai/DeepEP/blob/main/csrc/kernels/utils.cuh

공식 PTX 문서에는 ld 명령의 다양한 사용법이 정리되어 있다.

text
ld{.weak}{.ss}{.cop}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{.unified}{, cache-policy};
ld{.weak}{.ss}{.level::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{.unified}{, cache-policy};
ld.volatile{.ss}{.level::prefetch_size}{.vec}.type d, [a];
ld.relaxed.scope{.ss}{.level::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{, cache-policy};
ld.acquire.scope{.ss}{.level::eviction_priority}{.level::cache_hint}{.level::prefetch_size}{.vec}.type d, [a]{, cache-policy};
ld.mmio.relaxed.sys{.global}.type d, [a];

.ss = { .const, .global, .local, .param{::entry, ::func}, .shared{::cta, ::cluster} };
.cop = { .ca, .cg, .cs, .lu, .cv };
.level::eviction_priority = { .L1::evict_normal, .L1::evict_unchanged,
                              .L1::evict_first, .L1::evict_last, .L1::no_allocate };
.level::cache_hint = { .L2::cache_hint };
.level::prefetch_size = { .L2::64B, .L2::128B, .L2::256B };
.scope = { .cta, .cluster, .gpu, .sys };
.vec = { .v2, .v4 };
.type = { .b8, .b16, .b32, .b64, .b128,
          .u8, .u16, .u32, .u64,
          .s8, .s16, .s32, .s64,
          .f32, .f64 };

.weak의 경우 컴파일된 SASS는 기본값과 동일하게 LDG.E가 되고, .volatile은 LDG.E.STRONG이 된다. relaxed/acquire는 앞에서 자세히 다루었다. .mmio는 PTX 8.2에 추가되었고 SM_70(Volta) 이후 아키텍처에서만 지원되며, SASS는 LDG.E.MMIO.SYS로 나온다.

5.1 Cache Operation (.cop)

.cop는 성능 튜닝 시 매우 유용하다. 캐시 동작 정책을 정의한다.

• .ca: 모든 계층 캐시에 저장. 즉 L1/L2 모두 cache. 기본 동작.
• .cg: L2에만 cache하고 L1에는 cache하지 않는다.
• .cs: 한 번만 접근될 가능성이 높은 데이터에 사용. L1/L2 cache에서 evict‑first 정책을 적용한다. SASS에 EF 속성이 붙는다. 예를 들어 reduction 시 사용 가능.
• .lu: Last Use. spill된 레지스터 복원 또는 함수 스택 프레임을 pop할 때 불필요한 write를 줄인다. global address에 사용하면 .cs와 동일하게 동작.
• .cv: cache하지 않음.

흥미로운 질문 하나.

ld.weak.global.cv와 ld.volatile.global의 차이는?

없다. 둘 다 SASS에서 LDG.E.STRONG.SYS로 나온다.

그렇다면 ld.weak.global.cg는? SASS는 LDG.E.STRONG.GPU가 된다.

이외에도 L1 Cache의 eviction 정책, L1 할당 여부 등을 지정할 수 있다.

또한 L2 Cache Prefetch 크기 등도 설정할 수 있다. Scope 파라미터는 앞서 메모리 모델에서 살펴본 것과 동일하다.

5.2 Read‑Only Cache와 ld.global.nc

SM 내부에는 L1 Cache 외에 Read‑Only Memory가 하나 더 있다. 《CUDA Refresher: The CUDA Programming Model》[7]에서도 언급된다.

Read-only memory—Each SM has an instruction cache, constant memory, texture memory and RO cache, which is read-only to kernel code.

ld.global.nc를 사용하면 이 Read‑Only Cache를 선택적으로 사용할 수 있다. 텍스처 cache 크기가 충분히 크고, 지연이 충분한 병렬성으로 잘 가려질 수 있는 경우에 유용하다.

text
ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache-policy};
ld.global{.cop}.nc{.level::cache_hint}{.level::prefetch_size}.vec.type d, [a]{, cache-policy};
ld.global.nc{.level::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.type d, [a]{, cache-policy};
ld.global.nc{.level::eviction_priority}{.level::cache_hint}{.level::prefetch_size}.vec.type d, [a]{, cache-policy};

.cop = { .ca, .cg, .cs };
.level::eviction_priority = { .L1::evict_normal, .L1::evict_unchanged,
                              .L1::evict_first, .L1::evict_last, .L1::no_allocate };
.level::cache_hint = { .L2::cache_hint };
.level::prefetch_size = { .L2::64B, .L2::128B, .L2::256B };
.vec = { .v2, .v4 };
.type = { .b8, .b16, .b32, .b64, .b128,
          .u8, .u16, .u32, .u64,
          .s8, .s16, .s32, .s64,
          .f32, .f64 };

예를 들어 ld.global.nc를 사용하면 실제 SASS는 LDG.E.CONSTANT로 나온다. 여기에 L1 Cache 미할당, L2 prefetch 정책을 추가한

ptx
ld.global.nc.L1::no_allocate.L2::256B

를 쓰면 SASS는 LDG.E.NA.LTC256B.CONSTANT가 된다. DeepEP에서 사용하는 패턴이 바로 이 방식이다.

요약: 큰 프로그램에서 유연한 Cache 정책을 사용하면 Cache 활용률을 더 끌어올려 성능을 개선할 수 있다. 하지만 조합이 매우 다양하고 Memory Order와도 얽혀 있어, PTX 문서에 전부 상세히 정리되어 있지는 않다. 앞으로 더 세밀한 분석이 필요하다.

참고로, 고빈도(frequent) 트레이딩 쪽에서는 Cache 정책·Memory Model 관련 면접 질문이 거의 필수다. ns 단위로 시간을 다투는 프로그램을 작성해야 하니 당연한 일이다. DeepSeek이 이런 극단적인 최적화에 매달리는 것도 자연스러운 일이다.

5.3 Store 명령

Store 명령은 아래와 같다.

text
st{.weak}{.ss}{.cop}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};
st{.weak}{.ss}{.level::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};
st.volatile{.ss}{.vec}.type [a], b;
st.relaxed.scope{.ss}{.level::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};
st.release.scope{.ss}{.level::eviction_priority}{.level::cache_hint}{.vec}.type [a], b{, cache-policy};
st.mmio.relaxed.sys{.global}.type [a], b;

.ss = { .global, .local, .param{::func}, .shared{::cta, ::cluster} };
.level::eviction_priority = { .L1::evict_normal, .L1::evict_unchanged,
                              .L1::evict_first, .L1::evict_last, .L1::no_allocate };
.level::cache_hint = { .L2::cache_hint };
.cop = { .wb, .cg, .cs, .wt };
.sem = { .relaxed, .release };
.scope = { .cta, .cluster, .gpu, .sys };
.vec = { .v2, .v4 };
.type = { .b8, .b16, .b32, .b64, .b128,
          .u8, .u16, .u32, .u64,
          .s8, .s16, .s32, .s64,
          .f32, .f64 };

많은 부분은 ld와 동일하다. 차이는 cache operation policy(COP)에 있다. write‑back vs write‑through, L1/L2에 cache할지 여부, evict‑first 정책 사용 여부 등을 정의한다.

이런 옵션들은 L1/L2 Cache 점유를 최적화하는 데도 큰 도움이 된다.

마지막으로 레지스터 압력과 수명을 분석하는 방법을 하나 더 소개한다. 먼저 cuobjdump로 cubin을 추출한다.

bash
$ cuobjdump a.out -xelf all
Extracting ELF file 1: a.1.sm_86.cubin
Extracting ELF file 2: a.2.sm_86.cubin

그다음 nvdisasm로 분석한다.

bash
$ nvdisasm -plr ./a.2.sm_86.cubin

---

6. Scale‑UP/Scale‑Out 네트워크와 메모리 모델

컴파일/연산자 분해·스케줄링 관점에서 보면, Tile 기반 IR이 훨씬 다루기 쉬워진다. Cutlass의 layout 대수 추상화, Cutlass Distributed GEMM 관련 작업, 그리고 생태계 내 Triton 등도 있다. 특히 최근 ByteDance의 Triton‑Distributed[8]는 매우 뛰어난 작업이다.

응용 관점에서 보면, 향후 Scale‑UP/Scale‑Out에서 최소 통신 단위는 Tile 단위이거나 Token 단위가 될 것이다. 따라서 메시지 크기는 보통 2KB 이상이 된다.

한편 GTC25 세션이 제시한 로딩 최적화 권장 사항을 다시 보자.

칩 간 Scale‑UP/Scale‑Out 네트워크에서, 정말로 작은 size 접근을 위해 최적화를 해야 할까?

질문을 바꿔 보자. 만약 모든 메시지가 64B/128B 같은 작은 크기라면, 초대형 네트워크를 구성할 때 라우팅/CRC 등의 헤더에 얼마나 많은 비트가 필요할까? 특수 토폴로지를 지원하려면 소스 라우팅 헤더 등도 넣어야 한다. 그러면 작은 메시지의 실제 네트워크 전송 효율은 상당히 떨어진다.

UALink 프로토콜은 DataLink Flit 길이를 640B, Transaction Layer Flit을 64B로 정의한다.

이 설계는 GPU 쪽에서 온칩 네트워크를 달기에는 꽤 단순하지만, UALink Switch에서 매우 높은 처리량을 뽑아내기는 쉽지 않다.

각 Switch는 DL Flit을 풀어서, TL Flit 하나씩 처리해야 한다. 스위치의 룩업·포워딩 압력이 상당하다. 동시에 이 구조는 스위치 설계에도 제약을 준다. Shared Buffer Switch를 쓰면 TM의 MMU 설계를 51.2T/102.4T급 라인레이트의 PPS까지 감당하도록 만드는 것은 매우 어렵다. 결국 Port‑based Buffer 설계로 갈 수밖에 없고, 각 UALink 포트마다 작은 Tile 기반 Port Logic을 넣어야 한다. 하지만 이 경우 혼잡 제어가 또 하나의 난제로 떠오른다. UALink는 Credit‑based 방식을 채택했다.

이런 구조로도 수천 개 GPU를 연결하는 시스템을 만들 수는 있다. 단기적으로는 대역폭도 꽤 높게 뽑을 수 있을 것이다. 하지만 장기적인 진화(510년)를 생각하면 문제가 좀 있어 보인다. 또 다른 문제는 NVLink/UALink 같은 "대형기" 시스템이 510년 뒤에도 살아 있을지 자체가 의문이라는 점이다.

전반적으로 NVLink는 극단적 상황에서의 전송 효율이 UAL보다 낮을 수 있지만, 프로토콜 구조는 더 깔끔하고 스위치 확장이 수월하다.

사실 이더넷 테이프 한 장이면 충분하다.

앞에서 Little's Law로 메모리 접근과 inflight를 모델링한 장을 보면, 마치 항상 풀로드 상태로 시스템을 돌릴 수 있을 것처럼 느껴질 수도 있다. 하지만 실제 workload 변동을 고려하지 않은 모델이다. 그래서 필자는 시스템 scale 모델링에서는 반드시 Little's Law와 Kingman 공식을 함께 써야 한다고 강조해 왔다.

Kingman 공식의 관점에서 보면, 이용률이 100%에 가까워질 때 지연은 아래 곡선처럼 폭발적으로 증가한다.

그러나 실제로 지연을 측정·평가할 때는 대부분 무부하(빈 상태)에서만 본다. 물론 현업에서는 NCCL kernel launch 지연 같은 부분은 이미 문제로 보고 있다. 하지만 대부분은 cugraph로 kernel launch 오버헤드를 줄이는 측면만 보고 있다.

이 글에서는 당시 아키텍처에서 지연을 줄이는 방법을 논하면서, 이미 IBGDA를 제안한 바 있다.

필자가 Kingman 공식을 반복해서 강조하는 이유는, 많은 사람들이 RoCE에서 DeepEP benchmark만 보고 좋아하지만, 최종 E2E 성능 향상이 기대만큼 나오지 않는 근본 이유를 이해하지 못하기 때문이다.

네트워크 변동계수(CV) 관점에서 보면, DeepSeek이 왜 Adaptive Routing(AR)을 켜는지 이해할 수 있다. 연산 서비스의 변동계수 관점에서 보면, 왜 Low‑Latency Kernel에 hook을 걸고 GroupGEMM을 쓰는지, 왜 EPLB가 필요한지 알 수 있다. 최대한 부하를 고르게 만들어 GEMM 계산 지연의 jitter를 줄이기 위한 것이다.

시스템이 scale‑out 가능한지 여부는, 스위치 radix가 얼마인지, 어떤 토폴로지에서 이론상 몇 카드까지 묶을 수 있는지가 아니라, 높은 대역폭과 낮은 지연을 동시에 요구하는 상황에서, 시스템이 거의 만재 상태일 때 jitter를 어떻게 제어하느냐에 달려 있다.

이 관점에서 보면 PFC가 "at scale"에서 먹힌다는 말은 완전히 허구다. DCQCN은 복잡한 모델을 만들었지만 가장 기본적인 부분을 이해하지 못했다. 그래서 NV도 결국 DCQCN을 버렸다.

칩 간 네트워크에서 2KB~4KB 정도 데이터를 매번 전송하고, 비동기 접근으로 inflight가 늘어나는 상황에서는, 지연의 절대값보다 jitter 제어가 훨씬 더 중요해진다. 이 점을 잘 모르는 사람이 많다. 필자는 Cisco에서 10년 넘게 일하며, 온칩 네트워크 혼잡부터 데이터센터, WAN까지 수많은 문제를 다뤄 봤기 때문에, 지연과 jitter를 반복해서 강조할 수밖에 없다. eRDMA의 혼잡 제어 알고리즘을 설계할 때도 jitter를 1순위 최적화 목표로 두었다.

6.1 Scale‑UP/Scale‑Out 메모리 모델 설계

앞에서 메모리 모델을 한참 다뤘다. 그렇다면 Scale‑UP과 Scale‑Out의 메모리 모델은 어떻게 설계해야 할까?

동기 LD/ST, Cache 일관성 같은 순차 일관성 특성은 inflight bytes를 제약해 대역폭 활용을 떨어뜨린다. 한편 초대형 네트워크는 여러 계층 스위치를 거쳐 구성되므로, 다중 경로로 인한 Data Race 문제가 생긴다. 여러 패킷을 전송할 때, 네트워크가 손실·재전송·재전송으로 인한 순서 뒤바뀜을 허용할 것인지도 고민해야 한다.

현재 산업계 해법은 극단적이다.

• 한쪽 끝에는 Credit‑Based Flow Control로 zero‑loss를 목표로 하며 매우 낮은 BER(예: IB 1e‑15, Eth 1e‑12)을 요구하는 방식이 있다. Ethernet 진영은 최근 Low Latency Retrans(LLR)을 또 들고 나왔다.
• IBTA의 Reliable Connection은 STRONG ORDER를 요구해 lossless + GoBack‑N 재전송, 혹은 수신 측 reorder buffer를 요구한다.

다른 극단은 AWS SRD처럼 전송 순서를 전혀 보장하지 않고, 나머지 모든 복잡도를 소프트웨어에 떠넘기는 방식이다. 이 경우 통신 kernel에 엄청난 명령어 오버헤드가 생겨 연산 자원이 낭비된다.

또한 RDMA 방식을 Scale‑UP에 쓰지 말아야 하는 이유는 《HotChip2024 후기: 가속기 인터커넥트 및 Scale‑UP에 왜 RDMA를 쓰면 안 되는가》에서 이미 설명했다. 원래 한 번의 LD/ST로 끝날 일을 WQE를 만들고 온갖 일을 하는 복잡한 프로토콜로 바꾸기 때문이다. IBGDA는 현재 아키텍처 제약 하에서의 타협일 뿐이다. 그래서 DeepSeek 논문에서도 Unified Scale‑UP/Scale‑Out semantics가 필요하다고 지적한다.

NV 자체도 Async Proxy 하에서는 같은 주소(Same‑Address)에 대해 ordering을 보장하지 않는다.

즉 계산 측에서도 엄격한 ordering을 크게 요구하지 않는다. 전송 측에서 굳이 스스로를 옥죌 필요가 없다. 적절한 weak order만 잘 처리하면 된다. 이 지점이 바로 아키텍처 상의 trade‑off다.

필자가 지난 몇 년간 algebra 관점에서 commutative idempotent semigroup, semi‑lattice 등의 개념을 계속 언급한 이유도 여기에 있다. 예를 들어 3년 전 글 《向上，点亮未来：DPU的若干代数问题》를 참고하면 좋다.

Commutative idempotent semigroup + 부분 순서 집합(partially ordered set)에 통신 다중 경로, out‑of‑order를 얹으면, 선형 순서를 강제하는 건 오히려 시스템에 부담이 된다. 그렇다면 관점을 바꿔 메모리 쪽에서 해결할 수는 없을까?

앞에서의 모든 논의는 사실 이 질문을 던지기 위한 준비였다.

메모리 분포를 보자. 주소를 축으로 한 부분 순서 집합(Partially Ordered Set)이라 볼 수 있다. 이 메모리에 대한 연산이 교환 가능(Commutative), 멱등(idempotent)이고, 반군(Semigroup) 정의에 따라 폐합성과 결합법칙(associativity)을 만족한다면, 이 메모리 연산은 semi‑lattice 구조를 이룬다.

단순한 메모리 read/write는 멱등성을 만족한다. 결합법칙은 "연산의 단위(원자)가 무엇인가"에 따라 달라진다. 메시지 단위인가, byte 단위인가? 메모리에서 write와 read 연산의 주소 공간이 충돌하면 결합법칙이 깨진다. 하지만 메시지 수준의 semantic으로 둘을 구분해 두면 이 문제를 피할 수 있다. 그래서 분산 병렬 프로그래밍에서 Actor 모델과 CSP 모델이 자주 등장하는 것이다.

요약하면, 메모리 사용 단위를 메시지 semantic에 맞게 정의하고, 메모리 연산의 주소와 명령을 메시지와 묶으면, semi‑lattice 대수 구조를 구현할 수 있다. 그 결과, 대규모 통신의 난제를 해결할 수 있다.

• 교환법칙(Commutative): out‑of‑order 전송, 다중 경로로 혼잡을 해소해도 된다.
• 멱등성(idempotent): 패킷 손실 후 재전송을 언제든 할 수 있다.
• 결합법칙(associative): 여러 연산을 합쳐(merge) 먼 곳으로 보낼 수 있다.

하드웨어에서 semi‑lattice 구현은 의외로 비용이 적다.

• 교환법칙: 일부 relaxed order를 허용하면 된다.
• 결합법칙: TMA가 대표적인 예다. tensor_map descriptor로 여러 연산을 묶어 한번에 수행한다.
• 멱등성: 각 메시지에 Seq 필드를 두고, consumer 측에서 처리된 ci 포인터에 fence를 걸어 중복 write만 막으면 된다. 인‑네트워크 연산(in‑network computing)에서는 덧셈 멱등성 등 추가 고려가 필요하지만 원리는 비슷하다.

결국 프로그래밍 인터페이스는 아래처럼 단순해진다. GPU는 한 번의 명령을 issue하고 mbarrier의 completion_tx 카운터만 기다리면 된다.

cpp
cde::cp_async_bulk_tensor_2d_global_to_shared(tile_smem,
                                              &tensor_map, x, y, bar);
token = cuda::device::barrier_arrive_tx(bar, 1, sizeof(tile_smem));

// 다른 일 수행...

// 모든 데이터 도착 대기
bar.wait(std::move(token));

전송 계층의 요구 사항은 이렇게 바뀐다.

TMA 명령에서 발생하는 세밀한 LD/ST들을 하나의 메시지로 묶을 수 있어야 한다. 메시지 내부에서는 데이터가 순서 없이(commit order 무관하게) 도착해도 되고, 다중 경로 전송·손실·재전송도 허용된다. 메시지가 완료되면 mbarrier를 업데이트하면 된다.

축하한다. 20년 전 iWARP가 이미 만들어 둔 Direct Data Placement(DDP)를 다시 발명했다.

NV(Mellanox)는 이 개념을 완전히 이해했다고 보기 어렵다. AR에서 일부 DDP를 구현했지만, SEND/RECV에 대한 DDP는 지원하지 못한다. RoCEv2 프로토콜의 한계 때문이다. RoCEv2는 메시지 전송 방식을 First/Middle/Last flag 하나만으로 정의한다. 중간 패킷에는 메모리 연산 주소가 담기지 않으므로, 순서를 반드시 보장해야 한다.

후에 메시지를 여러 조각으로 나누고 각 조각마다 RETH를 실어 주소를 싣는 방식으로 발전하긴 했지만, SEND/RECV의 경우 수신 버퍼의 절대 주소를 정확히 알 수 없다. 또한 NIC에 포트가 두 개 있을 때, RoCEv2는 XOR 방식으로 한 포트만 골라 보내게 되어 있다. 단일 QP를 두 포트에 걸쳐 load balance하려면 어떻게 할 것인가?

반면 iWARP의 DDP는 아주 명확하게 정의되어 있다. Msg Seq Number(MSN)와 Msg offset(MO) 필드를 사용해, 상대 주소(relative offset)를 표현한다. 이렇게 하면 수신 측 reorder buffer가 필요 없다. 패킷이 도착하는 대로 offset을 기준으로 메모리에 직접 쓰면 된다. 그리고 작은 메시지 bitmap로 수신 완료 여부를 관리하고, 다 받으면 mbarrier를 업데이트하면 된다. 멱등성(중복 write 방지)은 이미 mbarrier에 반영된 MSN에 대해 재전송 도착 시 쓰기를 막는 방식으로 쉽게 보장할 수 있다.

이것이 eRDMA의 구현 방식이다. 매우 단순하고 하드웨어 비용도 작다. reorder buffer 덕분에 추가 지연·jitter가 생기지도 않는다. 동시에 fabric의 다중 경로를 충분히 활용하고, 일부 링크가 고장 나도 견딜 수 있다.

이제 eRDMA(Weak Order), NV(Mellanox)의 StrongOrder, AWS‑SRD의 RelaxOrder를 나란히 놓고 보면, 무엇이 best practice인지 자연스럽게 판단할 수 있을 것이다.

---

참고 자료

[1] CUDA Techniques to Maximize Memory Bandwidth and Hide Latency:
https://register.nvidia.com/flow/nvidia/gtcs25/vap/page/vsessioncatalog/session/1727709012449001X6PZ

[2] Advanced Performance Optimization in CUDA:
https://www.nvidia.com/en-us/on-demand/session/gtc24-s62192/

[3] Sequential Consistency and TSO:
https://www.cis.upenn.edu/~devietti/classes/cis601-spring2016/sc_tso.pdf

[4] SPCL_Memory Model:
https://spcl.inf.ethz.ch/Teaching/2019-dphpc/lectures/lecture4-memory-models.pdf

[5] CUDA C Programming Guide:
https://docs.nvidia.com/cuda/cuda-c-programming-guide/

[6] Cluster Group:
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cluster-group-cg

[7] CUDA Refresher: The CUDA Programming Model:
https://developer.nvidia.com/blog/cuda-refresher-cuda-programming-model/

[8] Triton‑distributed:
https://github.com/ByteDance-Seed/Triton-distributed

Embed anything (PDFs, Google Docs, Google Maps, Spotify…)