GPU 하드웨어와 성능 모델

추론 서빙을 공부하다 보면 계속 같은 벽에 부딪힌다.

• 디코드가 메모리 바운드라는데, 그 판정은 대체 어떻게 내리는 것일까
• FlashAttention은 왜 하드웨어 세대마다 처음부터 다시 짜야하는 것인가
• 서빙엔진이 Hopper에서는 FlashAttention을 Blackwell에서는 다른 커널을 고르는데, 그 분기는 무엇을 보고 결정된 것일까
• FP4로 바꾸면 정확히 어디서 시간이 절약되는 것일까?
• 배치를 키우면 처리량이 오른다고 하는데 어디까지 올려야 이득이 멈출까?

이 질문의 답들은 전부 GPU 하드웨어 안에 있다. 서빙 엔진은 하드웨어 성질을 이용하는 도구이고, 모델 아키텍처는 그 성질에 맞춰 역설계된 결과물이다.

하드웨어 층을 모르면 위의 두 층은 외어야하는 규칙덩어리로 남는다. 알고나면 대부분 유도할 수 있는 결론이 되는데.

그래서 이번 글에는 이 두가지의 목표를 갖는다.

1. 판단할 수 있게 되는 것 새로운 GPU가 나왔을때 스펙표만보고 이 워크로드는 어떤 의미고 계산할 수 있어야한다. 그 도구가 루프라인 모델이고 이런것들을
2. 커널을 읽을 준비를 하는것 워프나 공유 메모리, 뱅크 충돌, 텐서코어의 비동기 실행모델을 모르면 FlashAttention 소스는 암호문이라 알고나면 왜 여기서 이렇게했는지가 보인다.

GPU는 왜 이렇게 생겼는가 #

CPU와 GPU는 같은 목표를 다르게 푼게 아니라 애초에 다른 질문으로 다른 문제를 풀기위해 존재하는 하드웨어다.

CPU는 하나의 작업을 얼마나 빨리 끝낼 수 있는가? GPU는 단위 시간에 얼마나 많은 작업을 끝낼 수 있는가? 이다.

전자는 latency-optimized 후자를 throughput-optimized 라고 한다.

소포 배달로 비유를 해보면 CPU는 오토바이 퀵 서비스, 한 건을 빠르게 배달한다. 신호 예측을 하고 (분기문) 자주 가는 길을 외우고 (대용량 캐시), 막히면 순서를 바꿔간다 (비순차 실행). 대신 한 번에 한두건만 처리한다. GPU는 화물차다 한 건만 보내면 오토바이보다 느리지만 컨테이너 만 개를 보낼거라면 비교가 안된다.

이 차이는 실리콘 면적의 배분에서 나오는데 CPU는 다이 면적의 대부분을 제어 회로와 캐시에 쓴다. 분기 예측기와 비순차 실행 스케쥴러, 수십 MB의 L3 캐시가 존재하고 GPU는 같은 면적을 연산 유닛에 몰빵한 구조다.

H100 (132 SM)
  레지스터 파일: SM당 64K × 32bit = 256 KB  →  전체 약 33.8 MB
  L2 캐시:                                     50 MB
  → 레지스터 파일이 L2의 약 68%

일반적인 서버 CPU
  아키텍처 레지스터: 코어당 수백 바이트~수 KB
  L3 캐시:          수십 MB
  → 레지스터 파일이 캐시의 수천분의 1

CPU는 레지스터가 캐시에 비해 무시할만큼 작고 GPU는 레지스터가 캐시랑 맞먹는다.

왜 이렇게 설계되었냐면 상주하는 모든 스레드의 레지스터를 물리적으로 동시에 보유하기 어려우니 CPU의 컨텍스트 스위치는 레지스터를 메모리에 저장했다 복원하는 비용이 수백~수천사이클이라 무겁다.

GPU는 저장하고 복원할필요 없이 전환을 하는데 어느 레지스터 뱅크를 읽을지 바꾸는 것 뿐이라서 사실상 제로 오버헤드다

대신 레지스터파일이 아무리 커도 유한하기에 스레드 하나가 레지스터를 많이 쓰면 동시에 상주할 수 있는 스레드 수가 그만큼 줄어든다 이게 뒤에서 다룰 점유율의 물리적 근거인데

캐시로 지연을 줄이는 길이 막혀있기도 하다. H100의 L2 50MB를 상주 스레드 수 (132 SM x 2048 스레드 ≈ 27만)로 나누면 스레드당 약 185바이트로 CPU 코어가 MB를 독점하는 것과 비교가 되지 않는다. GPU는 캐시 적중률을 높이는 방식으로 메모리 지연에 대응할 수 없고, 다른 방법을 써야한다.

GPU가 지연을 숨기는 방식? #

GPU도 메모리에서 데이터를 가져와야하고, 그건 수백 사이클이 걸린다.

CPU처럼 캐시와 예측으로 이 지연을 줄이지 않으면 GPU는 어떻게 메모리 로딩 시간을 견딜까

답은 지연을 줄이지 않고 숨기는 것이다.

CPU의 전략 — 지연을 줄인다
  스레드 1: [계산][메모리 대기 300사이클.......][계산]
            ↑ 캐시 적중률을 높여 이 대기를 짧게 만든다

GPU의 전략 — 지연을 숨긴다
  워프 1:  [계산][────── 메모리 대기 ──────][계산]
  워프 2:        [계산][────── 대기 ──────][계산]
  워프 3:              [계산][───── 대기 ─────][계산]
  워프 4:                    [계산][──── 대기 ────]
           ↑ 한 워프가 기다리는 동안 다른 워프를 실행 → 연산 유닛은 계속 바쁨

이게 GPU 성능의 전부고, 계산하는 워프가 많으면 메모리 지연이 보이지 않고 부족하면 그대로 보인다.

뒤에 나올 점유율 occupancy라는 개념이 정확하게 이걸 재는 지표다

이 설계가 이따 나올 제약들의 뿌리인데, 레지스터를 전부 들고있어야하니 레지스터 파일이 거대해야하고 H100기준 SM당 64K개의 32비트 레지스터, 그런데도 유한하므로 스레드 하나가 레지스터를 많이 쓰면 동시에 살아 있을 수 있는 워프수가 줄어든다.

GPU 성능 #

GPU 성능을 볼땐 항상 세 가지를 봐야한다.

• 연산 능력: 초당 처리 가능한 부동 소수점연산 FLOPs, 부족하면 계산 바운드인것
• 메모리 대역폭: 초당 옮길 수 있는 바이트 GB/s, 부족하면 메모리 바운드인것
• 병렬성: 동시에 살아있는 워프수로, 부족하면 지연 노출 (둘 다 못채움)

세번째가 자주 잊히는데 연산 능력과 대역폭이 충분해도 병렬성이 부족하면 둘 다 못쓴다. 배치크기 1에서 GPU가 놀고있는 근본 원인이 이것이다.

GPU #

SM - GPU의 진짜 단위 #

GPU 코어 1만개라는 말을 보자. 스펙표에 CUDA 코어 16.896개 같은 숫자가 적혀있다.

이 숫자만 보면 만 개 이상의 독립적인 프로세서가 각자 다른일을 하는것처럼 들리지만 그렇게 이해해선 안된다.

실제 계층: GPU의 실질적 단위는 코어가 아니라 SM(Streaming Multiprocessor) 이다.

GPU (예: H100 SXM)
├── SM 0
│   ├── 파티션 0 ─ [FP32 유닛 ×32][INT 유닛][텐서 코어][워프 스케줄러][레지스터 16K]
│   ├── 파티션 1 ─ ...
│   ├── 파티션 2 ─ ...
│   ├── 파티션 3 ─ ...
│   ├── 공유 메모리 / L1 캐시 (최대 228 KB)
│   └── TMA 유닛 (Hopper 이상)
├── SM 1
├── ...
└── SM 131                      ← H100은 SM 132개

공용
├── L2 캐시 (50 MB급)
└── HBM (80~288 GB)

한 SM 안의 4개 파티션은 각각 자기 워프 스케줄러를 갖는다.

매 사이클 각 스케줄러가 준비된 워프 하나를 골라 명령을 발행한다.

즉 SM 하나의 사이클마다 최대 4개의 서로 다른 워프를 진행시킨다.

SM 단위인 이유는 세 가지 자원이 SM 단위로 배분된다.

1. 레지스터 파일: SM당 64K개(32비트 기준). 이 SM에 올라온 모든 스레드가 나눠 쓴다.
2. 공유 메모리: SM당 최대 228KB(H100). 이 SM에 올라온 블록들이 나눠 쓴다.
3. 워프 슬롯: SM당 최대 64개 워프

커널의 성능은 코어가 몇 개인가? 가 아니라 이 세자원을 얼마나 잘 나눠 썼는가로 결정된다.

그래서 SM 내부 구조를 아는것은 중요하다.

세대별 SM 구성 ex.

• A100(Ampere): SM 108개, SM당 최대 공유 메모리 164KB, register/SM 64K x 32bit
• H100(Hopper): SM 132개, SM당 최대 공유 메모리 228KB, register/SM 64K x 32bit
• B200(Blackwell): SM 148 x 2다이, SM당 최대 공유메모리 228KB,register/SM 64K x 32bit

B200은 물리적으로 두 개의 다이를 붙인 구조라 두 다이는 NV-HBI라는 10TB/s 링크로 연결되어 소프트웨어는 하나의 GPU로 보인다. 다이 경계를 프로그래머가 신경쓸필요가 없도록 설계한것인데 보이지 않는 경계가 성능 이상의 원인이 되는경우도 간간히 있다.

스레드, 워프, 블록 - 실행계층 #

만 개의 스레드를 어떻게 관리할 것인가

수 만개의 스레들르 각각 독립적으로 스케줄링하면 제어 회로가 감당하지 못한다. GPU가 면적을 아끼려고 제어 회로를 줄인 의미가 사라진다.

Warp는 32개의 스레드 묶음을 의미한다. GPU의 진짜 실행 단위인 것이다. 워프 안에 32개의 스레드는 같은 명령을 동시에 실행한다. 데이터만 다르다.

계층 구조

스레드(thread)      ─ 프로그래머가 코드를 쓰는 단위
   ↓ 32개
워프(warp)          ─ 하드웨어가 실행하는 단위 ★
   ↓ 여러 개
블록(block, CTA)    ─ 공유 메모리를 함께 쓰는 단위. 하나의 SM에 배치됨
   ↓ 여러 개
