LLM은 빠른 멀티 GPU 커널을 작성할 수 있을까?

URL: https://www.youtube.com/watch?v=pOvWgX7IJsc 날짜: 2026-08-28 원본 발행일: 2026-08-27 채널: aiDotEngineer 원문 제목: Can LLMs Write Fast Multi-GPU Kernels? — Simran Arora, Together AI

---

📌 핵심 질문 / 이 영상이 다루는 핵심 논점

LLM은 단일 GPU 코드와 일반적인 CUDA 문법을 넘어, 하드웨어 토폴로지·통신 순서·데이터 분할·스케줄링·전송 방식의 트레이드오프를 추론해 빠른 멀티 GPU 커널(Multi-GPU Kernel)을 새로 작성할 수 있는가? ==현재 최첨단 모델은 커널을 컴파일하고 일부 익숙한 패턴을 재현할 수 있지만, 멀티 GPU 성능을 결정하는 근본적인 설계 트레이드오프를 안정적으로 추론하지 못한다.==

• GPU 내부 커널을 개선해 병목이 멀티 GPU 통신으로 이동했으며, 통신을 커널 안에서 세밀하게 제어해야 한다.
• ParallelKittens는 소수의 원시 연산(Primitives)과 패턴으로 멀티 GPU 커널의 성능·버퍼링·동기화를 직접 제어하는 추상화다.
• Parallel Kernel Bench는 실제 AI 워크로드에서 나온 87개 문제로 LLM의 멀티 GPU 커널 생성 능력을 평가한다.
• 최상의 제로샷 모델은 87개 중 28개를 맞혔고 22개만 PyTorch+NCCL 기준보다 빨랐으며, 샘플을 늘려도 빠른 해의 비율은 약 31%에서 정체됐다.

Simran Arora는 Together AI의 수석 과학자이자 Frontier Performance Research 팀 리더이며, Stanford Hazy Research Lab에서 Chris Ré와 박사 학위를 받은 뒤 Caltech 교수로 부임할 예정이다. Arora의 연구팀은 최신 AI 하드웨어에서 가능한 성능을 끌어내기 위한 시스템·프레임워크·알고리즘을 만들고, 그 경험을 바탕으로 멀티 GPU 커널 개발의 구조적 원리와 LLM의 한계를 연결한다.

1. 멀티 GPU 통신이 핵심 병목이 된 배경

1.1. 발표의 출발점과 연구 배경

1. Simran Arora의 역할

   1. Together AI의 연구: Arora는 Together AI 수석 과학자로서 Frontier Performance Research 팀을 이끌며 시스템, 프레임워크, 알고리즘으로 현대 AI 하드웨어 성능을 최대한 끌어내는 일을 한다.
   2. 학문적 배경: Stanford의 Hazy Research Lab에서 Chris Ré와 박사 과정을 수행했고, Caltech의 차기 교수로 부임할 예정이다.

2. 현장의 짧은 농담과 문제 제기

   1. 예상보다 시끄러운 장소: Arora는 인사를 건넨 직후 장소가 시끄럽고 이런 상황을 예상하지 못했다며 사과한다.
   2. 핵심 질문: 단일 GPU 커널을 조금 고치는 수준이 아니라, 여러 GPU의 통신과 계산을 함께 설계하는 작업을 LLM이 수행할 수 있는지가 질문의 중심이다.

1.2. 단일 GPU 최적화에서 멀티 GPU 통신으로 이동한 병목

1. 과거의 GPU 활용률 제약

   1. 메모리 접근과 단일 GPU 커널: 몇 년 전 GPU 활용률은 GPU 내부 메모리 접근이 비효율적이고 연산이 단일 GPU 커널에 갇혀 있어 제한됐다.
   2. 커널·아키텍처 개선: FlashAttention 같은 고성능 커널, DeepSeek의 메모리 효율적 아키텍처, Sparse Attention, Mamba, 더 나은 DSL(도메인 특화 언어, Domain-Specific Language)이 이 병목을 크게 줄였다.

2. 새로운 병목의 위치

   1. 멀티 GPU 통신: GPU 내부의 계산·메모리 효율이 개선되면서 병목은 GPU 사이의 데이터 이동과 집단 통신(Collective Communication)으로 이동했다.
   2. 연구의 다섯 단계: GPU 네트워킹이 지금 중요한 이유, 멀티 GPU 커널의 문제 공간, 효과적인 커널 설계의 기본 원리, 프런티어 모델의 원리 이해 여부, Parallel Kernel Bench 결과를 순서대로 연결한다.

2. GPU 하드웨어와 연결 구조의 계층

2.1. H100 내부의 계산·메모리 계층

1. 계산 프로세서와 병렬 스레드

   1. 프로세서 타일: H100 그림의 무지개색 점들은 실제 계산이 발생하는 프로세서이며, 현대 AI GPU에는 대략 100~200개가 배치된다.
   2. 스레드 스케줄링: 병렬 스레드는 더 거친 단위로 묶이고, 스레드와 블록은 프로세서에 스케줄되어 AI 계산을 수행한다.

