GPU 커널부터 새로 짜서 GLM-5.2 속도 10배 올리기
GLM-5.2 753B 모델을 SASS 수준에서 최적화해서 초당 600~1000 토큰 생성하는법
용어설명
SASS(Streaming Assembler) 는 엔비디아 그래픽카드(GPU)에서 직접 실행되는 가장 낮은 수준의 기계어 코드입니다. 일반적으로 C++로 작성된 코드가 컴파일되면 SASS로 변환되며, GPU는 이 코드를 바탕으로 명령을 수행합니다.
tcgen05은 NVIDIA 블랙웰(Blackwell) 아키텍처(SM100) GPU에 도입된 5세대 텐서 코어(Tensor Core)의 핵심 행렬 곱셈(GEMM) 명령어입니다. 이전 세대의 워프 그룹(warp-group) 명령어 대신 단일 스레드가 명령을 발행하고, 결과값을 레지스터가 아닌 텐서 메모리(TMEM)에 저장하는 것이 특징입니다.
TMEM(Tensor Memory)은 NVIDIA 블랙웰(Blackwell) 아키텍처부터 도입된 텐서 코어 전용의 온칩(On-chip) 메모리입니다.
SFU(Special Function Unit)는 NVIDIA GPU등에 포함된 복잡한 수학 연산을 전담하는 보조 연산 장치입니다. 주 코어가 행렬 계산에 집중할 때, 삼각함수나 지수함수 같은 '비선형(non-linear)' 방정식을 빠르게 처리합니다. B300 Blackwell Ultra 부터는 직전 제품 대비 2배로 SFU 용량이 증가했습니다.
현대의 GPU는 아주 높은 연산력을 가지고 있습니다.
예를들어 사이오닉이 보유하고 있는 NVIDIA Blackwell Ultra B300은 1장은 15PFLOPS FP4 연산 성능을 가지고 있습니다.
일반적으로 쉽게 감이오지 않는 숫자 단위인데 단 한장의 카드가 1초에 1경 5000조 번 FP4 부동소수점 연산을 한다고 보면됩니다.
B300 한 장의 HBM3e 피크 대역폭은 ~8 TB/s입니다. 그리고 LLM에서 토큰 1개를 생성의 연산의 본질은
▎ "메모리 가중치 전부를 HBM에서 GPU 내부 SM으로 한 번 흘려보내면서, 그 위에 아주 적은 연산을 하는 것"
메모리를 전송량을 다소 줄이는 MoE 등 구조를 좀더 개선하는 방식이 있지만 구조상 크게 변함은 없습니다.
GPU 외의 다른 구조의 연산 가속기가 총 연산력이 더 적더라도 토큰 생성에서 GPU 보다 빠른 속도를 냈다는 많은 사례들은 GPU는 LLM 관점에서 연산력은 충분히 남고 메모리의 속도가 병목이기 때문입니다.
예를 들어보면
GLM-5.2는 753B 모델이지만 MoE 라 토큰당 top-8 전문가만 활성화됩니다. TP8로 8장에 나누면:
랭크(GPU 1장)당 토큰당 가중치 트래픽 ≈ 6.65 GB
÷ 피크
8 TB/s = 0.83 ms ← 물리적 한계
÷ 현실적 달성률 60%
4.8 TB/s ≈ 1.1 ms ← 실제 거의 최고 성능
즉 물리적인 한계상 모델의 배치 1 포워드 한번은 1ms 보다 빨라질 수 없습니다.
그런데 실측 포워드는 9ms — 대역폭 한계보다 약 8~11배 느립니다.
특수한 반도체를 사용하는것이 아닌 대부분의 프론티어 모델의 추론 속도가 초당 50~100 토큰 사이에 위치 하는것이 위와 같은 메모리 대역폭의 한계에 기인합니다.
실제로 GLM-5.2의 GPU 내부 실행 구조를 프로파일 해보면 포워드 1회 전체의 평균 대역폭에서
•
GEMM들의 실효 대역폭: ~2.1 TB/s = 피크의 26%
8 TB/s 짜리 HBM3e가 있지만 실제는 20% 정도 효율만 내고 있는 셈 입니다.
이유는 여러가지가 있는데 크게 다음 4가지 입니다
1. 낮은 SM 점유율(Occupancy)
실제 메모리 대역폭은 단순히 HBM의 스펙만으로 결정되지 않습니다. 얼마나 많은 SM이 동시에 메모리에 접근하느냐가 더 중요합니다.
GLM-5.2의 batch-1 gate_up GEMM에서는 CTA가 4개만 생성되어 B300의 148개의 SM 중 극히 일부만 HBM에 요청을 보내게 됩니다. 그 결과 이론적으로 1.5µs면 모두 읽어올 수 있는 가중치를 실제로는 53µs 동안 로드하게 됩니다.
여기에 GLM-5.2의 실시간 FP8 활성화 양자화(live activation quantization) 를 GEMM 직전에 단순히 융합하면 문제가 더욱 심각해집니다. SM103의 cp.async 경로에서는 양자화가 가중치 로드와 겹쳐 실행되지 못하고 직렬화되어, 실행 시간이 53µs에서 127µs 까지 증가했습니다.
사이오닉은 이 병목이 연산량이 아니라 SM 활용률에서 발생한다는 점에 주목했습니다.
SM 구조에 맞춰 Split-K를 적용하여 K축을 16개로 분할하고 병렬성을 높인 결과, 실행 시간은 127µs에서 37µs 까지 감소했습니다. 이는 실시간 양자화를 포함한 상태에서도 기존 비융합 커널(53µs) 보다 더 빠른 결과입니다.
이 최적화 덕분에 GLM-5.2의 Live Quantization을 추론 과정에서 직접 수행하면서도 성능 저하 없이 실행할 수 있게 되었으며, 이러한 SM 구조를 고려한 세밀한 최적화는 범용 GEMM 라이브러리인 DeepGEMM만으로는 구현하기 어려운 영역입니다.
2. 파이프라인 직렬화(Pipeline Serialization)
기존 HMMA 기반 커널은 K 48 반복마다 스케일 연산이 배리어 역할을 수행하면서, 다음 가중치를 미리 가져오는 비동기 파이프라인이 계속 끊어집니다.
그 결과 실제 메모리 대역폭은 0.49TB/s(6%) 수준까지 떨어집니다.
사이오닉은 Blackwell의 TMEM(Tensor Memory) 기반 tcgen05 활용하여 스케일 연산을 텐서코어 내부 메모리로 격리하였고, 로드와 연산이 중첩 파이프라이닝 되도록 만들었습니다.
그 결과 대역폭은 1.53TB/s(19%) 로 약 3.1배 향상되었습니다.
3. 레지스터 압박(Register Pressure)
누산기를 레지스터에 유지하면 레지스터 사용량이 증가하면서 SM당 동시에 실행 가능한 CTA 수가 감소합니다.
결국 메모리 요청 자체가 줄어들어 HBM 대역폭을 충분히 활용하지 못하게 됩니다.
TMEM 기반 누산 방식은 이러한 레지스터 압박을 제거하여 더 높은 메모리 수준 병렬성(MLP)을 유지할 수 있습니다.
4. 지나치게 많은 작은 커널
GLM-5.2의 포워드는 수백 개의 작은 커널로 구성되어 있습니다.
작은 커널은 실행될 때마다 파이프라인을 새로 채우고 비워야 하기 때문에, 충분한 메모리 대역폭에 도달하기도 전에 종료됩니다.
또한 활성화 데이터를 커널 사이마다 반복적으로 HBM으로 저장하고 다시 읽어오는 불필요한 트래픽도 발생합니다.
사이오닉은 PQ-GEMM(Quantization + GEMM Fusion) 을 통해 실시간 양자화를 GEMM 내부로 융합하여 이러한 왕복 메모리 접근 자체를 제거했습니다.
또한 충분한 프로파일링과 구조 검토를 통하여 지나친 메가 커널의 재구현은 지양하였습니다.
중간결과
사이오닉은 GLM-5.2의 Live Quantization을 추론 과정에서 직접 수행할 수 있도록 구현하면서도, SM 구조를 고려한 Split-K, TMEM 기반 Tensor Core 활용, Pipeline Fusion 등 B300 하드웨어 특화된 최적화를 적용했습니다.
이러한 접근은 단순히 GEMM 커널의 FLOPS를 높이는 것이 아니라, GPU의 실제 병목인 메모리 대역폭 활용률을 근본적으로 개선하는 데 초점을 맞춘 것입니다.
범용 GEMM 라이브러리인 DeepGEMM은 다양한 모델에서 동작하는 범용성을 목표로 하지만, GLM-5.2의 Live Quantization과 같이 모델 구조와 SM 아키텍처를 함께 고려해야 하는 세밀한 최적화는 적용하기 어렵습니다.
하지만 초당 600 토큰 이상의 아주 빠른 생성 성능을 내기에는 단순히 메모리 전송과 전용 레지스터 활용등 명령어 작은 최적화로는 부족합니다.
5. MTP — HBM의 물리적 한계를 알고리즘으로 우회
앞에서 소개한 Split-K, PQ-GEMM, TMEM, Pipeline Fusion과 같은 최적화는 모두 같은 목표를 가지고 있습니다.
가중치를 더 빠르게 읽고, 불필요한 메모리 왕복을 줄이며, 더 많은 SM이 동시에 일하도록 만드는 것
이러한 최적화 덕분에 GLM-5.2는 실시간 Live Quantization까지 포함하면서도 포워드 시간을 크게 줄일 수 있었습니다.
하지만 이것만으로는 초당 600 token/s 이상의 성능에 도달하기는 어렵습니다.
이유는 단순합니다.
아무리 GPU 커널을 잘 만들어도 HBM의 물리적인 대역폭 자체는 바꿀 수 없기 때문입니다.
앞서 말한것 처럼 이론적으로 토큰 하나의 생성에는 대략 9ms 수준의 메모리 속도 제한이 있습니다.
GLM-5.2와 같은 753B 규모의 MoE 모델에서는 토큰 하나를 생성할 때마다 활성화된 전문가의 가중치를 HBM에서 한 번 읽어야 합니다.
앞에서 계산한 것처럼 batch-1 기준으로 랭크당 약 6.65GB의 가중치가 토큰 하나마다 전송됩니다.
커널 최적화는 이 메모리 전송을 더 빠르게 만드는 기술입니다.
하지만 결국 메모리 전송 한번은 반드시 수행해야 합니다.
여기서 MTP(Multi-Token Prediction)는 접근 방식을 크게 바꿉니다.
"가중치를 한 번 읽어야 한다면, 그 한 번의 읽기로 토큰 하나가 아니라 여러 개를 확정하면 되지 않을까?"
이것이 MTP의 핵심입니다.
5. 포워드를 줄이는 것이 아니라, 한번에 여러 토큰을 생성한다.
일반적인 디코딩은 토큰 하나를 생성하고, 그 토큰을 다시 입력으로 넣어 다음 토큰을 생성합니다.
즉 토큰 하나마다 메인 모델의 포워드가 한 번씩 수행됩니다.
반면 MTP에서는 작은 Draft 모델이 먼저 미래 토큰 후보를 여러 개 생성합니다.
메인 모델은 이 후보들을 한 번의 포워드에서 동시에 검증합니다.
앞쪽 후보들이 메인 모델의 결과와 일치하면, 여러 개의 토큰이 한 번에 확정됩니다.
예를 들어 평균 예측 길이(Acceptance Length)가 3.5라면
기존 방식은
가중치 1회 전송 → 토큰 1개 확정
이지만,
MTP에서는
가중치 1회 전송 → 토큰 3.5개 확정
으로 바뀝니다.
예측 길이가 10이라면
가중치 1회 스트리밍 → 토큰 10개 확정
이 됩니다.
앞에서 계산했던 GLM-5.2의 토큰당 가중치 트래픽을 다시 계산하면 다음과 같습니다.
평균 예측 길이 | 토큰당 가중치 트래픽 |
MTP 미사용 | 약 6.65GB |
3.5 | 약 1.9GB |
10 | 약 0.67GB |
6. 포워드를 줄이는 것이 아니라, 한번에 여러 토큰을 생성한다.
배치 추론(batch inference)이 빠른 이유도 본질적으로 같습니다.
64개의 요청을 동시에 처리하면,
가중치 한 번을 읽어서 64개의 토큰 계산이 함께 사용합니다.
그래서 batch가 커질수록 토큰당 가중치 비용이 급격히 감소합니다.
MTP는 이 효과를 batch-1에서도 만들어냅니다.
여러 사용자의 요청을 기다려 배치를 키우는 대신,
한 사용자의 미래 토큰들을 동시에 검증하여 가상의 Batch를 만드는 것입니다.
그래서 latency를 크게 희생하지 않으면서 throughput은 batch가 커진 것처럼 증가할 수 있습니다.
실제로 GLM-5.2에서도 기본 Greedy Decoding은 약 100 tok/s 수준이었지만,
MTP를 적용한 이후에는 약 600 tok/s까지 크게 향상되었습니다.
7. 하지만 MTP를 적용한다고 끝이 아니다.
MTP와 Speculative Decoding은 새로운 알고리즘이 아닙니다.
이미 꽤 오래전인 2018년, 2023년 여러 논문으로 제안된 방식입니다.
문제는 MTP가 실제로 성능을 내도록 만드는 것입니다.
평균 예측 길이는 Draft 모델의 품질과 추론 경로의 아주 작은 변화에도 크게 흔들립니다.
예를 들어 활성화 양자화 때문에 logits 분포가 조금만 달라져도
실제 평균 예측 길이가
7.13 → 4.47
까지 감소할 수 있었습니다.
그러면 양자화로 줄어든 메모리 이득보다
MTP의 예측률이 크게 감소하면서
오히려 전체 throughput이 떨어지는 역전 현상이 발생합니다.
특히 하드웨어 최적화를 포함한 복잡한 구현 변경은 아주 작지만 누적된 양자화 계산오차를 발생시킬 수 있습니다.
즉 양자화와 MTP는 독립적인 최적화가 아닙니다. 속도를 개선하면서 동시에 수치적 오차까지 지켜내야 실제 성능이 증가합니다.
양자화는
토큰당 읽어야 하는 바이트 수를 줄이고
MTP는
한 번 읽은 바이트를 여러 토큰이 함께 사용합니다.
둘이 함께 동작해야 비로소 곱으로 성능이 증가합니다.
8. 사이오닉의 한일 — MTP 프로덕션 성능으로 만든 아래층
MTP를 발명한 것은 사이오닉이 아닙니다.
여러 논문에서 MTP의 구조와 기본 아이디어는 이미 공개 되어있고 GLM-5.2은 MTP를 기본적으로 사용가능한 구조를 가지고 있습니다.
사이오닉이 한일은
알려진 알고리즘이 B300과 GLM-5.2 753B MoE에서 실제 성능을 내도록 만든 그 아래층 전체
입니다.
주요 내용은 크게 3 가지가 있습니다.
a) MTP의 대한 실제 최적화
알고리즘은 공개되어 있지만,
언제 빨라지고 언제 느려지는지는 공개되어 있지 않습니다.
사이오닉은 GLM-5.2를 대상으로
•
temperature와 threshold에 따른 예측 길이 변화
•
코드·자연어·랜덤 워크로드별 Acceptance Length
•
양자화 수준와 MTP 예측률의 상관관계
•
품질을 유지하는 최적 파라미터
를 모두 실측했습니다.
덕분에 "빠른 설정"이 아니라
품질을 유지하면서 가장 빠른 레시피를 찾을 수 있었습니다.
b) 깊은 MTP가 만드는 새로운 병목을 제거
MTP가 깊어질수록
GPU보다 만큼 CPU와의 동기화 비용이 새로운 병목이 됩니다.
사이오닉은 서빙 기반인 SGLang에 20여 개 이상의 패치를 적용하여
•
CPU H2D 동기화 제거
•
CUDA Graph 적용
•
MoE AllReduce Fusion
•
Sampling Fusion
•
LMHead FP8
•
Custom Kernel Injection
등을 수행했습니다.
108 tok/s에서 600 tok/s까지 올라간 상당 부분은
이 런타임 계층의 병목 제거에서 나왔습니다.
c) 커널 수준 최적화
GLM-5.2의 Live Quantization은
단순히 GEMM 앞에 붙이면 오히려 53µs에서 127µs까지 느려졌습니다.
사이오닉은 SM103 구조에 맞춰 SM 구조에 최적화된 메모리 접근을 적용하여
라이브 양자화를 포함한 퓨전 커널을 127µs를 37µs까지 줄였습니다.
덕분에 GLM-5.2의 Live Quantization을
실제 추론에서 성능 저하 없이 사용할 수 있게 되었습니다.
또한
•
PQ-GEMM
•
tcgen05
•
TMEM
•
Blackwell Ultra Dispatch 수정
등을 통해 DeepGEMM과 같은 범용 GEMM 라이브러리만으로는 구현하기 어려운 모델 특화 최적화를 추가했습니다.
마치며
핀볼 게임을 단 10초만에 만들어내는 초당 600 토큰 GLM-5.2 최적화의 핵심은 GPU FLOPS를 높이는 것이 아닙니다.
핵심은
가중치 메모리 전송을 1회를 얼마나 최적화 하고 만들고, 그 한 번의 전송에서 몇 개의 토큰이 함께 사용할 것인가.
입니다.
커널 최적화는 포워드 1회의 비용을 줄입니다.
MTP는 포워드 1회를 여러 토큰에 나누어 사용합니다.
사이오닉은 이 두 축을 동시에 최적화하여, B300 기반 GLM-5.2에서 batch-1 환경에서도 600~1000 token/s 수준의 실제 추론 성능을 구현했습니다.
알고리즘과 실제 프로덕션 성능 사이에는 GPU 커널, 메모리 계층, 런타임, CPU 동기화, 예측률 특성화와 같은 수많은 최적화 계층이 존재합니다.
사이오닉의 주 연구 분야는 바로 그 아래층을 하나씩 제거하여, 알려진 알고리즘을 실제 프로덕션 성능으로 연결하는 것입니다.
단 하나의 GPU와, 단 하나의 모델만을 위한 커널 최적화와 조합으로 전용 반도체 수준의 추론 속도를 낼 수 있다는 부분이 중요한 포인트 입니다.
실제 서빙환경에서는 kv cache의 히트율과 세션등 여러 라우팅 기술들이 추가적으로 요구 됩니다. 실제 상단 라우팅 캐싱 환경에서 GLM-5.2를 고려한 프로덕션 환경의 MLOps 관점의 내용도 연이어 공유 드리겠습니다.