클러스터(cluster)   ─ Hopper 신설. 서로의 공유 메모리에 접근 가능
   ↓ 여러 개
그리드(grid)        ─ 커널 실행 전체

워프는 32명이 한 줄로 서서 조교의 구령에 맞춰 움직이는 소대같은 개념이다. 앞으로 가! 라는 명령에 32명이 동시에 움직이고 각자 다른위치에서 출발하지만 (다른 데이터) 동작은 같다.

32개씩 묶인 이유는 명령 하나를 32개 데이터에 적용하면 명령 디코딩 비용이 32분의 1로 줄어든다 이 방식을 SIMT(Single Instruction Multiple Threads)라고 한다. 32는 제어 회로 절감과 유연성 사이에서 NVIDIA가 고른 값이고 세대가 바뀌어도 변한적이 없다. (AMD에서는 CDNA에서 64를 쓰는데 이 차이가 두 플랫폼의 커널을 서로 옮기게 어렵게 만드는 원인 중 하나다.)

워프 다이버전스 - 설계의 대가 #

같은 명령을 함께 실행한다는 설계에는 명백한 대가가 있다. 워프 안의 스레드들이 서로 다른 분기를 타는 경우다.

if (threadIdx.x % 2 == 0) {
    A();          // 짝수 스레드 16개
} else {
    B();          // 홀수 스레드 16개
}

하드웨어는 이렇게 처리한다.

1단계: A() 실행 — 짝수 스레드만 활성, 홀수 16개는 놀고 있음
2단계: B() 실행 — 홀수 스레드만 활성, 짝수 16개는 놀고 있음
       → 총 실행 시간 = A 시간 + B 시간 (병렬이 아님!)
       → 실효 처리량 50%

이걸 워프 다이버전스 warp divergence라고 하고 분기가 32갈래로 갈리는 최악의 경우에는 처리량이 1/32가 되는것이다.

LLM Kernel에서는 이게 실제로 문제가 되는 지점이 있다. Attention 인과 마스킹이다.

시퀀스 위치에 따라 마스킹 여부가 달라지므로 순진하게 짜면 다이버전스가 생긴다. 그래서 FlashAttention 계열 커널은 마스크가 완전히 걸리는 블록은 아예건너뛰고, 경계에 걸친 블록만 마스킹 계산을 수행하는 식으로 분기를 타일 경계로 밀어낸다. 분기를 없앨 수 없으면 분기가 워프 경계와 정렬되도록 만드는것이 핵심이다

어텐션 인과 마스킹 Casual Masking: 인공지능이 문장을 만들거나 이해할때 아직 나오지 않은 미래의 단어들을 미리 보지 못하게 하는 기술.

GPGPU Tile: 타일 경계 Tile Boundary란 대규모 데이터를 작은 블록(타일) 단위로 쪼개어 병렬연산할 때, 각 타일과 인접한 타일이 만나는 구역, 느린 글로벌 메모리 접근을 줄이고 공유 메모리 활용률을 높이기 위해 전체 데이터를 작은 사각형 형태 타일로 나누는 기법. 즉 스레드들이 효율적으로 나누어가질 메모리 조각들이다.

스레드 블록 클러스터 - Hopper #

Hopper는 블록과 그리드 사이에 클러스터라는 층을 하나 넣었다. 최대 16개 블록이 하나의 클러스터를 이루고, 서로의 공유 메모리에 직접 접근할 수 있다.

이걸 분산 공유 메모리 (DSMEM, Distributed Shared Memory)라고 한다.

왜 필요했나: 공유 메모리는 SM안에 갇혀있어 블록 사이에 데이터를 주고받으려면 전역 메모리를 거쳐야 했고 왕복 비용이 크다.

얼마나 좋아지는가: 마이크벤치마크 측정에서 클러스터 내 SM간 접근 지연은 33~213 사이클로 L2 캐시나 전역 메모리 접근보다 확실히 낮다.

LLM 어디에서 쓰일까? >> 여러 블록이 같은 가중치 타일을 필요로할때 한 블록이 읽어서 클러스터 전체에 뿌린다. (TMA 멀티캐스트) L2 Traffic이 줄고, GEMM에서 A 행렬의 한 블록이 여러 N타일에 재사용 되는 상황이 이 경우다.

점유율 - 지연을 얼마나 숨길 수 있는가 #

커널 성능을 진단할 때 가장 먼저 보게 되는 지표들을 다뤄보자

워프가 부족하면 무슨 일이 일어나는가:
0장에서 GPU가 지연을 숨긴다고 했다. 숨기려면 대기중일때 실행할 다른 워프가 있어야한다. 없으면 그냥 멈춰있다.

워프가 충분할 때
연산 유닛: ████████████████████████  100% 가동

워프가 부족할 때
연산 유닛: ███░░░░░░░░███░░░░░░░░███  20% 가동
                ↑ 메모리 대기, 채울 워프 없음

점유율의 정의 #

점유율 occupancy = SM에 실제로 상주하는 워프 수 / SM이 수용가능한 최대 워프수

Hopper 기준 SM당 최대 64워프이므로, 32워프가 상주하면 점유율은 50%이다.

점유율은 각각 세 가지의 상한으로 만들어지며 그중 가장 낮은 값이 실제 점유율이다.

레지스터 압박 #

SM에는 64K개의 32비트 레지스터가 있다. 스레드 하나가 R개를 쓰게되면

상주 가능 스레드 수 = 65,536 / R
상주 가능 워프 수   = 65,536 / (R × 32)

R = 32일 때 → 64워프 (100%)
R = 64일 때 → 32워프 (50%)
R = 128일 때 → 16워프 (25%)
R = 255일 때 →  8워프 (12.5%)   ← 스레드당 최대치

복잡한 커널일수록 레지스터를 많이쓰고 그만큼 점유율이 떨어진다.

이게 텐서 코어 이야기에서 중요해지고 누산기를 레지스터에 들고있으면 레지스터가 부족해지는데, Blackwell이 누산기를 별도 메모리로 옮기는걸 적용하기도 했다. 더 아래에서 알아보겠다.

공유 메모리 압박 #

블록 하나가 공유 메모리를 S바이트 쓰게되면 SM의 공유 메모리 총량을 S로 나눈만큼만 블록이 올라간다.

H100: SM당 최대 228 KB
블록당 64 KB 사용 → 3개 블록만 상주
블록당 32 KB 사용 → 7개 블록 상주

FlashAttention은 타일을 공유 메모리에 올려놓고 작업하므로 타일 크기가 곧 점유율이다 타일을 키우면 데이터 재사용이 늘어 좋지만 점유율이 떨어진다 이 줄다리기가 커널 튜닝의 본질적 부분이다.

블록 크기 #

블록당 스레드 수가 어중간하면 자투리 낭비가 된다. 128,256 처럼 워프크기 32의 배수를 쓰는이유가 이거다.

그리고 점유율이 높다고 항상 빠른건 아니다. 여기서 자주 빠지는 함정인데 점유율 100퍼를 목표로 삼는것이 좋은건 아니다

점유율은 지연을 숨기는 능력을 나타낼 뿐이고 성능 자체가 아니라 실제로 고성능 커널중엔 의도적으로 점유율을 낮게 가져가기도한다.

왜냐면 레지스터와 공유메모리를 넉넉히 써서 데이터 재사용을 극대화하면 애초에 메모리접근자체가 줄어들고 숨겨야할 지연이 적어지므로 워프가 적어도 되기 때문에 핵심이 아니라는 것이다.

전략 A (고점유율): 작은 타일, 적은 레지스터 → 워프 64개로 지연을 숨김
전략 B (저점유율): 큰 타일, 많은 레지스터  → 메모리 접근 자체가 적어 숨길 게 없음

FlashAttention 계열은 전략 B에 가깝습니다.

점유율이 50%이상이면 대체로 충분하고 그 이상은 다른 병목을 먼저 확인하자

25% 미만이면 레지스터나 공유 메모리 사용량을 의심할 가치가 있다.

왜 의도적으로 점유율을 낮게 가져가냐

전략 A는 어차피 데이터를 가져오는데 오래걸리니 일꾼을 많이 투입하자라는 전략인데 애초에 B FlashAttention 방식은 애초에 느려터진 메모리에 가지말자는 전략이고 gpu 내부에는 아주 빠르지만 용량이 극도로 적은 Register, Shared Memory가 있고 하나의 스레드 묶음(블록)이 귀한 메모리를 많이 독차지해서 쓰면? 느린 VRAM을 갈필요없으니 데이터를 계속 재상요하여 초고속으로 계산을 끝낼 수 있다. 하지만 이 귀한 메모리는 총량이 있어서 한 블록이 레지스터를 많이쓰면 gpu내에 동시에 띄울수있는 블록의 수가 적어지니까 의도적으로 하드웨어적 점유율을 떨어트리고 메모리를 활용하자라는 전략인것이다.

점유율이 높다 = 일꾼이 많아서 땜빵 잘하고 있다 무난한 최적화

의도적 저점유율은 = 소수의 일꾼에게 최고급 장비 레지스터들을 몰아주며 땜빵없이 초고속으로 끝내자 (FlashAttention)

지연을 숨기려면 얼마나 필요한가 #

점유율얘기로 다시 넘어가서 점유율이 얼마나 필요한지는 리틀의 법칙으로 어림가능하다.

필요한 동시 진행 바이트 = 메모리 대역폭 × 메모리 지연

H100 예시:
  대역폭 3,350 GB/s, HBM 지연 약 600ns
  → 3,350 × 10⁹ × 600 × 10⁻⁹ ≈ 2 MB

즉 대역폭을 채우려면 항상 약 2MB의 메모리 요청이 "날아가는 중"이어야 합니다.

이 숫자가 커널 설계에 직접 들어온다. 워프 하나가 한 번에 요청하는 양이 작으면 워프가 많아야하고 요청 하나가 크면 (예. TMA로 큰 타일을 한 번에) 워프가 적어도 된다. Hopper가 TMA를 도입한 배경에 이 계산이 있다.

요약 #

• SM은 gpu 자원 배분의 단위로 레지스터, 공유메모리, 워프슬롯이 여기서 갈린다.
• 워프는 32개의 스레드를 가진 실행의 단위로 분기가 여기서 갈리면 처리량이 손실된다
• 워프 다이버전스는 워프 내 분기로 인한 직렬화로 어텐션 마스킹 설계에 영향이 간다.
• 클러스터와 DSMEM은 블록간 공유 메모리 접근으로 타일 재사용, TMA 멀티캐스트 기반이다.
• 점유율은 상주 워프 / 최대 워프로 지연 은닉 능력을 의미하며 높다고 빠른건아니다
• 레지스터 압박은 스레드당 레지스터가 점유율을 제한하는 것이고 텐서 코어 세대 진화 동기다.