2. 메모리 계층의 물리적 직관

   1. 레지스터(Register): 계산 유닛 바로 옆의 레지스터는 가장 빠르지만 가까운 공간이 작아 용량이 매우 제한적이다.
   2. L2 캐시와 HBM: 프로세서 사이의 사각형은 상대적으로 빠른 L2 캐시이고, 검은 상자는 대용량 가속기 메모리인 HBM(High Bandwidth Memory)이다. nvidia-smi에서 보이는 GPU의 80GB 같은 용량은 HBM을 가리킨다.
   3. 거리와 속도의 교환: 계산 유닛에 가까운 메모리는 접근이 빠르지만 양이 적고, 멀리 있는 메모리는 더 크지만 접근 시간이 길어진다.

3. H100의 구체적 수치

   1. 레지스터 대역폭: H100의 레지스터 접근 속도는 약 130TB/s로 제시되지만, 이 속도를 제공하는 메모리 공간은 매우 작다.
   2. 메모리 접근 설계: 커널은 빠른 소량의 저장공간과 느리지만 큰 저장공간 사이에서 데이터 재사용·타일링·이동량을 조정해야 한다.

2.2. GPU·노드·네트워크 사이의 연결

1. 연결 계층

   1. CPU와 GPU: CPU-GPU 통신은 PCIe를 통해 이뤄진다.
   2. 노드 간 통신: 멀티 GPU 또는 멀티 노드 통신은 InfiniBand나 TCP 같은 네트워크를 사용한다.
   3. GPU 내부 연결: NVIDIA 시스템에서 같은 스케일업(Scale-up) 도메인의 GPU 사이에는 NVLink가 지점 간 연결을 제공한다.

2. NVSwitch의 역할

   1. 비차단 패브릭: NVSwitch는 여러 NVLink 엔드포인트를 하나의 비차단(Non-blocking) 패브릭으로 연결해 GPU 간 전면 통신을 가능하게 한다.
   2. 네트워크 내부 연산: NVSwitch는 멀티캐스트(Multicast)와 리덕션(Reduction) 같은 통신 프리미티브를 네트워크 안에서 가속하는 오프 디바이스 연산도 제공한다.
   3. 대역폭: NVIDIA 세대의 NVLink는 원격 GPU 두 대 사이에 최대 약 900GB/s의 단방향(Unidirectional) 대역폭을 제공하는 것으로 제시된다.

3. 벤더별 토폴로지의 다양성

   1. AMD: AMD GPU는 스케일업 도메인에서 XGMI 인터커넥트를 사용해 GPU 사이의 지점 간 링크를 구성한다.
   2. TPU: TPU는 3D 토러스(3D Torus) 인터커넥트와 광학 랩어라운드(Optical Wraparound) 링크를 사용한다.
   3. NVIDIA: NVIDIA는 NVLink와 NVSwitch를 결합하며, NVSwitch는 연결 패브릭 자체에 리덕션 같은 계산 능력을 통합한다.

3. 왜 지금 GPU 네트워킹을 연구하는가

3.1. AI 연산 효율화가 통신 문제를 전면에 세우다

1. 소프트웨어 스택의 진화

   1. 효율적인 모델 구조: Mamba와 Sparse Attention은 필요한 계산량과 메모리를 줄인다.
   2. 하드웨어 인지 알고리즘: FlashAttention은 AI 알고리즘이 하드웨어 특성에 맞게 실행되도록 만든다.
   3. DSL과 프레임워크: TileLang, Mojo, Triton, Gluon, ThunderKittens는 AI 알고리즘을 하드웨어에 매핑하는 작업을 단순화한다.
   4. 연산 중첩: Mega-kernel은 여러 AI 연산자(Operator)의 실행을 긴밀하게 겹쳐 실행한다.
   5. 멀티 실리콘 지원: ThunderMittens는 Apple Silicon을, HipKittens는 AMD를 겨냥해 여러 벤더와 실리콘 플랫폼에서 실행할 수 있게 한다.

2. 실제 워크로드의 통신 비중

   1. 멀티 GPU 커널의 필요: 현대 AI 워크로드의 크기가 커지면서 여러 GPU에 걸쳐 실행되는 커널이 필요해졌다.
   2. 실행시간의 다수 비중: 실제 분산 학습(Distributed Training)과 추론(Inference)에서는 통신이 실행시간의 대부분을 차지하는 경우가 늘고, 규모가 커질수록 모델 FLOP 활용률(Model FLOP Utilization)이 낮아진다.

3.2. 하드웨어 발전 속도와 시스템 규모의 변화

