NVIDIA GPU를 중심으로 메모리 일관성 모델과 캐시, Async/Proxy 메모리 경로, Thread Block Cluster 및 DSMEM, TMA/비동기 복사, 그리고 Scale‑UP/Scale‑Out 네트워크 설계를 함께 분석한다.
사실 많은 사람들이 Scale‑UP/Scale‑Out 버스를 이야기할 때, 대부분 네트워크만 이야기하고 GPU Memory Model 관점은 잘 보지 않는다. 한편으로 황 CEO가 이야기한 “먼저 Scale‑UP, 그 다음 Scale‑Out”도, 사실은 세일즈용 화법이라고 볼 수 있다. 반대로 묻자. NV는 올해 GTC에서 Scale‑Out 쪽으로 내놓을 만한 게 뭐가 있었나? 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 based IR 등 일련의 요소를 합쳐 종합 분석을 해본다. 때마침 화웨이 UB, UALink 1.0 표준도 공개되었으니, 마지막에는 Scale‑Out·Scale‑UP에서 요구되는 메모리 모델을 이야기한다. 덤으로 eRDMA가 다중 경로(multipath)에서 메모리 모델을 어떻게 구현하는지, 그리고 표준 RC와 AWS SRD와의 비교도 다룬다.
또한 GTC25 세션에는 가치 있는 두 가지 주제가 있다. 저지연 Cluster 동기화, 그리고 메모리 대역폭 최대화 방법이다. 이 글에서 함께 소개한다. 구조는 다음과 같다.
일관성(consistency)의 근원은 폰 노이만 구조에서 “어떤 read도 가장 최근의 write 결과를 돌려주어야 한다”는 가정에서 온다. 하지만 분산 시스템이나 멀티코어 CPU 시스템에서는, 연산 지연 등으로 인해 결과가 예측 불가능해진다.
UPenn의 Sequential Consistency and TSO 튜토리얼[3]이 설명을 잘해 놓았다. SPCL의 Memory Model[4] 슬라이드도 좋다. 더 자세히 알고 싶다면 아래 책을 참고하면 된다.
우선 두 개의 코어가 있는 프로세서를 보자. 하나는 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를 유지하는 실질적 방법은 두 가지다. 단일 코어로만 실행하거나, 메모리 접근 시 순서 보장을 강하게 걸어 순차 접근을 강제하는 것.
즉, 매 타임스텝마다 스위치는 실행할 스레드를 선택하고, 그 스레드의 “다음 이벤트”를 완전히 실행한다. 이 모델은 순차 일관성 규칙을 보존하지만, 치명적인 단점이 있다. 너무 느리다는 점이다. 한 번에 하나의 명령만 실행하므로, 멀티스레딩으로 얻을 수 있는 병렬 실행 이점을 대부분 잃는다.
더 나쁜 점은, 각 명령이 완료될 때까지 기다렸다가 다음 명령을 시작해야 한다는 것이다. 현재 명령의 효과가 다른 스레드에 보이기 전에는 추가 명령을 실행할 수 없다.
단일 프로세서 관점에서 보면, 메모리 write가 끝날 때까지 직접 기다리면 Store가 너무 느려진다. 그래서 보통 Store buffer를 두어 지연을 숨기고 stall을 피한다. 멀티코어에서는 각 코어가 독립된 Store buffer를 갖는다.
하지만 이 경우 위 그림처럼, 두 코어 모두 오래된 값을 읽을 가능성이 생긴다.
하지만 이런 트레이드오프가 가져오는 성능 이득은 크다. 이것이 Total Store Order(TSO)가 등장한 이유다. 형식적 정의로는 Store→Load 보장(Store→Load ordering)을 포기하고, Store Buffer 설계를 허용한 모델이다.
Store→Load 문제는 FENCE로 해결할 수 있다. FENCE 구현도 간단하다. 예를 들면, store buffer를 비워(main memory에 flush) Read‑Write coherence를 보장하는 식이다.
더 나아가, 우리는 더 많은 Reorder를 허용해 실행 병렬성을 높일 수 있을까? 그리고 Fence(Memory Barrier)로 프로그램 순서를 논리적으로 다시 맞출 수 있을까?
산업계에서는 Relaxed Consistency에 대해 여러 정의가 있다. Total Store Order는 앞서 말했듯 하나의 제약을 포기했고, Partial Store Order(PSO)는 추가로 제약을 더 포기한다. 일부 Relaxed Memory Order(RMO)는 네 가지 제약 모두를 완전히 포기한다. 사실 많은 프로세서들이 이런 모델을 지원하고, GPU 자체도 Relaxed Consistency 시스템이다.
여기서 많은 사람들이 헷갈려 한다. Cache 일관성은 주로 Store를 다른 프로세서에 “어떻게” 전달할지에 관한 메커니즘이다. 즉, 쓰기(write)가 필요한 시점에 다른 프로세서에서 보이게 하는 메커니즘이다. 반면 Memory Model은 연산이 다른 프로세서에 “어떤 순서”로 전달되는지에 대한 경계를 정의한다.
NVIDIA GPU의 메모리 계층은 아래와 같다. 수천 개의 CUDA 코어에 대해 메모리 TSO(Total Store Order)를 보장하는 비용은 매우 크다. 그래서 NVIDIA GPU는 Partial Store Order 메모리 모델을 채택한다.
이 문제에 대해 아키텍처마다 조금씩 차이가 있다. 예를 들어, 인텔이 스스로 정리한 네 가지 메모리 모델 표현 방식 같은 것들.
단일 스레드 관점에서는 동일 주소에 대한 LD/ST는 순서를 보장한다. 예를 들어 다음과 같다.
다만 예외가 하나 있다. 아래 프로그램을 보고 출력 결과를 맞춰 보자. 사실 이 코드는 정의되지 않은 동작을 만든다.
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)); // 새 값 쓰기 *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가 있고, 상수(constant)가 그 공간에 배치되기 때문이다. 이 Read‑Only Cache는 L2 Cache와 데이터 경로가 분리되어 있어 문제가 생긴다.
이런 값을 수정하면 정의되지 않은 동작이 발생한다.
앞서 본 것처럼 순차 일관성(SC)을 유지하기 위해서는 다음 네 가지 규칙이 필요하다.
Total Store Order(TSO)는 Store buffer를 도입하기 위해 네 번째 규칙(Store→Load)을 포기한다. 단일 코어 내부에서는 Load가 Write buffer의 값을 bypass해서 읽을 수 있고, 코어 간에는 fence로 순서를 맞춘다.
하지만 GPU 내부에는 수많은 CUDA 코어가 존재한다. 수천 개 코어의 모든 메모리 연산 순서를 강하게 보장하면, 명령 및 데이터 병렬성에 막대한 성능 저하가 온다. GPU에서 TSO를 유지하는 비용이 너무 크므로, 실제로는 Relaxed Order를 지원하고 ATOMIC, FENCE로 필요한 부분만 순서를 맞추는 것이 더 적절하다.
NVIDIA GPU는 네 가지 메모리 오더를 지원한다.
아래 그림처럼, 순차 일관성은 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에서는 다음을 볼 수 있다.
ptxfence.sc.sys
이 fence는 Prior load/store가 atomic 이후로 넘어가는 것을 막는다. 동시에 atomic load는
ptxld.acquire
형태로 발행되어, 이후 LD/ST가 이 acquire 이전으로 당겨지는 것을 막는다.
ptxld.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]
위 코드에서
cppval = a.load(cuda::std::memory_order_seq_cst);
를
cppval = a.load(cuda::std::memory_order_acquire);
로 바꾸면, PTX를 다시 봤을 때 fence.sc.sys가 제거된 것을 볼 수 있다.
ptxld.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는 여전히 acquire에 의해 막힌다.
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가 앞으로 당겨지는 것은 허용한다.
ptxld.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
마지막은 가장 느슨한 Relaxed 메모리 모델로, 앞뒤 LD/ST 모두 freely reorder될 수 있다.
코드는 다음과 같다.
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:
ptxld.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
앞서 본 명령들은 모두 .sys 속성을 갖고 있다. 실제로는 필요에 따라 다양한 범위(scope)를 선택할 수 있다. CUDA C++에서는 다음과 같은 scope를 정의한다.
PTX에서 scope는 아래와 같이 표현된다.
NVIDIA GPU 메모리 계층 구조를 떠올려 보자.
간단한 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로 나온다.
ptxConsumer: 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는 다른 데이터 경로를 타고, block scope 밖이라 서로 막지 않는다.
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:
ptxConsumer: 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 슬라이드에서처럼 “잘” 동작하지는 않는다. 왜일까?
Scope가 GPU 전체라고 해도, 슬라이드 작성자가 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가 수반되므로 느리지만 여러 스레드가 다수의 값을 교환할 때 유용하다.
Ampere 세대부터 Async Thread 프로그래밍 능력이 도입되었다. 핵심은 GMEM→SMEM(Load) 또는 SMEM→GMEM(Store) 방향의 비동기 복사다. 이는 레지스터·L1 사용을 피하면서 대용량 데이터 처리의 throughput을 높인다.
우선 비동기 prefetch를 수행하고,
그다음 계산을 진행한다.
대략적인 계산 흐름은 다음과 같다. 자세한 코드는 다음을 참조.
https://github.com/zartbot/tensorcore_gemm/blob/main/05_pipeline_gmem_to_smem.cu
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; // TensorCore 계산
}
}
}
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 완료를 기다릴 수 있다.
구체적인 흐름은 다음과 같다. 우선 SMEM에 mbarrier를 하나 할당한 뒤, 한 스레드가 TMA 명령(UBLKCP)을 발행한다. 이때 SMEM에 복사되는 데이터는 16바이트 정렬(alignas(16))이어야 한다.
완료 여부는 completion_tx 카운터로 관리한다.
TMA‑2D 예시는 아래와 같다.
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에 쓴 데이터 완료 대기 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; } // GPU 6에 메모리 할당 및 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); // GPU 7에서 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를 지원한다.
Hopper의 Cluster 프로그래밍에 대한 공개 자료는 많지 않다. cuda c programming guide[5]에 약간의 설명이 있다. Hopper에서 새로 도입된 계층 구조로, 로컬 SM‑to‑SM 데이터 경로를 만들고 Distributed Shared Memory(DSMEM) 개념을 제공한다.
소프트웨어 측에서는 Grid와 Block 사이에 Cluster 계층이 하나 더 생긴 모양새다.
아래는 간단한 예시 코드이다. Kernel 정의에
cpp__cluster_dims__(x, y, z)
를 붙여 Cluster 모양을 정하고, cg::this_cluster() 함수로 현재 cluster descriptor를 얻을 수 있다. 이식성을 고려해 단일 Cluster에는 최대 8개의 Thread Block만 지원하도록 되어 있지만, Hopper 데이터시트를 보면 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 dimension을 결정할 수도 있다.
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] 절에 정리되어 있다.
Cluster의 가장 중요한 활용은 Distributed Shared Memory이다. Cluster 내부 SM 간의 shared memory 접근을 L2를 우회(bypass)해 저지연으로 수행하며, 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 범위 동기화: 모든 스레드가 SMEM 선언 및 초기화 완료 보장 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; }
실행 결과(예시):
threadIdx 0 blockIdx 6 clusterRank 6 smem:60700
...
threadIdx 3 blockIdx 5 clusterRank 5 smem:50403
Thread Block Cluster의 장점은 SM‑to‑SM 네트워크로 데이터를 직접 교환해, L2/GMEM에 쓰는 비용을 피할 수 있다는 것이다.
그러나 cluster::sync() 구현에는 병목이 있다. 이는 cluster 전체를 동기화하고, cluster 내 다른 thread에서 LD/ST 데이터를 볼 수 있게 한다. 이로 인해 데이터가 L2를 거쳐가며 비용이 커진다. PTX에서는 cluster::sync()가 아래 두 명령으로 분해된다.
ptxbarrier.cluster.arrive; barrier.cluster.wait;
하지만 PTX 수준에서 arrive와 wait를 분리하고, release/relaxed를 선택적으로 써서 LD/ST visibility를 조절할 수 있다.
예를 들어 barrier 초기화 시 cluster::sync()를 써도 되지만, L2를 통과하기 때문에 느리다.
대신 release‑acquire 패턴을 이용할 수 있다.
Cluster 내부 SM‑to‑SM 통신에는 비동기 store와 local mbarrier를 사용해 저지연 동기화를 구현할 수 있다.
전체 코드는 다음과 같다.
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 visible 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); }
여기서 async store에는 cluster‑scope의 relaxed를 사용하고, local barrier wait에는 acquire를 쓴다.
H20에서 실제 테스트해 보면 mbarrier_arrive_expect_tx(sem_release) 대비 46% 정도 빨랐다.
low latency kernel elapsed 2.068736
async kernel elapsed 3.714880
TMA에는 멀티캐스트 기능도 추가되어, 한 번에 여러 block으로 데이터를 로드할 수 있다. 아래는 예시 코드이다. 컴파일 시 아래와 같은 경고가 뜬다. multicast::cluster는 sm_90a / sm_100a / sm_101a 아키텍처에서 제대로 성능이 나오고, sm_90에서의 성능은 떨어질 수 있다는 의미다.
ptxas ... warning : Advisory: '.multicast::cluster' modifier on instruction
'cp.async.bulk{.tensor}' should be used on .target 'sm_90a/sm_100a/sm_101a'
...
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() { 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_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; }
실행 결과 예시는 다음과 같다(멀티캐스트가 선택된 block에만 데이터가 들어감).
cluster 6 smem[0000]
cluster 7 smem[48 49 50 51]
cluster 0 smem[48 49 50 51]
...
이제 GTC25 세션의 또 다른 주제인 "Maximizing Memory Bandwidth"를 보자.
NVIDIA GPU에는 매우 많은 CUDA 코어가 있으며, 이에 따라 메모리 역시 깊은 계층 구조를 가진다.
코어에 가까운 cache일수록 접근 지연을 크게 줄일 수 있다. Shared Memory는 세대별로 용량이 늘어왔지만, 반도체 면적 상 한계에 가까워졌다. Hopper와 Blackwell에서 SMEM 용량이 늘지 않은 것도 같은 이유다. 대신 소규모 on‑chip 네트워크를 구성하여, SMEM들을 묶어 Distributed SMEM을 만들고 데이터 locality를 향상시켜 GMEM 접근 지연을 줄이는 방향을 택했다.
그래서 3장에서 본 것처럼, Hopper와 Blackwell에서 DSMEM과 Cluster 추상화를 잘 사용하는 것이 점점 더 중요해진다.
Blackwell은 또한 Cluster 혼재 배치(mixed placement) 및 스케줄링을 최적화해 활용도를 높인다.
그러나 on‑chip 면적 제약으로 SM 수 증가는 둔화되었다. 더 큰 SMEM, 더 큰 연산 성능, 더 많은 SM을 한꺼번에 요구하는 것은 근본적으로 상충한다. 결과적으로 FP64를 잘라내야 하는 상황까지 온다…
한편 HBM3e/HBM4 등으로 메모리 용량과 대역폭은 계속 증가하지만, compute와 memory access 간 튜닝 난도 역시 커지고 있다.
분산 시스템의 효율을 최적화할 때, 미시적으론 메모리 접근 지연, 거시적으론 부하 균형과 네트워크 혼잡 제어 등을 고려해야 한다. 이를 분석하는 가장 올바른 시각 중 하나가 대기행렬 이론이다. Kingman 공식 등으로 잘 알려져 있다. 여기서는 그보다 간단한 모델 하나를 보자. 안정된 시스템에서, 장기 평균 고객 수 L은, 도착률 λ와 평균 체류 시간 W의 곱이다. 이는 매우 직관적인 Little의 법칙이다.
예를 들어 에스컬레이터 위에 평균 2초마다 사람 1명(또는 1초에 0.5명)이 도착하고, 에스컬레이터를 타는 데 40초가 걸리면, 시스템이 동시에 감당할 수 있는 사람 수 L은 20명이다.
메모리 접근도 비슷하다. 메모리 대역폭과 평균 접근 지연으로 in‑flight bytes를 계산할 수 있다. Hopper에선 약 32KB의 in‑flight가 있어야 메모리 대역폭을 꽉 채울 수 있고, Blackwell은 거의 두 배인 64KB 정도가 필요하다.
단순한 kernel 하나를 생각해 보자. 한 스레드가 메모리를 몇 번 access하는지, 한 번에 몇 바이트를 읽는지, block 내 스레드 수, 하나의 SM에 올라가는 block 수를 알면, SM당 in‑flight bytes를 계산할 수 있다.
in‑flight bytes를 늘리는 방법은 보통 세 가지다.
예를 들어 loop unrolling으로 명령어 병렬성을 키우는 방법이 있다.
한편 vector load를 사용해 데이터 병렬성을 키울 수 있다.
하지만 ILP와 DLP를 늘리면 레지스터 압박이 크게 늘어난다.
그래서 레지스터 사용을 피하기 위해 비동기 복사가 등장했다.
동시에 비동기 로드는 데이터 복사와 계산을 overlap시켜준다.
Producer‑Consumer 패턴에 warp specialization을 결합하는 방식도 가능하다.
예를 들어 일부 thread를 Producer로 두어 메모리 복사에만 전념하게 한다.
Consumer는 계산을 수행하면서, 동시에 일부 thread로 미리 데이터를 prefetch한다.
GTC25 세션은 마지막으로 로드 최적화 가이드를 제시한다.
ILP·DLP와 대량의 비동기 복사가 동시에 사용되면, 레지스터와 cache 압박을 줄이는 것이 필수 과제가 된다. Ampere에서 도입된 Async Copy는 레지스터/L1을 우회해 GMEM에서 SMEM으로 바로 데이터를 가져온다. Hopper에서는 이를 확장해 nD tensor에 대한 issue 수를 줄이는 TMA를 도입했다. 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 명령의 다양한 사용법이 정리되어 있다.
ptxld{.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로 나타난다.
cop는 캐시 정책을 정의하는 옵션으로, 성능 튜닝에서 매우 중요하다.
.ca: 모든 캐시 계층에 저장(L1/L2 모두 cache). 기본 동작..cg: L2에만 cache, L1에는 cache하지 않음..cs: 데이터가 한 번만 사용될 때 유리. L1/L2에서 evict‑first 정책 사용(SASS에 EF 플래그). 예를 들어 reduction 등에 사용할 수 있다..lu: last use. spilled register 복구나 함수 스택 프레임 pop에 사용해 불필요한 write를 줄인다. global address에 사용하면 .cs와 동일..cv: cache하지 않음.재미있는 문제 하나.
ld.weak.global.cv와 ld.volatile.global의 차이는?LDG.E.STRONG.SYS.ld.weak.global.cg는?LDG.E.STRONG.GPU.또한 L1 캐시의 eviction 정책, allocate 여부 등을 정의할 수 있다.
L2 캐시에 대해서는 prefetch size도 설정할 수 있다. scope는 앞에서 설명한 메모리 오더의 scope와 같다.
SM 내부에는 L1 외에도 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를 선택적으로 사용할 수 있다. texture cache 용량이 크고, 지연은 충분한 병렬성으로 숨길 수 있을 때 유용하다.
ptxld.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 }; // cache operation .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 no‑allocate, L2 prefetch 256B를 더한 ld.global.nc.L1::no_allocate.L2::256B는 SASS에서 LDG.E.NA.LTC256B.CONSTANT로 나온다. DeepEP에서 사용한 바로 그 패턴이다.
요약: 큰 프로그램에서는 다양한 cache 전략을 유연하게 적용해 cache 활용도를 높이고 전체 성능을 올릴 수 있다. 다만 조합이 매우 많고, 메모리 오더와도 얽혀 있다. PTX 문서에 모든 조합이 상세히 기술되어 있지는 않으므로, 앞으로 더 세밀한 분석이 필요하다.
참고로 고빈도 트레이딩(high‑frequency trading) 쪽에서는 cache 전략과 memory model이 면접 단골 주제다. 프로그램이 ns 단위로 시간을 다툴 때 필수 스킬이다. DeepSeek 쪽 사람들이 이런 extreme optimization을 파고드는 것도 자연스럽다.
Store 명령은 다음과 같다.
ptxst{.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/write‑through, L1/L2 cache 여부, evict‑first 적용 여부 등을 제어한다.
이 역시 L1/L2 점유 최적화에 유용하다.
추가로, 레지스터 압박과 라이프타임을 분석하는 방법을 하나 더 보자. 먼저 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
이러한 컴파일/연산자 분할 및 스케줄링 관점에서 보면, Tile based IR가 왜 중요한지가 더 잘 보인다. CUTLASS의 layout algebra 추상, CUTLASS Distributed GEMM 관련 작업, Triton 생태계, 그리고 최근 ByteDance의 Triton‑Distributed[8] 같은 훌륭한 작업도 모두 이 맥락에 있다.
응용 관점에서 보면, 미래의 Scale‑UP/Scale‑Out에서 최소 통신 단위는 tile 단위나 token 단위가 되어 갈 가능성이 높다. 즉, 메시지 size는 보통 2KB 이상이 될 것이다.
한편 GTC25 세션에서 제안한 로드 최적화 가이드를 다시 떠올려 보자.
칩 간 Scale‑UP/Scale‑Out 네트워크에서, 작은 size 접근을 위한 별도 최적화가 필요할까?
다른 식으로 묻자. 만약 모든 메시지가 64B/128B와 같이 작다면, 초대형 네트워크를 구성할 때 라우팅/CRC 등 오버헤드용 header는 얼마나 필요할까? 특수 토폴로지를 지원하면 source routing header 등의 추가 정보도 필요하다. 이 경우 작은 메시지의 실제 네트워크 효율은 크게 떨어질 수 있다.
UALink 프로토콜은 Data Link Flit 길이를 640B, Transaction Layer Flit 길이를 64B로 정의한다.
이런 설계는 GPU 측에서 칩 내부 네트워크를 그대로 매달아 쓰기엔 편하다. 하지만 UALink 스위치에서 고처리량을 구현하는 것은 쉽지 않다.
각 스위치는 DL Flit을 풀고, TL Flit을 하나씩 처리해야 한다. 이는 스위치의 lookup 및 전송 부담을 키운다. 또한 스위치 설계에 제약을 건다. Shared buffer switch를 택하면 TM MMU가 51.2T/102.4T line‑rate의 PPS를 처리해야 하는데, 이는 매우 어렵다. 따라서 실질적으로는 port‑based buffer 설계를 택해, 각 UALink 포트에 작은 tile‑based port logic을 만들어야 할 것이다. 그러나 이 경우 혼잡 제어가 또 다른 난제가 된다. UALink는 credit‑based 방식을 택했지만 말이다.
물론 이러한 방식으로 수천 개 GPU를 상호 연결하는 것도 가능하고, 단기적으로는 대역폭도 꽤 높게 만들 수 있다. 그러나 장기적 진화(5~10년 스케일)를 생각하면 여러 의문점이 생긴다. NVLink/UALink로 구축한 이런 “대형기” 시스템이 장기적으로 남을지 역시 논의할 만한 주제다.
전체적으로 보면 NVLink는 극단적인 경우 UAL보다 전송 효율이 낮을 수 있지만, 프로토콜이 더 깔끔하고 스위치 용량을 늘리기에도 수월하다.
사실 이 모든 것을 덮을 수 있는 것은 한 장의 이더넷 “테이프(duct tape)”일지도 모른다.
앞에서 GTC25 세션에서는 Little의 법칙으로 메모리 접근과 in‑flight bytes를 모델링했는데, 마치 항상 full‑load로 잘 돌아갈 듯하게 보인다. 그러나 실제 워크로드 변동을 고려하면, Little만으로는 부족하다. 필자는 시스템 scale을 모델링할 때 Little과 Kingman 공식을 함께 써야 한다고 줄곧 강조해 왔다.
Kingman 시각에서 보면, 자원 활용률이 100%에 가까워질수록 지연은 다음과 같은 곡선으로 폭발한다.
하지만 많은 경우, 우리가 지연을 측정·평가하는 상황은 공회전(idle)에 가까운 상황뿐이다. 물론 엔지니어들은 NCCL kernel launch 오버헤드 등 실제 일부 문제를 보고 있긴 하다.
위 글에서도 기존 아키텍처에서 지연을 줄이는 방법을 논하며, 당시부터 IBGDA를 제안했었다.
필자가 Kingman 공식을 반복해서 언급하는 이유는, 많은 이들이 RoCE 하에서 DeepEP benchmark는 좋게 나오는데 E2E 성능 향상은 기대에 못 미치는 이유를 잘 이해하지 못하기 때문이다. 네트워크 변동계수(coefficient of variation)의 관점에서 보면, DeepSeek가 왜 Adaptive Routing(AR)을 켜는지 알 수 있다. compute 서비스 변동계수 관점에서 보면, 왜 low‑latency kernel에 hook을 걸고 GroupGEMM과 EPLB를 하는지, 즉 GEMM 계산 지연의 jitter를 줄이기 위한 것임을 이해할 수 있다.
시스템이 scaling이 가능한지 여부는, 스위치 radix가 몇인지, 이론상 몇 장까지 토폴로지를 짤 수 있는지 같은 표면적 질문에 답하는 것으로 결정되지 않는다. 대역폭과 latency를 동시에 요구하는 상황에서, 시스템이 거의 full‑load 상태일 때 전체 jitter를 어떻게 제어할지가 latency를 낮추는 핵심이다. 이런 관점에서 보면, PFC를 내세워 “at scale”을 떠드는 건 사실상 헛소리에 가깝다. DCQCN처럼 복잡한 모델을 만들어 놓고도 기본적인 부분을 못 잡은 셈이다. … 그래서 NVIDIA도 결국 DCQCN을 버려 가는 것이다.
메시지 하나당 전송 size가 2KB~4KB 정도이고, 비동기 접근으로 in‑flight가 커진 상황에서는, 정적인 지연보다 jitter 제어가 훨씬 중요해진다. 많은 이들이 이 점을 잘 모른다. 필자가 Cisco에서 십수 년 동안 on‑chip 네트워크 혼잡부터 데이터센터, WAN까지 수많은 문제를 다루며 이 부분을 강조해 온 이유다. eRDMA의 혼잡 제어 알고리즘을 설계할 때도 이 jitter 최소화를 1차 목표로 두었다.
앞에서 메모리 모델을 길게 이야기했는데, 그렇다면 Scale‑UP/Scale‑Out 메모리 모델은 어떻게 설계해야 할까?
동기식 LD/ST와 cache 일관성 같은 strong sequential consistency를 칩 간에도 그대로 요구하면, in‑flight bytes가 줄고 대역폭 활용이 나빠진다. 또 초대형 네트워크는 보통 다단 스위치로 구성되므로, 여러 경로를 거치면서 Data Race 문제가 발생한다. 다수 패킷 전송 시, 네트워크가 packet loss 및 재전송을 허용하는지, 재전송으로 인한 out‑of‑order를 어떻게 처리할지도 고민해야 한다.
현업의 해법은 양극단이다.
다른 한 극단은 AWS SRD처럼 전송은 전혀 ordering을 보장하지 않고, 나머지 문제를 전부 소프트웨어에 떠넘기는 쪽이다. 그러나 이는 통신 kernel에 많은 명령 오버헤드를 유발해 compute 자원을 낭비한다. 또한 이런 명령 오버헤드 관점에서 보면, 이전 글 "HotChip2024 후기: 가속기 인터커넥트와 Scale‑UP에 RDMA를 쓸 수 없는 이유"에서도 왜 가속기 Scale‑UP에 기존 RDMA 방식을 쓰기 어렵다고 했는지 설명했다. 원래 LD/ST 한 번이면 될 일을 WQE 준비 등으로 부풀려 버리기 때문이다. IBGDA는 기존 아키텍처 제약 아래에서의 타협일 뿐이고, DeepSeek 논문에서 Unified Scale‑UP/Scale‑Out semantics 요구를 언급한 이유도 여기 있다.
NVIDIA 스스로도 Async Proxy 아래에서는 same‑address ordering조차 보장하지 않는다.
본질적으로 compute 측에서도 그렇게 엄격한 ordering을 요구하지 않는데, 전송 측에서 굳이 스스로 족쇄를 찰 필요가 없다. 적절한 weak order만 잘 설계하면 된다. 이런 아키텍처적 trade‑off 때문에, 필자가 몇 년 전부터 반격(algebraic) 관점에서 semilattice semantics를 주장해 온 것이다. 자세한 내용은 3년 전 글 "向上，点亮未来：DPU的若干代数问题"을 참고하라.
A commutative idempotent semi‑group, 그리고 부분 순서 집합(partially ordered set)…
많은 통신에서 multipath를 쓰려면 out‑of‑order 전송이 필요하다. 강한 ordering은 큰 부담이다. 그렇다면 시각을 바꿔, “메모리” 쪽에서 이 문제를 풀 수는 없을까?
앞서 모든 논의의 핵심은 이것이다.
메모리 배치는 메모리 주소를 기준으로 한 부분 순서 집합(Partially Ordered Set)이다. 메모리 연산이 교환법칙(Commutativity)과 멱등(idempotent), 그리고 반군(Semigroup)이 요구하는 폐쇄성과 결합법칙을 만족하면, 이 메모리 연산은 semilattice 구조가 된다. 단순한 메모리 read/write는 멱등성을 만족한다. 결합법칙은 이 연산의 단위 원(幺元, identity)을 어떻게 보느냐에 따라 달린다. 메시지 단위를 기본 원소로 볼지, byte 단위를 볼지에 따라 달라진다. 단순 write/read에서 주소 공간이 겹치면 결합법칙이 깨지지만, 메시지 단위 semantics는 이를 자연스럽게 분리해 준다. 그래서 분산 병렬 프로그래밍에서 Actor/CSP 모델이 자주 등장하는 것이다.
따라서 메시지 semantics를 메모리 사용의 단위 원으로 두고, 메모리 연산의 주소와 명령을 메시지와 결합하면 semilattice 대수 구조를 만들 수 있다. 이렇게 하면 대규모 통신 문제를 해결할 수 있다. 교환법칙은 멀티패스 out‑of‑order를 허용해 혼잡을 완화하고, 멱등성은 packet loss 재전송을 자유롭게 해 준다. 결합법칙은 여러 연산을 “merge”해 원격 노드에 전송할 수 있도록 한다.
하드웨어에서 semilattice를 구현하는 비용은 크지 않다.
ci 포인터에 fence를 걸어 재작성(rewrite)을 막으면 된다. in‑network computing에서의 덧셈 멱등 처리는 별도 이슈지만, 기본 틀은 같다.결국 프로그래밍 인터페이스는 아래와 같이 단순해진다. GPU는 한 번의 명령을 issue하고, mbarrier의 completion_tx 카운터만 기다리면 된다.
cppcde::cp_async_bulk_tensor_2d_global_to_shared(tile_smem, &tensor_map, x, y, bar); auto token = cuda::device::barrier_arrive_tx(bar, 1, sizeof(tile_smem)); // 다른 연산 수행... bar.wait(std::move(token));
전송 계층에 요구되는 것은, TMA가 내부적으로 쪼개는 미세한 LD/ST들을 하나의 메시지로 묶어, 메시지 내부에서는 데이터가 서로 순서 없이 도착해도 되고, 여러 경로로 나뉘어 전송·재전송 가능하게 만드는 것이다. 끝에 mbarrier만 올바르게 갱신되면 된다. 이러면 네트워크 BER 요구도 낮아진다.
이렇게 설명하면, 20년 전 iWARP가 이미 정의했던 Direct Data Placement(DDP)를 다시 발명한 셈이 된다.
하지만 NV(Mellanox) 쪽은 이 개념을 충분히 이해하지 못한 듯하다. AR에서 일부 DDP를 구현하긴 했지만, SEND/RECV에 대해 제대로 된 DDP를 지원할 수 있을까? 사실상 불가능하다. RoCEv2 프로토콜의 한계 때문이다.
RoCEv2는 메시지 전송 방식을 First/Middle/Last flag 하나로 정의한다. 중간 메시지는 메모리 주소 정보를 포함하지 않으므로, 반드시 in‑order로 전송되어야 한다. 이후 개선으로 메시지를 더 잘게 쪼개고 각 메시지에 RETH를 붙여 주소를 담게 했지만, SEND/RECV은 수신 buffer의 절대 주소를 알 수 없다.
또한 NIC에 포트가 2개 있는 경우, RoCEv2는 XOR 방식으로 하나를 선택해 전송하게 된다. 한 QP를 두 포트에 걸쳐 load balance하려면 어떻게 해야 할까?
반면 iWARP의 DDP는 아주 명료하다. MSN(Message Sequence Number) + MO(Message Offset)만 있으면 된다. relative offset을 쓰면 receiver 측 reorder buffer가 필요 없다. packet이 도착하는 대로 offset을 계산해서 바로 메모리에 써 넣으면 된다. 그리고 간단한 bitmap으로 하나의 메시지가 다 도착했는지 확인하고, 끝나면 mbarrier를 update하면 된다. 멱등성을 유지하기 위해 이미 mbarrier를 update한 MSN에 대해서는 이후 재전송 packet의 write를 막으면 된다.
이것이 eRDMA 구현 방식이다. 매우 단순하고 하드웨어 비용도 작다. reorder buffer가 없어 예측 불가능한 지연과 jitter를 일으키지 않는다. 동시에 fabric 전체의 multipath를 적극 활용할 수 있고, 일부 링크 장애에도 잘 견딘다.
이제 이런 관점에서 eRDMA(Weak Order), NVIDIA(Mellanox)의 Strong Order, AWS SRD의 Relaxed Order를 비교해 보면, 어떤 방식이 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