TMA (Tnesor Memory Accelerator)는 엔비디아 Hopper, SM 90+ 아키텍처부터 도입된 글로벌 메모리와 공유 메모리간 대용량 다차원 데이터를 비동기로 전송해주는 하드웨어 엔진 및 명령어 세트를 의미한다.

메모리 계층 #

GPU가 지연을 숨기는 방법을 봤다. 이 장은 그 지연이 어디소오고 어떻게 하면 애초에 덜 발생시킬 수 있는가? 를 다룬다. 커널 최적화 작업의 대부분이 이 장의 내용이다.

Layer #

메모리는 빠를수록 비싸다. 정확히는, 빠른 메모리는 연산 유닛에 가까이에 물리적으로 붙어있어야하는데 그공간이 유한하기에 그래서 몯느 컴퓨터는 작고 빠른것부터 크고 느린것까지 계층을 만들어서 쓴다.

H100 기준 실제 수치

용량(SM당/전체)   지연          대역폭       비고
┌──────────────┐
│  레지스터      │   256 KB / SM     ~1 사이클     ~수십 TB/s   스레드 전용
├──────────────┤
│ 공유메모리/L1  │   228 KB / SM     ~30 사이클    ~수 TB/s     블록 내 공유, 수동 관리
├──────────────┤
│  L2 캐시      │   50 MB (전체)    ~200 사이클   ~수 TB/s     GPU 전체 공유, 자동
├──────────────┤
│  HBM (VRAM)  │   80 GB           ~600 사이클   3.35 TB/s    ★ 여기가 병목
├──────────────┤
│ CPU DRAM     │   TB급            ~수 μs        ~64 GB/s     PCIe 5.0 경유
├──────────────┤
│  NVMe SSD    │   수십 TB          ~수십 μs      ~7 GB/s      
└──────────────┘

HBM과 공유 메모리 사이의 격차가 20배다, 지연도 20배 대역폭으로도 비슷한 배수차이가 난다.

이 격차가 커널 최적화의 유일한 큰 게임이다. 정리하면

HBM에서 한 번 가져온 데이터를 공유 메모리 레지스터 안에서 여러번 최대한 재사용하라

FlashAttention이 하는 일을 한문장으로 축약하면 정확하게 위 표현이다.

어텐션 스코어 행렬을 HBM에 쓰고 다시 읽는 대신에 타일 단위로 잘라서 공유 메모리 안에서 처리하고 끝낸다

알고리즘이 빠귄게 아니라 데이터가 머무는 위치가 바뀐것이다.

KV Cache (HBM -> CPU DRAM -> SSD)는 위 표 아래를 한칸 더 연장시킨걸로 PCIe가 HBM의 2% 수준이라는 사실이 그 계층의 손익분기를 결정하는 사실은 이 표에서 알수있다.

cpu 메모리 갖다쓰는건 그 대역폭이랑 잃는것들 메모리 총량등 잘 계산해서 정하자.

병합 저근 - 32개의 스레드가 메모리 읽는 방식 #

워프가 흩어진 주소를 읽으면: 워프가 32개 스레드가 각각 메모리를 읽을 때, 이 32개의 요청이 어떻게 묶여 처리되는지 실효 대역폭을 결정한다.

메모리는 바이트 단위로 읽히지 않는다. **트랜잭션이라는 고정 크기 단위 (보통 32~128바이트 섹터)**로 읽히는데 4바이트만 필요해도 32바이트를 가져온다.

병합 Coalescing 이 일어나는 조건

좋은 경우 — 연속 접근
스레드:   0    1    2    3   ...  31
주소:   0x00 0x04 0x08 0x0C ... 0x7C
        └────────── 연속된 128 바이트 ──────────┘
        → 128바이트 트랜잭션 1번으로 32개 요청을 모두 처리 ✅

나쁜 경우 — 스트라이드 접근 (32바이트 간격)
스레드:   0      1      2         31
주소:   0x000  0x020  0x040  ... 0x3E0
        → 각 스레드마다 별도 트랜잭션 → 32번의 트랜잭션 ❌
        → 필요한 건 128바이트인데 실제로 1,024바이트를 읽음 (효율 12.5%)

병합 coalescing은 워프의 요청들이 연속된 주소에 몰려서 적은 수의 트랜잭션으로 처리되는 상태를 뜻한다.

HBM에서 대역폭 3.35TB/s는 완벽하게 병합된 경우의 숫자다. 접근이 흩어지면 그 절반의 4분의 1 심하면 8분의 1까지 실효 대역폭이 떨어질 수 있다.

메모리 바운드인 LLM 디코드에서 실효 대역폭이 절반이 된다는 것은 속도가 절반이 된다는 뜻이다. 알고리즘을 바꾸지 않고 데이터 배치만 바꿔서 2배의 성능을 얻거나 잃을 숭 ㅣㅆ다.

이는 KV Cache 레이아웃에서 나타나는데, Paged KV Cache는 블록이 물리적으로 흩어져있어 PagedAttention이 순진하게 구현하면 접근이 병합되지 않는다. 그래서 Paged KV를 다루는 어텐션 커널은 블록 내부는 연속으로 배치하고 블록 단위로만 점프하도록 레이아웃을 설계한다 블록 크기가 16토큰인것도 이 트랜잭션 크기와 관련이 있다.

진단하는 방법 #

Nsight Compute에서 다음 지표를 볼 수 있다.

ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./my_program

섹터수 / 요청수가 병합 효율이다. 이상적이면 요청 1개당 4섹터 128바이트 = 32바이트 x 4이고 비율이커지면 흩어진 접근이다.

공유 메모리 뱅크 충돌 #

공유 메모리도 완전한 병렬은 아니다

공유 메모리는 빠르지만 32개 스레드가 동시에 아무 주소나 읽을 수 있는것은 아니다.

32개의 뱅크 #

공유 메모리는 32개의 뱅크로 나뉘어 있고, 각 뱅크는 4바이트 폭이다.

주소가 뱅크에 배정되는 규칙은 단순하다.

뱅크 번호 = (바이트 주소 / 4) % 32

주소:  0    4    8   12  ...  124  128  132
뱅크:  0    1    2    3  ...   31    0    1   ← 128바이트마다 한 바퀴

한 사이클에 각 뱅크는 하나의 접근만 처리한다. 워프 32개 스레드가 32개의 서로 다른 뱅크를 건드리면 한 사이클에 전부 끝난다. 두 스레드가 같은 뱅크의 다른 주소로 건드리면 두 사이클로 직렬화 된다.

32개의 창구가 있을때 32명 모두가 다른 창구로 가면 바로 끝나지만 한창구로만 몰리면 32명이 순차적으로 대기하게 되는 상황이 발생한다.

// 32×32 float 타일을 공유 메모리에 두고 열 방향으로 읽는 경우
__shared__ float tile[32][32];

// 스레드 i가 tile[i][0]을 읽는다 (열 접근)
// 주소 = i × 32 × 4바이트 = i × 128바이트
// 뱅크 = (i × 128 / 4) % 32 = (i × 32) % 32 = 0
//        → 32개 스레드가 전부 뱅크 0! 32-way 충돌 ❌ 32배 느려짐

행 방향 접근 tile[0][i] 은 완벽하게 분산되는데 열 방향 접근은 최악이 된다.

그리고 행렬 곱셈 커널은 한 행렬은 행 방향, 다른 행렬은 열 방향으로 읽는게 자연스러운 구조라서 이 문제를 반드시 만나게 된다,

패딩과 스위즐 #

패딩은 간단하고 고전적인 방식인데

__shared__ float tile[32][33];   // 33! 한 칸 더

// 뱅크 = (i × 33 × 4 / 4) % 32 = (i × 33) % 32 = i % 32
//        → 32개 스레드가 32개 뱅크에 고르게 분산 ✅

배열의 행 길이를 1늘려줘서 충돌을 없앤다 메모리를 3% 더쓰고 32배의 손실을 막는 기법이다.

스위즐링은 현대 커널의 표준인데 패딩은 메모리를 낭비하고 무엇보다 TMA와 텐서 코어가 요구하는 정렬 조건을 깨뜨린다 그래서 최신 커널은 주소를 XOR 연산으로 뒤섞는 스위즐 방식을 쓴다.

스위즐된 주소 = 원래 주소 XOR (행 번호에서 유도한 패턴)
-> 데이터의 논리적 위치는 그대로, 물리적 뱅크 배정만 재분산

CUTLASS, FlashAttention이 이 방식을 쓰고, Hopper의 TMA는 스위즐 모드를 하드웨어 기능으로 지원한다. TMA 디스크립터에 스위즐 패턴을 지정하면 전송과정으로 자동 적용된다.

브로드캐스트는 충돌이 아니다

한가지 중요한 예외, 여러 스레드가 같은 뱅크의 같은 주소를 읽으면 충돌이 아니다. 하드웨어가 한 번 읽어서 모두에게 뿌린다 (브로드캐스트)

이 성질이 실전에서 유용한데 모든 스레드가 참조하는 상수나 스케일 펙터를 공유 메모리에 두어도 비용이 없다.

HBM #

병목의 정체기도한데 일단 메모리 대역폭이 왜 성능을 지배하는지 알아보자

디코드는 메모리 바운드 작업이다. 그 근거를 확인해보자

70B 모델을 BF16으로 디코드 한 스텝을 돌리게 되면

읽어야 하는 가중치: 70B × 2바이트 = 140 GB
H100 대역폭:      3,350 GB/s
→ 최소 소요 시간 = 140 / 3,350 ≈ 41.8 ms

같은 스텝의 계산량 (배치 1):
2 × 70B × 1 = 140 GFLOP
H100 BF16 연산능력: 989 TFLOPS
→ 최소 소요 시간 = 140 × 10⁹ / 989 × 10¹² ≈ 0.14 ms

41.8 ms vs 0.14ms 약 300배 차이다. 즉 계산은 0.14ms만에 끝났는데 41.8ms동안 데이터 옮기는데 쓰게 되는것이다. gpu의 연산 유닛은 그 동안 99,7% 놀고있다.

이 계산 하나가 나타내는 바는