1. 계산과 통신의 속도 차이

   1. BF16 Tensor Core: NVIDIA A100(2020)에서 B200(2024)으로 가는 동안 BF16 Tensor Core 처리 속도는 7.2배 향상됐다.
   2. 상대적으로 느린 네트워크: 같은 기간 노드 내부 통신은 3배, 노드 간 통신은 2배 향상에 그쳤다.
   3. 이질성의 핵심: Tensor Core와 메모리 계층은 AI 벤더와 실리콘이 달라도 서로 비슷해지는 반면, 네트워킹 스택은 여전히 벤더마다 크게 다르다.

2. 더 커지는 스케일업 도메인

   1. 대규모 시스템: 차세대 칩에서는 하나의 스케일업 도메인에 72개 GPU가 들어갈 수 있다.
   2. 2027년 전망: NVIDIA는 2027년에 단일 시스템에서 576개 GPU를 연결하는 구상을 제시하고 있다.
   3. 새로운 프로그래밍 기회: Tensor Memory Accelerator(TMA)는 디바이스 측에서 비동기 네트워크 전송을 수행해 이런 도메인을 더 세밀하게 제어할 수 있게 한다.

3. KV 캐시와 이질적 실행

   1. 분산된 KV 캐시: 현대 추론 시스템의 KV 캐시는 GPU, CPU, 디스크, 원격 머신 사이에 걸쳐 배치된다.
   2. 단계별 하드웨어 분리: 추론은 Speculative Decoding, Decode, Prefill을 서로 다른 하드웨어 백엔드에서 실행할 수 있다.
   3. 연구 과제: 모델은 변화하는 네트워킹 구조를 활용해야 하고, 시스템은 그 구조를 효율적으로 프로그래밍할 수 있어야 한다.

4. 멀티 GPU 커널 개발의 현재 선택지와 한계

4.1. 범용 통신 라이브러리의 장단점

1. NCCL과 RCCL의 역할

   1. 대중적인 기본 도구: NVIDIA의 NCCL과 AMD의 RCCL은 각 회사가 멀티 GPU 작업을 쉽게 만들기 위해 많은 엔지니어링을 투자한 대표적 통신 라이브러리다.
   2. 대량 전송 최적화: 두 라이브러리는 크고 연속적인 데이터 덩어리를 한 번에 옮기는 벌크 전송(Bulk Transfer)에 맞춰져 있다.

2. 세밀한 커널에서 발생하는 한계

   1. 유연성 부족: 최대 성능을 요구하는 세밀한 통신이나 여러 비자명 집단 통신을 하나로 융합하는 상황에서는 기본 설계가 충분히 유연하지 않다.
   2. 동기화 비용: Megatron-LM, FlexFlow, NanoFlow 같은 시스템은 주로 NCCL을 통해 벌크 집단 통신을 오케스트레이션하고, 전송 전후에 동기화를 요구한다.
   3. 커스텀 커널의 기회: 애플리케이션에 맞춘 통신 커널은 필요한 전송·계산·집단 통신을 직접 결합해 더 높은 성능을 낼 수 있다.

3. 기본선의 성능 여지

   1. PyTorch+NCCL 기본선: Parallel Kernel Bench에서 PyTorch와 NCCL로 작성한 단순한 기본선 다수는 통신을 고려한 루프라인 상한(Communication-aware Roofline Bound)의 50% 아래에 머문다.
   2. 해석: 기본 라이브러리가 기능적으로 편리해도 하드웨어가 제공하는 통신 한계에 가까이 가려면 전용 커널과 통신 순서 설계가 필요하다.

4.2. 컴파일러·DSL과 수작업 최적화의 한계

1. 분산 DSL의 적응 문제

   1. Triton Distributed와 TileLink: 멀티 GPU 커널을 표현하는 컴파일러와 DSL로 Triton Distributed와 TileLink가 존재한다.
   2. 빠른 네트워크 변화: 네트워킹 기술이 빠르게 바뀌기 때문에 프레임워크가 새 하드웨어 특성을 신속하게 지원하기 어렵다.
   3. 아키텍처 민감성: H800 GPU 8개를 중심으로 조정된 Triton Distributed가 H100 같은 다른 아키텍처에서는 효율적으로 적응하지 못하는 사례가 보고됐다.

2. 연산자별 수작업 튜닝

   1. 대표 사례: DPP, Comet, Ring Attention, FLUX, FlashDecoding, CUTLASS의 분산 GEMM 커널은 특정 연산자에서 최고 성능을 달성한다.
   2. 정밀도 이식 비용: 한 정밀도로 설계한 커널을 다른 정밀도로 확장하는 데 5~6개월이 걸리는 경우가 있다.
   3. 확장성 문제: 연산자와 정밀도 조합마다 사람이 세밀하게 코드를 작성하는 방식은 빠른 하드웨어·모델 변화에 확장되기 어렵다.