• 양자화가 효과적인 이유: 140GB, 70GB로 줄여서 시간이 절반이 된다. 계산량은 그대로인데 2배 빨라진다.
• 배치크기를 키우는게 공짜라고 느껴지는 이유: 배치 32여도 가중치는 140GB만 읽기에 계산은 32배 늘지만 0.14ms가 4.5ms가 될 뿐, 41.8ms안에 들어가기에, 사실 별 지연적으로 문제는 못느끼는 것이다.
• 왜 투기적 디코딩이 좋은가?: 남는 계산 자원으로 여러 토큰을 미리 검증하기에 메모리를 덜 쓰면서 계산자원을 더 쓰는 트레이드오프를 한것

HBM의 세대별 스펙 변화

HBM4부터 인터페이스 폭이 스택당 1024비트에서 2048비트로 두 배가 되었다.

핀 속도로 11Gb/s를 넘겨서 스팩당 대역폭이 3TB/s를 상회한다. Rubin GPU는 이런 스택을 8개 달아 288GB를 구성한다.

여기서 실무적으로 알아둘 것은 Rubin의 22TB/s는 원래 목표였고 메모리 공급자들 사이에 이 속도를 맞추지못하면 초기 물량은 20TB/s 수준으로 조정되었다는 보도가 있다. 스펙표 숫자 출하 시점에 바뀌는 일은 이 분야에서 드물지않고 용량 계획을 세울때는 벤더 발표를 그대로 믿기보다 실측 벤치마크가 나올때까지 여유를 두는편이 안전하다.

용량이 대역폭만큼 중요한 이유

대역폭 이야기를 하면 용량도 잊기 쉬운데, 서빙에서 용량 당연히 중요하다 동시 사용자 수를 결정하기 때문에 서빙 노트에서 KV 캐시 계산을 보자

GPU 메모리 = 모델 가중치 + KV 캐시 + 활성값 + 여유분
             └ 고정 ┘   └ 사용자 수에 비례 ┘

가중치가 차지하는 비율이 낮을수록 → KV 캐시에 쓸 공간이 많음 → 동시 사용자 많음

B200(192GB)에서 B300(288GB)으로 50% 용량 증가와 왜 큰 사건인지 설명한다. 1조 파라미터 모델을 FP4로 올리면 약 500GB이므로 NVL72 랙의 72장 중 2장이면 가중치가 들어가고, 나머지 70장을 전부 KV 캐시와 동시 요청에 쓸 수 있다.

L2 캐시 #

자주 잊히는 층인데 L2는 자동으로 동작하므로 프로그래머가 신경쓸일이 적다.

근데 LLM 서빙에서 의외지점에서 영향을 준다.

여러 요청이 같은 가중치를 읽을때 인데 배치 안에서 32개의 요청이 모두 같은 레이어를 통과하므로 첫 요청에 읽어온 가중치가 L2에 남아 나머지가 재사용한다.

H100에는 50MB L2는 레이어 하나의 가중치를 담기에 충분한 경우가 많다.

그래서 배치를 키우면 대역폭 요구가 선형으로 늘지 않는다. 앞의 계산에서 배치 32여도 가중치는 140GB만 읽는다고 한 것이 성립하는 물리적 근거가 L2와 캐시 계층이다.

커널 수준의 활용 #

L2를 명시적으로 제어하는 수단도 있다.

// 특정 데이터를 L2에 상주시키도록 힌트 (CUDA 11+)
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

자주 재사용되는 작은 텐서 (ex. 라우터 가중치, 정규화 파라미터)를 L2에 고정하는 식으로 쓸 수 있다. 다만 효과가 워크로드에 의존하므로, 다른 최적화를 다 하고 난 뒤에 시도할 항목이다.

요약 #

커널을 볼때 메모리 최적화 체크리스트는 아래 순서대로 보자

1. HBM 접근 횟수를 줄였는가? 잘못됐을때 수 배의 손실이 있으며 진단 방법은 총 이동 바이트를 계산한다
2. 전역 메모리 접근이 병합되었는가? 2~8배 정도의 손실이 있을 수 있고 섹터/요청 비율을 확인하자
3. 공유 메모리 뱅크 충돌이 없는가? 최대 32배 손실이 날 수 있으며 shared_ld_bank)confilict 지표를 보자
4. 데이터 타입이 필요 이상으로 크지 않은가? 최대 2~4배 손실이 날 수 있으며 양자화를 검토하자.
5. L2 재사용을 살렸는가? 수십 %의 손실이 날 수 있으며 배치 구성과 접근순서를 확인하자

1번이 압도적으로 중요하다. 2~5번은 1번을 제대로 한 뒤에 의미가 있고 아무리 완벽하게 병합해도 100번 읽을거 10번 줄이는 것 보다 못하니까.

텐서 코어 - 왜 세대마다 커널을 다시 짜는가 #

FlashAttention은 하드웨어 세대마다 재작성이 되어야 한다.

왜 그럴까? 재작성해야되는 이유를 알아보자.

텐서 코어가 등장한 배경 #

일반 연산 유닛으로는 행렬 곱이 비효율적이다.

딥러닝 연산의 90% 이상이 행렬 곱셈이다. 일반적인 FP32 유닛으로 행렬곱을 하면 이런 흐름이 된다.

매 곱셈-덧셈마다:
  1. 레지스터에서 A의 원소 읽기
  2. 레지스터에서 B의 원소 읽기
  3. 곱한다
  4. 누산기에 더한다
  5. 레지스터에 쓴다
→ 실제 연산(3, 4)보다 데이터 이동(1, 2, 5)의 비용이 큼

한 번의 명령이 하나의 곱셈-덧셈만 처리하므로 명령 발행, 레지스터 접근의 오버헤드 연산량 대비가 크다.

작은 행렬 곱 전체를 한 명령으로: 텐서코어는 작은 행렬 블록의 곱셈 누산 (MMA, Matrix Multiply-Accumulate)을 하나의 명령으로 어리하는 전용 유닛이다.

D = A x B + C

Volta 세대: 4x4 행렬 단위
현재 세대: 훨씬 큰 타일 (예: 64 x 256 x 16)

한 명령이 수천 번의 곱셈-덧셈을 수행하므로 명령 오버헤드가 분산되고, 내부 데이터가 경로로 전용 배선이라 레지스터 왕복이 줄어든다.

얼마나 좋아지냐면 H100 기준 비교표다

FP32 일반 연산 유닛:      67 TFLOPS
BF16 텐서 코어:          989 TFLOPS   ← 약 15배
FP8 텐서 코어:         1,979 TFLOPS   ← 약 30배

딥러닝 성능의 대부분이 텐서 코어에서 나온다. 일반 연산 유닛을 쓰는 커널은 시작부터 15배의 손해를 안고 간다. 이게 텐서 코어를 쓰느냐가 커널 설계의 첫번째 질문인 이유 사실 당연한거

쓰기가 까다롭긴하다 텐서코어는 데이터가 정확한 레이아웃으로 정확한 위치에**** 있어야 동작한다.

어느 스레드가 어느 원소를 들고있어야 하는지 하드웨어적으로 정해져있는 것이다.

이 레이아웃 요구가 세대마다 바뀌고 그래서 커널을 다시 짜야한다.

세대별 진화 #

이 절이 이 장의 핵심이다. 각 세대가 직전 세대의 어떤 병목을 풀었는지 따라가면 진화의 논리가 보인다.

Volta(V100) - 텐서 코어의 등장 #

mma.sync 명령, 워프 단위 동기 실행 이다. 워프의 32개 스레드가 협력해서서 하나의 작은 행렬곱을 수행하고 명령이 끝날 때까지 기다린다.

남은 문제: 동기 실행이므로 텐서 코어가 계산하는 동안 그 워프는 아무것도 못한다. 데이터를 미리 가져오는 일도 다른 워프에게 맡겨야 한다.

Ampere(A100) - 비동기 복사의 시작 #

cp.async 도입, 전역 메모리에서 공유 메모리로 복사를 레지스터를 거치지 않고 비동기로 수행한다.

왜 중요한가: 이전에는 전역 -> 레지스터 -> 공유 메모리 순으로 옮겨야해서 레지스터를 잡아먹고 스레드가 대기했다. 이제 복사를 걸어놓고 계산을 계속할 수 있다.

소프트웨어 파이프라이닝이 여기서 표준이 되고 다음 타일을 가져오면 현재 타일을 계산하는 구조다.

남은 문제: 복사 자체는 여전히 모든 스레드가 주소 계산에 참여해야한다. 스레드당 30~40개의 레지스터가 주소 산술에서 소모된다.

Hopper (H100) - 두 개의 큰 변화 #

WGMMA - 워프그룹단위 비동기 MMA

wgmma.mma_sync. 두가지가 바뀌었다.

• 워프그룹 단위: 4개의 워프 (128스레드)가 하나의 단위로 협력한다. 더 큰 타일을 한 명령으로 처리한다.
• 비동기: 명령을 발행하고 나면 결과를 기다리지 않고 다른 일을 할 수 있다.

효과가 큰데, 마이크로벤치마크 연구에 따르면 WGMMA는 Hopper 이론 성능의 95%에 도달하는 반면, 하위 호환용 mma는 62.9%에 그친다. 같은 하드웨어에서 명령을 어떤 걸 쓰느냐로 3분의 1이 갈리는 것이다.

또 하나 중요한 변화가 있다. WGMMA는 피연산자 A를 공유메모리에 직접 읽을 수 있다. 레지스터로 옮길 필요가 없어 레지스터 압박이 줄고 앞 장에서 본 점유율이 개선된다.

TMA - 텐서 메모리 가속기

이게 Hopper의 진짜 혁신인데, cp.async도 주소 계산은 스레드가 했다. 다차원 텐서에서 타일 하나를 가져오려면 각 스레드가 자기 담당 주소를 계산해야하고, 경계 검사도 해야한다.

해법으로는 TMA 전용 하드웨어 엔진을 쓰는것이다. 호스트쪽에서 텐서모양 스트라이드, 타일크기, 스위즐 모드를 담당하는 디스크립터를 만들어두고 커널에서는 스레드 하나가 타일 좌표만 주면서 cp.async.bulk.tensor 명령을 발행한다. 나머지는 TMA 엔진이 필요하다.

Ampere 방식
  128개 스레드 전부가 주소 계산 → 각자 복사 요청 → 30~40 레지스터 소모

Hopper TMA 방식
  스레드 1개가 "이 좌표의 타일 가져와" 한 마디 → TMA 엔진이 전부 처리
  나머지 127개 스레드는 그동안 계산 수행 → 레지스터 해방

비유 - 이사할 때 직원 128명이 각자 짐을 나르던 것을, 전담 물류팀에게 3층 안방 짐 가져와 한마디로 맡기는 것이다. 직원들은 그동안 본업을 한다.

얼마나 좋아지나: 타일 크기가 4KB를 넘으면 TMA의 처리량이 H100의 3.35TB/s HBM 대역폭을 포화시킨다, 스레드당 30~40개의 레지스터도 해방된다. 비동기 파이프라인 측정에서는 작은 블록 기준 최대 39.5%의 처리량 향상이 보고되었다.

둘을 합치면 - 워프 특화 Warp Specialization #

TMA, WGMMA는 서로 다른 하드웨어 유닛이다. 따라서 동시에 돌릴 수 있고, 돌려야 최대 성능이 난다. 여기서 현대 커널의 표준 구조가 나온다.

생산자 워프그룹 (Producer)          소비자 워프그룹 (Consumer)
       │                                    │
  TMA로 타일 N+1 가져오기              WGMMA로 타일 N 계산
       │                                    │
  mbarrier로 "도착했음" 신호  ──────→  mbarrier 대기 후 진행
       │                                    │
  (다단계 순환 버퍼로 반복)

→ 데이터 이동과 계산이 완전히 겹침

같은 커널안에서 워프그룹마다 역할을 다르게 주는 것이 워프 특화다. Hopper 이전에는 모든 워프가 같은 루프 본문을 실행해야 했으므로 이 구조가 불가능했다.

이게 왜 FlashAttention-3가 FlashAtteion-2와 완전히 다른 코드인지의 답이다.

알고리즘은 같지만 **실행 구조가 다르기에 FA3가 H100에서 840TFLOPS(이론치 약 85%)**를 달성한 것은 이 구조 덕이고 FA2는 A100에서 이론치 50~73% 수준이다.

Blackwell B200/B300 - tcgen05와 텐서메모리 #

FP4/FP6 같은 초저정밀 텐서 코어에서 처리 속도가 너무 빨라지자.

레지스터 파일 대역폭이 병목이 되었다. 텐서 코어가 요구하는 속도로 피연산자를 공급할 수가 없다.

또 하나, 타일이 커질수록 누산기 레지스터를 잡아먹는다. 앞 장의 레지스터 압박 문제로 극단적으로 나간다.

해법1. TMEM (Tensor Memory) #

누산기를 레지스터에서 빼내어 전용 온칩 메모리로 옮겼다.

Ampere/Hopper: 누산기 = 레지스터 (스레드가 조각조각 나눠 보유)
Blackwell:     누산기 = TMEM (CTA 범위의 별도 2차원 메모리 공간)

sm_100 기준 CTA당 128개 레인 x 512개 컬럼 구조고 SM당 256KB이다. 레지스터 파일이 MMA 실행중에는 해방되어, 발행과 마무리 단계에만 쓰이면 된다.

유용한 사고 모델, TMEM은 텐서 코어에게 있어 L1 캐시가 ALU에게 갖는 관계와 같다.

대신 프로그래머가 직접 관리를 해야한다. 할당, 해제, 데이터 복사를 명시적으로 해야하고 에필로그에셔 결과를 읽으려면 tcgen95.ld로 레지스터에 가져와야한다. 각 워프가 TMEM에 4분의 1만 접근할 수 있어 에필로그에 워프그룹 전체가 필요하다.

CTA는 Cooperative Thread Array의 약자로 엔지디아 cuda모델등에서 스레드 블록을 부르는 말이다 수십~수백개의 스레드가 모인 하나의 그룹이며 gpu 하드웨어 코어인 스트리밍 멀티프로세서에 통째로 할당되어 실행된다.

해법2 단일 스레드 발행 #

tcgen05.mma는 CTA 전체를 대표해 스레드 하나를 발행한다. Hopper의 워프그룹 범위에서 한 단계 더 나아간 것인데, 피연산자 전부 SMEM, TMEM 이라는 CTA 공유 공간에 있으므로 가능해진 설계다.

CTA 페어 #

클러스터 안에서 랭크의 마지막 비트만 다른 두 CTA(0과 1, 4와 5 등) 짝을 이뤄 하나의 MMA를 수행할 수 있다.

1SM 모드:  CTA 하나가 M×N 타일 처리
2SM 모드:  CTA 페어가 2M×N 타일 처리
           A 타일은 각자 다른 것을 로드
           B 타일은 분할해서 DSMEM으로 공유 → 로드량 절반

이렇게 하면 같은 계산에 필요한 SMEM 용량과 대역폭이 줄어든다. 갖아 큰 타일 m256n128k64는 CTA 하나의 TMEM 예산으로는 담을 수 없어 2SM 모드가 필수다.

얼마나 좋아지냐면 tcgen05는 데이터타입에 따라 Hopper WGMMA의 2~4배 처리량을 낸다.

history:

1. Volta: 워프단위 동기 발행이고 누산기 위치는 레지스터며 스레드가 수동으로 데이터를 이동해 행렬곱 전용에 기본 형태다
2. Ampere: 워프단위 동기이며 누산기 위치는 레지스터고 데이터 이동이 스레드 수동에서 cp.async 비동기로 발전하여 파이프라이닝 문제를 해결했다.
3. Hopper: 워프 그룹단위이며 비동기고 누산기는 레지스터에 있다. 데이터 이동은 TMA 엔진으로 레지스터 해방과 워프 특화로 해결했다.
4. Blackwell: 스레드 한개가 발행범위가 되었으며 누산기 위치는 TMEM, 데이터 이동은 TMA + CTA 페어로 레지스터 대역폭 병목을 해결했다.

gpu의 발전이 한 방향으로 일관되게 움직이고 있다. 데이터 이동과 누산을 스레드의 레지스터에서 점점 떼어내어 전용 하드웨어와 전용 메모리로 옮기는 방향이고 스레드는 점점 명령을 발행하고 결과를 받는 역할만 남는다.

실무에서 쓰이는 지점 #

왜 커널이 하드웨어마다 다른 경로를 탈까?

SGLang의 백엔드 선택 로직을 예로 들어보자. 분기가 있는데

Hopper (SM90)  → FlashAttention 3 (WGMMA + TMA 워프 특화 활용)
Blackwell (SM100/103) → TRT-LLM MHA 또는 CUTLASS 경로 (tcgen05 활용)
그 외 / 폴백    → FlashInfer 또는 Triton

워크스테이션 blackwell 주의점: 워크스테이션급 Blackwell(SM120, RTX PRO)계열은 클러스터 크기1을 넘길 수 없어서 tcgen05의 CTA페어 MMA를 아예 쓸 수 없다.

그 결과 FlashInfer의 CUTLASS-Blackwell 경로 tcgen05사용하는 곳을 못타고 Triton 경로로 떨어지게 된다. 즉 처리량이 낮아지게 된다.

SM100 전용 커널을 SM120으로 옮기는 것은 기계적으로는 가능하지만 가장 큰 타일을 기준으로 PTX 명령 수가 약256배 늘어나며 달성 처리량은 SM120 최적치 40~70% 수준에 그친다 그리고 SM120 최적치 자체가 SM100의 일부에 불과하다

Blackwell 지원이라는 표기는 SM100(데이터센터) SM120(워크스테이션)에서는 전혀 다른 의미로 벤치마크를 볼 때 어느쪽인지 확인해야한다.

PTX는 CUDA PTX (Parallel Thread Execution)으로 gpu 프로그래밍에서 ptx는 저수준 가상 머신 및 명령어 집합 아키텍처를 뜻한다.

데이터센터는 기업 클라우드 서비스 제공사의 수천대의 서버를 모아 대규모 클라우드 온라인 서비스를 24시간동안 처리하는 시설이며 워크스테이션은 1인이 책상 아래나 두는 타워형 또는 고성능 모바일 형태라고 생각하면 된다

커널 소스를 읽을때 찾아볼만한거 ㅊㅊ #

CUTLASS, FlashAttention 소스를 열었을 때 이런것들이 눈에 들어오면 이 장을 이해한 것이다.

// TMA 디스크립터 생성 (호스트 측)
make_tma_copy(...)               → 텐서 모양·스트라이드·스위즐 지정

// 다단계 순환 버퍼
Stages = 4                       → 파이프라인 깊이

// 배리어
cutlass::arch::NamedBarrier / mbarrier   → 생산자-소비자 동기화

// 워프 역할 분기
if (warp_group_idx == 0) { /* 생산자 */ } else { /* 소비자 */ }

// Blackwell 전용
tcgen05.mma / make_fragment_C()  → TMEM 누산기
cta_group::2                     → CTA 페어 모드

수치 정밀도 #

양자화를 알아볼때 이게 왜 통하는 지 하드웨어 수준에서 알아보면 도움이 될 것이다.

부동 소수점의 구조 #

형식 이름부터 알아야하는데

FP8에서 E4M3과 E5M2 두 종류가 있고, NVFP4, MXFP4가 다르다고 하는데, 이 이름들이 무엇을 뜻하는지 모르면 어느 것을 골라야할지 판단할 수 없다.

부동소수점 숫자는 세 부분으로 이루어진다.

[부호 S][지수 E][가수 M]

값 = (-1)^S × 1.M × 2^(E - bias)

• 부호(sign): 1비트, 양수/음수
• 지수(exponent): 표현 가능한 범위를 결정하고 비트가 많을수록 아주 크거나 작은 수를 담을 수 있음.
• 가수(mantissa): 정밀도를 결정하고. 비트가 많을수록 값을 촘촘하게 표현한다.

지수는 자릿수, 가수는 유효숫자로 "약 3×10²³"과 "3.14159×10²³"의 차이가 가수 비트의 차이, "10²³까지 표현 가능"과 "10³⁸까지 표현 가능"의 차이가 지수 비트의 차이이다. 걍 소수점 얼마나 많이가 가수 숫자 얼마나 크게가 지수

E4M3는 지수 4비트 가수 3비트라는 뜻이다. 부호 1비트를 더하면 8비트가 된다.

이 규칙만 알면 아래 표만 봐도 이해가 될거다.

왜 BF16이 FP16을 밀어냈는가

이 대목이 딥러닝에서 실제로 벌어진 중요한 사건이다.

FP16은 가수가 10비트라 정밀도가 좋지만 지수가 5비트뿐이라 범위가 좁다. 표현 가능한 최댓값이 약 65,504라서 학습 중 그레디언트나 활성값이 이 값을 넘으면 무한대가 되어버린다. (오버플로우)