3. 연구 질문의 정리

   1. 소수의 지배 원리: 멀티 GPU 커널 작성 전체를 지배하는 트레이드오프와 기본 원리가 실제로 소수의 패턴으로 압축되는지 확인한다.
   2. 자동화 가능성: 그런 원리가 존재한다면 모델이 원리를 활용해 커널 개발을 단순화할 수 있는지 검증한다.

5. 멀티 GPU 커널을 결정하는 세 가지 설계 트레이드오프

5.1. 데이터 전송 메커니즘 선택

1. GPU별 Copy Engine

   1. 호스트 주도 전송: Copy Engine은 호스트 또는 CPU가 시작하는 작업이다.
   2. 큰 메시지에 적합: 전송 데이터가 큰 경우 통신 대역폭을 포화시켜 높은 처리량을 얻기 쉽다.
   3. 계산 자원 보존: 귀중한 레지스터를 빼앗지 않고, 계산을 수행하는 GPU 프로세서도 사용하지 않으므로 다른 파이프라인 단계가 그 자원을 활용할 수 있다.

2. Tensor Memory Accelerator(TMA)

   1. 디바이스 주도 전송: TMA는 GPU가 시작하는 비동기 전송을 제공하며, 장치 측에서 네트워크 데이터를 이동한다.
   2. 작은 메시지의 대역폭: 비교적 작은 메시지 크기로도 NVLink 대역폭을 포화시킬 수 있어 세밀한 통신에 적합하다.
   3. 낮은 자원 사용량: 레지스터를 매우 적게 사용하고 소수의 프로세서만으로 높은 활용률을 달성하므로 계산과 통신을 겹치기 좋다.
   4. 기능적 제약: NVSwitch가 제공하는 네트워크 내부 리덕션 같은 연산을 효과적으로 활용하지 못한다.

3. 레지스터 수준 명령어

   1. PTX 기반 제어: LDST, red, multimem 계열과 같은 레지스터 수준 PTX 명령어는 전송을 더 직접적으로 제어한다.
   2. 네트워크 내부 리덕션 활용: 이 방식은 NVSwitch의 인-네트워크(In-network) 리덕션을 활용하는 데 유리하다.
   3. 선택의 기준: 큰 벌크 전송에는 Copy Engine, 세밀한 전송과 중첩에는 TMA, NVSwitch의 네트워크 내부 계산에는 레지스터 명령어가 적합하다.

5.2. 계산·메모리·통신을 겹치는 스케줄

1. Intra-SM 스케줄

   1. 워프 전문화: 하나의 SM 안에서 일부 워프 또는 스레드는 계산을, 다른 워프는 통신을 동시에 담당한다.
   2. 데이터 정렬 조건: 계산과 통신이 같은 데이터를 입력으로 사용하고 실행 패턴이 서로 잘 맞을 때 효과적이다.
   3. 장점: 같은 SM 자원 안에서 계산과 통신을 긴밀하게 겹쳐 실행할 수 있다.

2. Inter-SM 스케줄

   1. 프로세서 전문화: 계산·통신·메모리를 서로 다른 SM 또는 GPU 프로세서에 분리 배치한다.
   2. 불일치 패턴 대응: 커널이 레지스터 파일이나 공유 메모리를 계산과 통신 단계에 서로 맞지 않는 방식으로 나눠 써야 할 때 유리하다.
   3. NVLink 활용: Intra-SM 중첩으로 NVLink 이동을 최대화하기 어려운 경우 Inter-SM 배치가 더 적합할 수 있다.

3. 대표 연산별 선택

   1. GEMM+Reduce-Scatter: GEMM과 Reduce-Scatter를 결합한 패턴에서는 Intra-SM 중첩 스케줄이 효과적이다.
   2. GEMM+All-Reduce: GEMM과 All-Reduce를 결합한 패턴에서는 NVSwitch의 네트워크 내부 리덕션을 활용하는 Inter-SM 스케줄이 효과적이다.
   3. 트레이드오프의 본질: 어느 스케줄이 항상 우월한 것이 아니라 데이터 재사용, 자원 분할, 링크 이동 패턴의 정렬 여부가 선택을 결정한다.

5.3. 버퍼링과 동기화 추상화

1. 개발자가 통제해야 할 상태

   1. 버퍼링: 데이터를 보내는 주체와 받는 주체 사이에서 언제 어느 정도의 데이터를 버퍼링할지 정할 수 있어야 한다.
   2. 동기화: 송신자와 수신자가 어떤 시점에 서로를 기다릴지 세밀하게 제어해야 한다.

2. 추상화의 조건

   1. 유연성과 단순성의 결합: 추상화는 통신 구현의 반복 작업을 감추면서도 전송 메커니즘·스케줄·버퍼·동기화 정책을 개발자가 선택하게 해야 한다.
   2. 소수 원리의 재사용: 이 세 가지 축의 조합을 작은 프리미티브와 패턴으로 표현하면 연산자별 수작업 튜닝을 줄일 수 있다.

6. ParallelKittens: 멀티 GPU 커널을 위한 작은 프로그래밍 층

6.1. 수작업으로 원리를 먼저 정립하다

1. LLM보다 먼저 만든 기반

   1. 기본 원리의 수동 검증: Together AI 팀은 곧바로 LLM에 문제를 던지기보다 멀티 GPU 커널의 트레이드오프를 직접 연구했다.
   2. ParallelKittens: 그 결과 멀티 GPU 커널을 위한 최소한의 프리미티브와 패턴 집합인 ParallelKittens를 구축했다.

2. ParallelKittens의 목적

   1. 트레이드오프 학습: 여러 통신·스케줄 조합이 성능에 미치는 영향을 이해하는 공통 기반으로 사용한다.
   2. 고성능 커널 축적: 데이터 병렬성(Data Parallelism), 시퀀스 병렬성(Sequence Parallelism), 전문가 병렬성(Expert Parallelism) 등 다양한 병렬화 방식에서 최고 수준의 커널을 모은다.
   3. LLM 평가의 컨텍스트: 먼저 사람이 패턴을 정립한 뒤 모델이 그 원리를 새 커널에 일반화할 수 있는지 검사한다.

6.2. 코드 양과 실사용 성과

1. 작은 변경량

   1. 추가 코드: 단일 GPU 커널에 멀티 GPU 프리미티브를 삽입하는 데 대략 12줄의 코드만 추가하면 된다.
   2. 추상화의 효과: 작은 코드 증가로 데이터 이동·집단 통신·버퍼링·동기화 패턴을 표현할 수 있다.

2. 실제 도입

   1. Together AI: ParallelKittens는 Together AI의 프로덕션 환경에서 사용된다.
   2. 외부 파트너: Cursor와 AI 분야의 다른 회사들도 사용한다.
   3. 성능: 데이터·시퀀스·전문가 병렬화 커널은 강력한 참고 기본선과 비교해 최첨단(State-of-the-art) 결과를 보인다.

7. Parallel Kernel Bench: 멀티 GPU 커널 생성 평가

7.1. 각 문제의 입력과 요구사항

1. 모델이 받는 정보

   1. 비최적화 기준 구현: 각 과제는 torch.distributed와 NCCL 연산으로 작성된 최적화되지 않은 PyTorch 기준 구현을 제공한다.
   2. 시스템 토폴로지: 랭크(Rank) 수와 노드 내부 하드웨어 구성을 명시한 토폴로지를 함께 제공한다.

2. 모델의 출력

   1. CUDA 커널 재작성: 모델은 기준 구현을 성능 지향 CUDA 커널로 바꿔야 한다.
   2. Unified Virtual Addressing: 생성 커널은 GPU 메모리 주소를 통합적으로 다루는 Unified Virtual Addressing(UVA)을 사용해야 한다.
   3. 측정 대상: 단순히 컴파일되는 코드가 아니라 지정된 토폴로지에서 정확하고 기준보다 빠른 커널이어야 한다.

7.2. 조합적 문제 공간과 87개 실제 과제

1. 병렬화 차원의 조합

   1. Transformer 계층의 축: 하나의 Transformer 계층은 데이터, 시퀀스, 텐서, 컨텍스트, 레이어, 파이프라인, 전문가 차원으로 병렬화할 수 있다.
   2. 통신 패턴의 변화: 병렬화 차원을 어떻게 조합하느냐에 따라 필요한 통신 패스와 데이터 이동 패턴이 달라진다.
   3. 조합적 폭발: 단일 GPU 문제보다 멀티 GPU 문제 공간이 조합적으로 훨씬 빠르게 커진다.

2. 분류 체계와 출처

   1. 대표성 확보: Parallel Kernel Bench는 멀티 GPU 문제의 대표 유형을 넓게 포괄하도록 분류 체계를 만들었다.
   2. 실제 코드 기반: GitHub 저장소, 최적화 라이브러리 구현, 멀티 GPU 커널을 작성한 DSL에서 문제를 뽑았다.
   3. 실사용 범위: 87개 과제는 추론, 강화학습(Reinforcement Learning), 포스트 트레이닝(Post-training)에서 실제로 나타나는 패턴을 반영한다.

3. 인공 문제를 피한 설계

   1. 생산성 목표: 벤치마크를 푸는 일이 인위적 장난감 커널이 아니라 새롭고 유용한 프로덕션 커널로 이어지도록 구성했다.
   2. 평가의 의미: 모델이 인터넷에 흔한 단일 GPU 패턴을 복사하는지, 하드웨어와 통신 제약을 감안해 새로운 구현을 만드는지 구분한다.

7.3. Pass@K와 Fast@K

1. 정확성 지표

   1. Pass@K: K번 생성한 시도 안에서 올바른 커널을 하나 이상 얻는지를 측정한다.
   2. 컴파일 이상의 정확성: 통신 순서와 결과 검증까지 맞아야 올바른 해로 인정된다.