그래서 FP16 학습에는 손실 스케일링 이라는 번거로운 기법이 필요하다.

손실에서 큰 수를 곱해 그레디언트를 표현 가능한 범위로 끌어올린뒤에 다시 나누는 번거로운 기법이고 스케일 값을 잘못 고르면 학습이 터진다

BF16은 정 반대로 가수를 7비트로 줄이는 대신 지수를 8비트로 유지해 FP32와 똑같은 범위를 갖게했다. 정밀도는 나빠졌지만 오버플로우가 없으니 손실 스케일링을 안해도 되는것이다

결론적으로 딥러닝에서 정밀도보다 범위(다이내믹 레인지)가 중요하고 신경망은 값이 조금만 부정확해도 학습으로 보정되지만 무한으로 발산하면 복구가 불가능해 이 원칙이 뒤에 나올 마이크로스케일링 형식 설계에서도 그대로 이어진다.

마이크로스케일링 - 4비트를 쓸 수 있게 만든 아이디어 #

4비트로는 아무것도 표현할 수 없다

FP4 E2M1은 지수 2비트 가수 1비트 표현가능한 값이 총 16개 뿐이다. 부호 포함

E2M1이 표현하는 값들:
0, ±0.5, ±1, ±1.5, ±2, ±3, ±4, ±6

신경망 가중치를 이 16개의 값으료 표현하라는 것은 무리이고, 어떤 레이어의 가중치는 0.0001 근처에 몰려있고 다른레이어는 10근처에 있는데 같은 16개값으로는 둘 다 담을 수 없다.

작은 그룹마다 스케일을 따로 준다?
-> 마이크로스케일링 microscailing의 아이디어는 단순하다.

값들을 작은 블록으로 묶고, 블록마다 공유하는 스케일 팩터를 하나로 둔다. 실제 값 = 4비트 값 x 블록 스케일

블록 (16개 값)
[3, -2, 1, 4, ...]  ← 각각 4비트로 저장
        × 0.0037    ← 블록 스케일 (별도 저장)
= [0.0111, -0.0074, 0.0037, 0.0148, ...]  ← 복원된 실제 값

신경망 가중치는 국소적으로 비슷한 크기를 가져서 인접한 16개의 가중치가 서로 1000배씩 차이나는 경우는 드물기에, 블록 단위로 스케일을 맞추면 4비트의 16단계로도 그 블록안의 분포를 꽤 잘 근사할 수 있다.

이게 위에서 무한으로 발산하는건 못잡아도 표현력이 부족한건 커버가 가능하다는 의미.

MXFP4, NVFP4의 차이 #

같은 아이디어인데 블록 크기와 스케일 형식이 다르다. 이 두가지가 정확도와 오버헤드를 거른다.

1. MXFP4(OCP표준): 원소 형식은 FP4(E2M1) 블록크기는 32개 블록 스케일은 FP8 E8M0으로 전역스케일은 없고 스케일 오버헤드는 낮다.
2. NVFP4(NVIDIA): 원소형식 FP4(E2M1)로 같고, 블록크기는 16개 FP8 E4M3로 블록스케일하며 전역스케일은 FP32 하나 추가, 스케일 오버헤드는 2배다.

위 두개 차이를 알아보면 일단 블록크기가 32 vs 16인데 블록이 작을수록 그 안의 값들이 더 비슷하므로 스케일이 잘 맞는다. 이상값(outlier) 하나가 블록 전체에 스케일을 망치는 확률도 줄어든다.

차이2는 E8M0 vs E4M3 인데 E8M0은 지수 8비트에 가수 0비트라 즉 2의 거듭제곱만 표현할 수 있다 스케일이 0.5, 1, 2, 4, 8 ... 이런값만 가능하다 반면에 E4M3는 가수가 3비트 있어 1.5, 2.5배같은 중간값 표현도 가능하다.

어떤 블록의 최댓값이 3.7이라고 합시다.

E8M0 스케일: 4밖에 못 씀 (2^2)     → 3.7/4 = 0.925, 표현 범위의 92.5%만 사용
E4M3 스케일: 3.75를 쓸 수 있음      → 3.7/3.75 = 0.987, 거의 꽉 채워 사용

대부분의가중치 분포는 E8M0의 극단적인 다이나믹 레인지를 필요로 하지 않는다. 오히려 자기 작업 범위 안에서 스케일이 촘촘한 편이 이득이다. 그래서 같은 블록 크기로 비교해도 NVFP4가 낫다.

그 다음으로는 2단계 스케일링 차이다. NVFP4에서는 블록 스케일 E4M3 위에서 텐서 전체를 위한 FP32 전역 스케일을 하나더 둔다. 텐서 전체의 크기 수준을 먼저 맞춰놓고 블록 스케일을 국소 조정을 하는 구조다.

둘의 차이가 얼마나 나는 것인가 #

FP4 학습 연구에서 350M 규모 Llama 스타일 모델로 스케일 형식을 전수 비교한 결과가 명확하다.

블록 크기 16 고정에서 E1M6은 완전히 발산했고, E3M4와 E4M3가 가장 좋았다. 블록 크기는 8, 16, 32, 64, 128을 비교했을때 작을수록 좋지만 16아래로는 수확 체감이었다.

즉 NVFP4가 고른 블록 16, E4M3는 실험적으로 확인된 지점이고 일반적인 경향은 아래처럼 정리된다.

• 소규모 블록 조건에서는 NVFP4는 정확도와 안정성 모두에서 INT4, MXFP4를 상회한다
• 다만 MXFP4도 보정기법을 쓰면 격차가 크게 좁혀지고 MR-GPTQ처럼 양자화 전에 블록 단위 하다마드 회전을 적용하면 이상값 블록 내 모든 채널로 분산되어 32개 블록의 최악 조건이 완화되기도한다.
• 소수의 민감한 레이어나 이상값 채널만 FP8/BF16으로 남기는 혼합 정밀도가 흔하고 이 방식으로 가중치 메모리 30%이상 절감에 정확도 손실 1%미만을 보고한 사례들도 있다.

하드웨어와 툴체인 묶음 #

어떤 형식을 고르느냐가 어떤 하드웨어에서 돌릴 수 있는가를 결정할 수 있다.

• NVFP4는 NVIDIA Blackwell의 5세대 텐서 코어가 네이티브로 처리한다. 원소 그룹화, 동적 스케일링, 4비트 행렬 연산이 하드웨어에서 자동으로 처리된다.
• MXFP4는 OCP 표준이라, AMD CDNA4(MI350 시리즈)도 네이티브 지원한다.
• 지원없는 하드웨어에서는 소프트웨어 에뮬레이션이 오히려 느려진다.

체크포인트를 배포할 때 이 선택이 사용자의 하드웨어를 제약한다 오픈웨이트 모델이 어떤 형식으로 양자화 버전을 내놓느냐가 생태계에 영향을 주는 이유다.

정밀도가 성능으로 이어지는 두 개의 경로 #

비트를 줄이면 정확히 어디서 빨라지는가를 더 알아보자.

메모리 이동량 감소 #

BF16 가중치 70B: 140 GB
FP8  가중치 70B:  70 GB   → 읽는 시간 절반
FP4  가중치 70B:  35 GB   → 읽는 시간 4분의 1

**메모리 바운드 구간(디코드)**에서는 이 경로가 거의 전부고 계산량은 그대로인데 데이터 이동만 줄여서 그만큼 빨라진다.

텐서 코어 처리량 증가 #

낮은 정밀도일수록 텐서코어가 더 많은 연산을 동시에 처리하기에 대체로 비트폭이 절반이 되면 처리량의 2배이다.

B200 (dense 기준)
BF16:  2.25 PFLOPS
FP8 :  4.5  PFLOPS   (2배)
FP4 :  9    PFLOPS   (4배)

계산 바운드 구간 프리필, 대배치에서는 이 최적화를 볼 수 있다.

결론적으로 GB200 실측 기준 BF16 어텐션 + FP8 MoE에서 FP8 어텐션 + NVFP4 MoE로 바꾸게 된다면 프리필은 18,471 -? 26,156 디코드 9,087 → 13,386 토큰/초로 개선이 되는것을 볼 수 있는데.

프리필 개선 1.42배 디코드 개선 1.47배의 결과를 각각 텐서 코어 처리량 증가와 메모리 이동량 감소로 인해 얻게 된 결과이다.

정밀도도를 낮추면 메모리 이동량이 줄고 연산 능력도 함께 늘어난다.

그러면 메모리 바운드에서 계산 바운드로 넘어가는 지점도 옮겨갈 것 같지만 실제로는 움직이지는 않는다. 왤까?

루프라인 모델 #

지금까지 알아본 개념은 준비과정이었고 이 장에서 그 개념들이 하나의 계산 가능한 모델로 합쳐지게 된다.

루프라인을 익히고나면 프로파일러없이도 이 연산이 무엇에 막히는가를 판정할 수 있게 된다.

산술 강도 #

무엇이 막혀있는지 병목을 찾는 방법: 커널이 느리다고 쳐보자, 원인이 게산부족일까 데이터를 못 가져와서일까? 이 판정을 하지않으면 최적화 방향을 정할 수 없다. 계산문제인데 메모리를 최적화해봤자 아무 효과 없으니까.

바이트당 몇 번 계산하는가? 산술강도 AI(Arithmetic Intensity)를 정의해야한다.

수행한 연산 수 (FLOPs)
산술 강도 =  ─────────────────────────────
              이동한 데이터 양 (Bytes)

단위: FLOP/Byte

쉬운 비유로 재료창고를 주방으로 나르는것과 요리하는 것을 생각해보면 산술강도는 한 번 나른 재료로 몇인분을 만드는가다. 강도가 낮으면 재료 하나로 한입, 요리사가 놀고 강도가 높으면 짐꾼이 놀게 된다.

ex. 벡터 덧셈

C[i] = A[i] + B[i]   (FP32)

연산: 1 FLOP
이동: A 4바이트 읽기 + B 4바이트 읽기 + C 4바이트 쓰기 = 12바이트
AI = 1/12 ≈ 0.083 FLOP/Byte     ← 극단적으로 메모리 바운드

GEMV - 디코드 핵심 연산

배치 B, 가중치 행렬, 파라미터 수 P, 바이트/파라미터 b라고 치자

연산: 2 × P × B FLOP   (곱셈 1 + 덧셈 1)
이동: P × b 바이트     (가중치가 압도적. 활성값은 무시할 만함)

AI = 2PB / (Pb) = 2B / b