2. 성능 지표

   1. Fast@K: 올바르면서 PyTorch+NCCL 기준선보다 빠른 해를 얻는지를 센다.
   2. 1배속 기준: 기준선보다 1배 이상 빠른 경우를 성능 성공으로 간주한다.
   3. 성능 지표의 엄격성: 문법적으로 맞는 커널과 실제 하드웨어에서 이득을 내는 커널 사이의 간극을 드러낸다.

8. LLM의 멀티 GPU 커널 생성 결과

8.1. 제로샷과 테스트 타임 스케일링

1. 제로샷 성능

   1. 정확성: 평가한 프런티어 모델 중 최상의 모델은 87개 문제 중 28개를 해결했다.
   2. 속도: 그 28개 중 22개가 PyTorch+NCCL 기준선보다 빨랐다.

2. 샘플 수를 늘린 결과

   1. 정확한 해 증가: 여러 샘플을 생성하는 테스트 타임 컴퓨트(Test-time Compute) 스케일링으로 정확한 해 수는 약 36개까지 늘었다.
   2. 빠른 해의 정체: 하지만 Fast@K 성능은 약 31%에서 정체됐다.
   3. 추가 샘플의 한계: 병렬 생성 수를 계속 늘려도 빠른 커널 비율이 크게 개선되지 않아 단순 샘플 확대만으로는 부족하다.

3. 모델이 얻는 성능 향상의 원천

   1. NCCL 스테이징 제거: 정답 커널에서 나타나는 속도 향상은 대개 NCCL 스테이징 오버헤드를 없애고 NVLink 직접 로드·스토어를 사용한 데서 나온다.
   2. 익숙한 패턴 편중: 성공은 집단 통신 프리미티브, 텐서 병렬 GEMM, Ulysses 스타일 컨텍스트 병렬성에 집중된다.
   3. 일반화에 대한 의문: 이런 패턴은 인터넷에 많이 등장한 유형이며, 모델이 새로운 트레이드오프를 스스로 추론했다는 증거와는 다르다.

8.2. 성능 임계값이 높아질 때의 급격한 하락

1. GPT-5.5와 DeepSeek V4 Pro 비교

   1. 최고 모델의 하락: 벤치마크에서 가장 좋은 모델로 측정된 GPT-5.5도 기준선 대비 요구 속도 향상 폭이 커지면 빠르게 통과 문제 수가 줄어든다.
   2. 차트의 의미: 가로축은 PyTorch+NCCL 기준선 대비 요구되는 속도 향상 임계값이고, 세로축은 그 임계값 이상으로 빠른 올바른 커널 수다.
   3. 비교선: GPT-5.5는 주황색 선, DeepSeek V4 Pro는 아래쪽 청록색 선으로 표시된다.

2. 문제의 성격

   1. CUDA 문법만의 문제가 아님: 오류를 보고 다시 생성하거나 다중 샘플링을 하면 모델은 커널을 컴파일 가능한 상태로 고치는 경우가 많다.
   2. 성능 추론의 실패: 컴파일 성공 이후에도 집단 통신 순서, 데이터 분할, Intra-SM 대 Inter-SM 스케줄, 전송 메커니즘 선택을 함께 추론하는 데 어려움을 겪는다.
   3. 고급 기능의 미사용: 모델은 레지스터 전송 명령어와 TMA를 커널에 넣지 않는 경우가 많다.

8.3. 코딩 에이전트와 멀티턴 실험

1. 실험 구성

   1. 에이전트 하니스: Claude Code와 비슷한 환경을 만들기 위해 mini-swe-agent 멀티턴 하니스를 사용했다.
   2. 모델과 도구: 성능이 좋았던 Gemini 3 Pro에 로컬 Bash 환경을 제공해 코드를 실행하고 오류를 확인하게 했다.

2. 에이전트의 개선 폭

   1. 정확한 문제 수: 단일 시도에 가까운 조건에서 24개를 풀던 모델이 에이전트 루프를 통해 87개 중 35개를 해결했다.
   2. 빠른 문제 수: 그중 26개가 기준선 대비 1배 초과 속도를 달성했다.
   3. 시간 확장의 한계: 에이전트에 더 많은 시간을 줘도 성능은 정체됐고, 계속 확장하려면 추가적인 탐색·검증 기법이 필요했다.

9. 새로 발견된 프로덕션 후보와 남은 과제

9.1. 벤치마크가 만든 새로운 커널

1. 실제 병목 해결 가능성

   1. 추론과 RL: Parallel Kernel Bench에서 생성된 커널은 현재 추론과 강화학습의 중요한 병목을 해결하는 새로운 프로덕션 커널로 이어질 가능성이 있다.
   2. 비인공적 문제: 실제 저장소와 라이브러리에서 문제를 골랐기 때문에 벤치마크 점수가 곧바로 실용적 구현 탐색으로 연결될 수 있다.