BF16 (b=2): AI = B
FP8  (b=1): AI = 2B
FP4  (b=0.5): AI = 4B

배치 1의 BF16 디코드는 AI = 1 FLOP/Byte

GEMM - 프리필 핵심 연산

M x K , K x N 행렬 곱

연산: 2 × M × N × K FLOP
이동: (MK + KN + MN) × b 바이트

M = N = K = 4096, BF16이면
연산 = 2 × 4096³ ≈ 137 GFLOP
이동 = 3 × 4096² × 2 ≈ 100 MB
AI ≈ 1,370 FLOP/Byte      ← 강하게 계산 바운드

같은 행렬곱인데 GEMV는 1이고 GEMM은 1,370으로 이 1,000배 차이가 프리필 디코드 성격을 가른다.

GEMM에서는 가중치 하나를 읽어 N번 재사용하고 GEMV(N=1)에서는 한 번 쓰고 버리니 재사용횟수가 곧 산술 강도가된다.

GEMM은 GEneral Matrix Vector multiplication 2차원 행렬, GEMV는 1차원 벡터로 General Matrix-Vector Multiplication을 뜻한다. GEMM이 프리필 GEMV가 디코딩에 쓰인다.

어텐션은 예외인데 여기서 가장 중요한 부분이 나온다.

디코드 단계의 어텐션 연산을 보면 시퀀스 길이 S, 배치 B라고 할때

연산: 각 요청이 자기 KV 캐시 S개와 내적 → 2 × S × d × B FLOP
이동: 각 요청의 KV 캐시를 읽어야 함 → S × d × b × B 바이트
                                        ↑ 배치에 비례해서 함께 증가!

AI = 2Sd·B / (Sd·b·B) = 2/b      ← B가 약분되어 사라짐

배치를 아무리 키워도 어텐션의 산술강도는 안오르고 이유는 명확한데 가중치는 모든 요청이 공유하지만 KV는 요청마다 다르기에 배치를 키우면 재사용되는것은 가중치고 KV캐시는 함께 증가하는것이다. 재사용을 하지 못하기 때문이다.

배치 크기 →     4      16      64     128
FFN(GEMV) AI:   4      16      64     128    ← 오름 ✅ 가중치
어텐션 AI:      1       1       1       1    ← 그대로 ❌

이 사실은 아주 많은것을 설명할 수 있는데

• 대배치에서도 어텐션이 병목인 이유: FFN은 계산 바운드로 넘어가지만 어텐션은 메모리 바운드로 남기 때문
• MLA와 GQA가 중요한 이유: KV 캐시를 줄이는 것이 어텐션의 이동 바이트를 줄이는 유일한 방법이라
• 왜 KV 캐시는 양자화에 효과적인가: b를 줄이면 AI가 오른다.
• 왜 프리필과 디코드를 분리하는가: 두 단계가 서로 다른 자원에 막혀있으므로 최적설정이 다르다.
• B300이 어텐션 연산 2배를 별도 강조하는이유: 어텐션이 남는 병목이라는걸 하드웨어 설계에서 반영한 것이다.

루프라인 그리기 #

아 루프라인은 컴퓨터 구조 AI 성능 분석을 시각적으로 보여주는 성능 분석 모델입니다.

어떤 커널도 두 가지를 동시에 넘을 수 없다.

1. 연산 천장: 하드웨어 최대 FLOPS
2. 메모리 천장: 대역폭 x 산술 강도

2 메모리 천장이 왜 그런지 확인해보자 대역폭이 3,350 GB/s이고 AI가 2FLOP/Byte라면 초당 3,350GB를 가져와서 바이트당 2번 계산하므로 최대 6,700 GFLOPS이다.

즉 달성가능한 성능은 대역폭 x AI를 넘을 수 없다.

성능(FLOPS)
   ▲
   │           ┌─────────────────────  ① 연산 천장 (수평)
   │          ╱
   │        ╱   ← ② 메모리 천장 (기울기 = 대역폭)
   │      ╱
   │    ╱
   │  ╱
   └────────┬──────────────────────→ 산술 강도 (FLOP/Byte)
         릿지 포인트
      ← 메모리 바운드 │ 계산 바운드 →

두 천장이 만나는 지점을 ridge pint 라고 한다.

릿지 포인트 = 최대 연산 성능 ÷ 메모리 대역폭   [FLOP/Byte]

이건 하드웨어의 속성이지 워크로드 속성은 아니고 같은 GPU와 모든 커널이 같은 릿지 포인트를 공유하고 각 커널은 자기 AI에 따라 좌우 어느쪽에 놓일지 결정된다.

하드웨어는 물리적 장치, 워크로드는 특정 작업을 처리하기위한 소모하는 연산량과 자원의 양

주요 gpu의 릿지 포인트

A100  BF16:  312 TFLOPS / 2.04 TB/s ≈ 153 FLOP/Byte
H100  BF16:  989 TFLOPS / 3.35 TB/s ≈ 295
H200  BF16:  989 TFLOPS / 4.8  TB/s ≈ 206   ← 대역폭만 늘어 릿지가 내려감
B200  BF16: 2250 TFLOPS / 7.7  TB/s ≈ 292
B200  FP4 : 9000 TFLOPS / 7.7  TB/s ≈ 1169
MI355X BF16: 2500 TFLOPS / 8.0 TB/s ≈ 313
MI355X MXFP4:10100 TFLOPS / 8.0 TB/s ≈ 1263

H100, 200 비교해보라 연산 능력은 같은데 대역폭이 늘어서 릿지 포인트가 295에서 206으로 내려갔다.

릿지가 내려간다는 것은 더 낮은 산술강도에서 연산 천장에 닿는다라는 뜻으로 메모리 바운드 워크로드에게 유리한 기계다, H200이 디코드 중심 서빙에서 H100보다 체감이 큰 이유다/

여기서 얻을 수있는 일반 원칙은 프리필/학습 처럼 산술 강도가 높은 워크로드는 연산 천장이 높은 기계를 원하고, 디코드처럼 산술 강도가 낮은 워크로드는 릿지포인트가 낮은 즉 대역폭이 상대적으로 큰 기계를 원한다. 총 FLOPS가 크다고 디코드가 빨라지는 것은 아니다.

임계 배치 크기 #

배치 값을 얼마로 설정해야할까.

앞선 서빙노트에서 배치를 키우려면 처리량이 오르지만 지연이 나빠진다고 했다.

그런데 처리량이 오르는것도 무한하지는 않다. 어느 지점부터 계산 바운드가 되어 이득이 급격히 줄어든다.

그 지점을 계산할 수 있어야한다.

전환점으로 식 구하기 #

디코드 FFN의 산술 강도가 AI = 2B/b 였다. 이게 릿지 포인트와 같아지는 배치 크기가 전환점이다.

2B_crit / b = 릿지 포인트
B_crit = 릿지 포인트 × b / 2

H100 BF16에 넣어 보면:

B_crit = 295 × 2 / 2 = 295

배치 295 근처에서 FFN이 메모리 바운드에서 계산 바운드로 넘어간다.

양자화와 배치의 임계값을 바꿀 수 있지는 않다

같은 H100에서 FP8로 계산해보면

FP8 릿지 포인트 = 1,979 TFLOPS / 3.35 TB/s ≈ 591
b = 1

B_crit = 591 × 1 / 2 = 295      ← BF16과 똑같다.

비트폭은 절반으로 줄어들었으면 이동 바이트가 절반이되어 AI가 2배가 되고 동시에 텐서코어 처리량도 2배가되어 릿지포인트도 2배가 되니 둘이 정확히 상쇄가 되어 배치가 늘진 않는다.

양자화 효과 정리

바뀌는 것:   같은 배치에서의 절대 속도 (2배, 4배 빨라짐) ✅
             메모리 사용량 (절반, 4분의 1) ✅
바뀌지 않는 것: 메모리 바운드 → 계산 바운드 전환점 ❌

FP4로 바꿨으니 이제 배치 더 키워야지는 연산 측면에서는 틀렸고 다만 메모리가 남아서 배치를 키울 수 있게 되는 것은 맞다.

두 효과를 구분해야한다. 정확히 말하면 양자화 KV 캐시 공간을 확보해서 배치 상한을 올려주는거지 계산 효율의 전환점 자체는 옮기지 않는 것이다.

하드웨어별 임계 배치는

A100:  약 153
H100:  약 295
H200:  약 206
B200:  약 292
MI355X: 약 313

• 동시 요청이 임계 배치보다 훨씬 작다면 -> 메모리 바운드 구간이고 양자화와 투기적 디코딩이 잘 통한다 배치를 더 키우는것도 거의 공짜다.
• 동시요청이 임계 배치 근처거나 그 이상이라면? -> 계산 바운드로 넘어가니 투기적 디코딩 이득이 사라지고 배치 32를 넘으면 EAGLE-3를 끄라는 권고가 있는데 이론상 임계값보다 낮은 이유는 투기적 디코딩이 스텝당 여러 토큰을 검증하느라 실효배치를 몇 배로 부풀리기 때문이다. 배치 32에 5토큰 투기면 실효 160이니 여기에 어텐션 병목과 드래프트 모델 비용이 더해져 최악

최적화 루프라인 #

각 최적화 기법이 그래프 위에서 무엇을 하는지 정리해보겠다.

성능
  ▲
  │              ┌────────── 연산 천장
  │         ╱────┘
  │    ╱───┘
  │ ╱
  └──────────────────────→ AI

천장에서 한참 아래에 있는데 GPU를 바꿀필욘 없고 최적화 생각을 해야한다.

GPU를 더 사면 해결된다는 오해다 애초에 그 gpu에도 공간이 많은데 잘 쓰고있지 못하는 것 뿐

아래 절차를 꼭 밟자

1. 이 연산의 AI를 계산한다
        ↓
2. 하드웨어의 릿지 포인트와 비교한다
        ↓
   AI < 릿지 → 메모리 바운드
        → 이동 바이트를 줄여라 (양자화, 타일링, 캐시 재사용)
        → 연산을 더 해도 공짜다 (투기적 디코딩)
   AI > 릿지 → 계산 바운드
        → 텐서 코어를 제대로 쓰고 있는지 확인해라
        → 저정밀 포맷을 검토해라
        ↓
3. 실측 성능이 해당 천장의 몇 %인지 확인한다
        ↓
   70% 이상 → 이 방향은 거의 다 왔다. 다른 병목을 찾아라
   30% 미만 → 병렬성 부족(점유율), 병합 실패, 뱅크 충돌을 의심하라

루프라인 한계 #

이 모델은 강력하지만 단순화된 모델이다 다음을 무시하기 때문이다.

1. 캐시 효과: 실제로 L2에 데이터가 남아 재사용되므로 이동 바이트가 계산보다 적을 수 있다 그래서 실측이 이론 예측보다 좋게 나오는 경우가 있다.
2. 지연과 병렬설: 루프라인은 충분한 병렬성이 있어 지연이 완전히 숨겨진다고 가정한다 점유율이 낮으면 두 천장 어디에도 못 닿는다. 이 경우 루프라인 위의 점은 천장 한참 아래에 놓이고 원인은 AI가 아니라 병렬성이다.
3. 명령 조합: 전용 유닛(텐서 코어)를 쓰는지 일반 유닛을 쓰는지에 따라 실효 천장이 다르기에 WGMMA (이론치의 95퍼)와 하위 호환 mma(62,9%)의 차이를 알아두자
4. 커널 실행 오버헤드: 작은 커널이 많으면 발행 비용이 지배한다 CUDA 그래프가 필요한 이유

그래소 루프라인은 어디를 볼지 결정하는 도구이지 정확한 예측 도구는 아니기에 계산으로 방향을 잡고 프로파일러로 확인하는 순서가 맞다

릿지 포인트 성능 분석 모델인 루프라인 모델에서 하드웨어 최대 연산 성능 Peak FLOPS와 최대 메모리 대역폭 Peak Memory Bandwidth 한계선이서로 만나는 전환점(꼭짓점)이다. Peak FLOPS / Peak Memory Bandwidth

2026년 하드웨어 지도 #

앞 장의 도구를 실제 하드웨어에 적용한다. 이 장의 숫자들은 빠르게 낡으므로 숫자 자체보다 어느축이 어떻게 움직였는가를 주목해서 보자

세대별 요약표 #

6.1 세대별 요약표 #

PF와 TF는 AI의 연산성능을 나타내는 단위로 TF는 TeraFLOPS로 초당 1조번 연산 PF는 PetaOPS로 초당 1000조 연산을 뜻함

값은 벤더 공개 자료 기준이며 SKU, 냉각 방식에 따라 달라진다 특히 B200 180GB(SXM6 실측 기준)과 192GB 표기가 혼용된다.

세대별로 바뀐 값 #

Hopeer (H100 -> H200) 대역폭만 늘린 리프레서 #

H200은 H100과 연산 유닛이 동일하다. 메모리만 HBM3에서 HBM3e로 바꿔 용량 80 -> 141GB, 대역폭 3.35 -> 4.8TB/s가 되었다

루프라인으로 보면 연산 천장은 그대로인데 메모리의 천장 기울기만 가팔라졌다. 릿지 포인트가 295 에서 206으로 내려간다.

이 의미는 학습이나 프리필 중심이면 H100 대비 이득이 크지 않지만 디코드 중심 서빙이면 이득이 직접적이다. 대역폭이 43% 늘었으니 메모리 바운드 구간에서 그만큼 빨라지고 용량이 76% 늘었으니 KV 캐시에 쓸 공간도 그만큼 늘어난다. 대신 그만큼 컴퓨트할게 많아질거라 컴퓨터와 트레이드오프다.

Blackwell (B200) 세 가지의 변화 #

1. 듀얼 다이 구조: 레티클 한계에 도달한 다이 두 개를 붙였는데, 총 2,080억 트랜지스터, 두 다이는 NV-HBI로 10TB/s로 연결되어 소프트웨어에는 단일 GPU로 보인다.
2. FP4 네이티브 지원: 5세대 텐서 코어와 2세대 트랜스포머 엔진, 앞 장에서 본 NVFP4 마이크로스케일링을 하드웨어가 직접 처리한다.
3. tcgen05 프로그래밍 모델: TMEM, 단일 스레드 발행, CTA 페어

복스비 tcgen05는 blackwell 아키텍처에 도입된 텐서코어 전용 PTX 명령어 세트 프로그래밍 모델이고 TMEM(Tensor Memory)는 연산 결과를 레지스터가 아닌 텐서코어 전용 하드웨어 캐시 메모리에 직접 기록하고 관리하며 단일 스레드 발행은 이전세대 명령어 mma.sync, wgmma와 달리 단일 스레드 수준에서 텐서코어 행렬 연산 명령 내리기가 가능해지는것 CTA페어라는 뜻은 동일한SM에서 두 CTA(스레드 블록)이 협력해 텐서 코어 명령어를 실행해 데이터를 공유할 수 있게 해주는 병렬 컴퓨팅 기능이다.

루프라인으로 보면 BF16은 릿지 292로 H100(295)와 거의 같다. 연산과 대역폭이 둘다 비슷한 배수로 올랐기 때문이다, FP4를 쓰면 1169로 릿지가 치솟는다, 즉 FP4는 연산능력을 크게 올리지만, 그만큼 그 능력을 다쓰려면 산술 강도도 높아야한다.

Blackwell Ultra (B300) - 추론에 맞춘 재조정 #

B300은 B200에서 방향을 틀었다

• 메모리 12-high 스택으로 288GB(B200대비 50퍼 증가)
• FP4 dense 15 PFLOPS (B200 대비 1.5배)
• FP64 성능을 의도적으로 낮춤
• 어텐션 연산 능력을 별도 강화

추론 편향, 메모리 밀집 프로세서로 설계되었고 FP64를 희생해서 저정밀 효율과 어텐션, 메모리 대역폭에 실리콘을 몰아준 것이다.

왜 어텐션을 따로 강화했냐면 앞 장에서 본 대로 어텐션 배치를 키워도 산술 강도가 오르지 않아 끝까지 병목으로 남는다. 하드웨어 설계자가 루프라인 분석 결과를 그대로 반영한 사례다.

GB300 NVL72랙 기준으로 72장의 20.7TB의 HBM3e를 이루고 NVLink 도메인 대역폭 130TB/s, FP4 1.1 엑사 플롭스, 소비 전력 약 120kW이데.

Rubin (VR200) #

2026년 2월에 첫 샘플이 출하되었고 하반기부터 양산 물량이 클라우드 사업자들에게 들어가는 일정이다.

• Rubin GPU 하나가 FP4 50 PFLOPS (B300의 3.3배)
• HBM4 288GB, 목표 대역폭 22 TB/s (실제 초기 물량은 약 20TB/s로 조정 보도)
• 레티클 크기 다이 2개 + HBM4 스택 8개
• Grace를 대체하는 Vera CPU (88 Armv9.2 Olympus 코어)
• NVLink 6, 랙 인터커넥트 260 TB/s
• 랙 전력 약 190230 kW (Blackwell 120130 kW에서 크게 상승)

루프라인으로보면 FP4 릿지 포인트가 50 PFLOPS = 50,000 TFLOPS, 20 TB/s = 20,000 GB/s로 릿지는 2,500 FLOP/Byte다. B200의 FP4 릿지 1,169에서 두배 이상 ㄷ

이게 뜻하는바는 연산 능력이 대역폭보다 빠르게 늘었다 이고 메모리 바운드 워크로드 입장에서는 이론 FLOPS의 증가분을다 쓰기 어려워진다 디코드 성능은 주로 대역폭 증가분 (8 -> 20 TB/s 2.5배) 만큼 개선되고 FP4 FLOPS 3.3배는 프리필과 대배치에 쓰이게된다.

공급 현실 - HBM4는 2026년 물량이 사실상 소진이고 신규 생산능력이 의미있게 늘어나는 시점은 2027년으로 전망된다. 전력 측면에서도 190230kW는 800VDC 급전과 전면 액체 냉각을 요구해 하드웨어 로드맵과 실제 조달 가능성은 다른 문제고 20262027년에는 Blackwell이 여전히 주력일 가능성이 크다

AMD CNDA 4 (MI355X) #

• 256 컴퓨트 유닛, 1,024 메트릭스 코어 (AMD는 텐서코어대신 메트릭스 코어라고함)
• 288GB HBM3e, 8 TB/s 8,192 비트 인터페이스
• 64MB Infinity Cache (L3에 해당하고 NVIDIA에 L2 보다 큼)
• MXFP4/MXFP6/MXFP8 네이티브 지원으로 OCP 표준 계열이라 NVFP4가 아니다.
• 8개 컴퓨트 다이 + Infinity Fabric 4

포지셔닝: B200과 비교하면 메모리 용량에서 앞선다 288 vs 196이라, FP16기준 192GB를 넘는 모델 (대략 110B 이상) 이 MI335X 한 장에 들어가지만 B200은 두 장 이상이 필요하다 샤드 수가 줄면 통신 오버헤드도 줄어든다.

주의할점은 AMD가 홍보하는 MI300X 대비 35배 추론 향상같은 수치는 MI300X의 FP8을 기준선으로 MI355X의 FP4를 비교한거라 FP4 vs FP4 비교가 아니다. 벤더 수치를 읽을때 항상 기준선 확인을 해야한다.

워프 크기가 64인데 NVIDIA는 32로 전제하고 짠 커널이라 그대로 옮겨지지 않는다 ROCm 생태계가 성숙했다해도 커널 수준에서는 별도 작업이 필요하다.

하드웨어 선택기준 - 루프라인으로 #

스펙표를 보고 우리는 어느 워크로드에 어느쪽이 나은가를 확인하는 절차다.

1단계 — 워크로드의 지배적 산술 강도를 구한다
        디코드 중심?     → AI가 낮음 (1~수십)
        프리필/학습 중심? → AI가 높음 (수백~수천)
        혼합?            → 토큰 비율로 가중평균

2단계 — 후보 GPU의 릿지 포인트를 계산한다
        릿지 = 해당 정밀도 peak FLOPS ÷ 대역폭

3단계 — 판정
        AI ≪ 릿지  → 대역폭이 성능을 결정. 대역폭당 가격을 비교
        AI ≫ 릿지  → 연산이 성능을 결정. FLOPS당 가격을 비교

4단계 — 용량 제약을 확인한다
        모델 가중치 + 목표 동시 사용자의 KV 캐시가 들어가는가?
        (이 계산은 앞선 서빙 노트의 KV 캐시 공식을 사용)

5단계 — 소프트웨어 지원을 확인한다
        쓰려는 엔진이 그 하드웨어에서 최적 커널 경로를 타는가?
        (SM100인가 SM120인가, ROCm 지원 성숙도는 어떤가)

5단계 무조건 확인하자 이론 성능이 좋아도 커널이 없으면 그 성능을 못쓰기 때문이다.