2. 구체적 사례

   1. Nemo vocabulary-parallel filtering: Nemo의 어휘 병렬 필터링을 위한 새로운 멀티 GPU 커널이 만들어졌다.
   2. Hyena: Hyena 아키텍처의 컨텍스트 병렬성을 위한 커널이 새로 나왔다.
   3. SAM 3: SAM 3 비디오 세그멘테이션 모델의 IoU suppression 커널이 새 후보로 등장했다.

9.2. 연구의 결론과 다음 설계 방향

1. 현재 모델의 한계

   1. 패턴의 수는 적음: 효과적인 멀티 GPU 커널 작성에 관여하는 패턴 자체는 많지 않으며 ParallelKittens의 작은 프리미티브 집합으로 캡슐화할 수 있다.
   2. 추론의 부재: 모델은 그 프리미티브와 트레이드오프를 컨텍스트에 제공해도 현재 제대로 추론하지 못한다.
   3. 벤치마크의 역할: 모델이 단일 GPU 중심의 과거 벤치마크에 과적합된 것인지, 새로운 멀티 GPU 원리를 배울 수 있는지는 Parallel Kernel Bench 같은 평가가 드러낸다.

2. 하드웨어·모델 공동 진화

   1. 네트워크 스택 적응: 네트워킹 스택이 변하는 속도와 방향에 맞춰 성장하는 모델·시스템 아키텍처가 필요하다.
   2. 스케일업으로의 이동: 더 큰 스케일업 도메인은 기존 스케일아웃(Scale-out) 중심 설계의 비중을 낮춘다.
   3. 온칩 메모리: 거대한 온칩 메모리 구조는 데이터 이동과 커널 설계의 전제를 바꾼다.

3. 실용적 시사점

   1. 개발자의 출발점: 먼저 하드웨어 토폴로지와 메모리 계층을 확인하고, 메시지 크기와 NVLink 경로에 맞춰 Copy Engine·TMA·레지스터 명령어를 선택해야 한다.
   2. 스케줄 선택: 계산·통신 데이터가 잘 정렬되면 Intra-SM 중첩을, 자원과 이동 경로가 어긋나거나 NVSwitch 리덕션이 중요하면 Inter-SM 스케줄을 검토해야 한다.
   3. LLM 활용 방식: LLM은 CUDA 문법 초안·컴파일 오류 수정·익숙한 패턴 탐색에는 유용하지만, 성능 설계의 최종 판단은 토폴로지 인지 프로파일링과 사람이 정립한 원리 검증을 거쳐야 한다.
   4. 평가 방식: 정확성만 보지 말고 기준선 대비 실제 속도, 통신 오버헤드, 다양한 병렬화 차원, 하드웨어 아키텍처 간 이식성을 함께 측정해야 한다.

주요 발언 모음

“GPU 활용률을 제한하던 병목은 멀티 GPU 통신으로 이동했다.”

“멀티 GPU 커널 작성에는 많은 패턴이 필요한 것이 아니라, 소수의 핵심 트레이드오프를 이해하는 일이 필요하다.”

“모델은 커널을 컴파일할 수 있지만, 집단 통신 순서와 데이터 분할, Intra-SM과 Inter-SM, 전송 메커니즘 사이의 트레이드오프를 아직 잘 추론하지 못한다.”

“정확한 커널을 얻은 뒤의 속도 향상은 대개 NCCL 스테이징 오버헤드를 제거하고 NVLink 직접 로드·스토어를 사용한 데서 나온다.”

“더 많은 테스트 타임 컴퓨트를 사용해도 빠른 해의 성능은 약 31%에서 정체됐다.”

핵심 데이터 & 수치

• H100 레지스터 대역폭: 약 130TB/s이며, 가장 빠른 대신 용량이 작다.
• 현대 AI GPU의 프로세서 수: 대략 100~200개다.
• NVLink 대역폭: 원격 GPU 사이 최대 약 900GB/s 단방향이다.
• 스케일업 규모: 차세대 도메인 72 GPU, NVIDIA의 2027년 단일 시스템 구상 576 GPU다.
• A100(2020)→B200(2024): BF16 Tensor Core 7.2배, 노드 내부 통신 3배, 노드 간 통신 2배 향상이다.
• 기본선 품질: 단순 PyTorch+NCCL 기본선 다수가 통신 인지 루프라인 상한의 50% 미만이다.
• ParallelKittens 코드 증가량: 단일 GPU 커널에 대략 12줄을 추가한다.
• Parallel Kernel Bench 규모: 실제 AI 워크로드에서 추출한 87개 과제다.
• 제로샷 최고 성능: 87개 중 28개 정답, 그중 22개가 PyTorch+NCCL보다 빠르다.
• 다중 샘플링: 정답 수는 약 36개까지 늘었지만 Fast@K는 약 31%에서 정체됐다.
• 코딩 에이전트: Gemini 3 Pro+mini-swe-agent가 24개에서 35개 정답으로 개선했고, 26개가 기준선보다 1배 초과 빨랐다.
• 정밀도 이식 기간: 일부 수작업 커널은 다른 정밀도로 확장하는 데 5~6개월이 걸린다.

결론 및 시사점

1. 멀티 GPU 성능의 본질: 계산량보다 통신 토폴로지, 데이터 분할, 집단 통신 순서, 자원 중첩의 선택이 성능을 좌우하는 단계에 들어섰다.
2. ParallelKittens의 의미: 소수의 프리미티브와 패턴으로 복잡한 멀티 GPU 커널 설계를 재사용 가능한 프로그래밍 층으로 정리했다.
3. LLM의 현재 위치: LLM은 문법 오류를 고치고 인터넷에 흔한 커널 패턴을 재현하지만, 새로운 하드웨어·통신 트레이드오프를 성능 있게 일반화하지 못한다.
4. 에이전트의 역할: 로컬 실행·컴파일·오류 피드백을 결합하면 정확성은 개선되지만, 실행 시간만 늘리는 전략은 성능 정체를 넘지 못한다.
5. 실행 전략: 개발자는 Parallel Kernel Bench처럼 실제 토폴로지와 기준선을 포함한 테스트로 LLM 생성물을 검증하고, TMA·레지스터 명령어·NVSwitch 리덕션을 명시적으로 탐색하게 해야 한다.
6. 미래의 방향: 더 큰 스케일업 도메인과 온칩 메모리에 적응하는 시스템·모델 아키텍처가 멀티 GPU AI 효율의 다음 단계가 된다.

핵심 요약 (20줄)

1. 멀티 GPU AI의 병목은 단일 GPU 내부 계산에서 GPU 사이 통신으로 이동했다.

2. FlashAttention, Mamba, Sparse Attention, DeepSeek식 메모리 효율화가 GPU 내부 비효율을 줄였다.

3. 현대 분산 학습과 추론은 통신에 실행시간 대부분을 쓰며 규모가 커질수록 모델 FLOP 활용률이 낮아진다.

4. A100에서 B200으로 BF16 Tensor Core 속도는 7.2배 늘었지만 노드 내부 통신은 3배만 빨라졌다.

5. NVIDIA GPU 시스템은 PCIe, NVLink, NVSwitch, InfiniBand와 TCP가 서로 다른 통신 계층을 이룬다.

6. NVSwitch는 GPU 연결을 비차단 패브릭으로 만들고 멀티캐스트와 리덕션을 네트워크 안에서 가속한다.

7. GPU별 Copy Engine은 레지스터와 계산 프로세서를 아끼면서 큰 메시지를 높은 대역폭으로 옮긴다.

8. TMA는 작은 메시지로 NVLink를 포화시키고 적은 레지스터와 프로세서로 세밀한 통신 중첩을 구현한다.

9. 레지스터 수준 PTX 명령어는 NVSwitch의 인-네트워크 리덕션을 활용하는 데 유리하다.

10. Intra-SM 스케줄은 데이터가 계산과 통신에 잘 정렬될 때, Inter-SM 스케줄은 자원과 경로가 어긋날 때 효과적이다.

11. Together AI는 이 트레이드오프를 소수의 프리미티브와 패턴으로 정리한 ParallelKittens를 만들었다.

12. ParallelKittens는 단일 GPU 커널에 약 12줄을 더해 멀티 GPU 기능을 표현하며 실제 프로덕션에서 사용된다.

13. Parallel Kernel Bench는 토폴로지와 최적화되지 않은 PyTorch+NCCL 구현을 주고 CUDA 커널 재작성을 요구한다.

14. 데이터, 시퀀스, 텐서, 컨텍스트, 레이어, 파이프라인, 전문가 병렬성의 조합은 문제 공간을 조합적으로 확장한다.

15. 실제 저장소에서 뽑은 87개 과제는 추론, 강화학습, 포스트 트레이닝의 통신 패턴을 포괄한다.

16. 최상의 제로샷 모델은 87개 중 28개를 맞혔고 22개는 PyTorch+NCCL보다 빨랐다.

17. 샘플을 늘리면 정답은 약 36개로 증가하지만 빠른 해의 비율은 약 31%에서 멈췄다.

18. 모델은 컴파일 오류를 고칠 수 있어도 집단 통신 순서와 데이터 분할의 성능 트레이드오프를 자주 놓친다.

19. Gemini 3 Pro와 로컬 Bash를 결합한 멀티턴 에이전트는 정답을 24개에서 35개로 늘렸지만 시간 확장 뒤 성능이 정체됐다.

20. LLM은 커널 초안과 오류 수정에 유용하지만 토폴로지 인지 프로파일링과 사람이 정립한 원리 검증이 여전히 필요하다.