Can LLMs Write Fast Multi-GPU Kernels? — Simran Arora, Together AI

한 줄 핵심

멀티 GPU 커널의 병목이 계산에서 네트워크로 옮겨간 지금, 핵심 전송·스케줄링 원리를 잘 정리하면 성능은 크게 올릴 수 있지만 frontier 모델은 아직 그 trade-off를 안정적으로 추론해 고성능 코드를 만들지 못한다.

시간별 요약

1. 0:00
   단일 GPU 커널에서 GPU 간 네트워크로 이제 성능 병목은 단일 GPU 내부의 메모리 접근보다 여러 GPU 사이의 통신으로 이동했다. 더 빠른 커널과 효율적인 아키텍처가 보편화되면서, 멀티 GPU 커널을 쉽게 개발하는 문제가 핵심 과제로 떠올랐다.
2. 1:29
   H100 내부 구조와 메모리 거리의 중요성 H100의 연산 유닛, L2 cache, HBM, register는 속도와 용량이 서로 다른 계층을 이룬다. 계산에 가까운 메모리일수록 빠르지만 작고, 멀수록 크지만 느리기 때문에 데이터가 어디에 있느냐가 성능을 좌우한다.
3. 4:44
   상호연결 계층: PCIe, NVLink, NVSwitch GPU 시스템에는 CPU-GPU용 PCIe, 노드 간 통신용 InfiniBand·TCP, GPU-GPU용 NVLink와 NVSwitch 같은 계층이 있다. 특히 NVSwitch는 비차단 패브릭과 multicast, reductions 같은 in-network 연산을 지원해 멀티 GPU 통신의 기반이 된다.
4. 5:46
   왜 지금 GPU 네트워킹인가 모델은 더 크고 복잡해졌지만 하드웨어 스택은 벤더마다 달라 네트워킹이 성능의 결정 변수가 되었다. 분산 학습과 추론에서는 통신이 런타임 대부분을 차지하는 경우가 많아, 네트워크 효율이 곧 AI 효율이 된다.
5. 8:27
   분산된 워크로드와 576 GPU 시스템 현대 추론 시스템은 KV cache를 GPU, CPU, disk, 원격 머신에 나눠 다루고, speculative decoding·prefill·decode를 서로 다른 하드웨어로 분리한다. 동시에 72 GPU급 scaleup domain과 576 GPU 시스템 같은 거대 구성이 등장하면서, device-side 비동기 전송과 같은 세밀한 제어가 중요해졌다.
6. 10:24
   연산은 7.2배, 네트워크는 3배 A100에서 B200으로 오면서 BF16 tensor core 성능은 7.2x 향상됐지만, intra-node 통신은 3x, inter-node 통신은 2x에 그쳤다. 연산 발전 속도에 비해 네트워크가 뒤처지면서, 멀티 GPU 커널에서 통신 최적화의 비중이 커졌다.
7. 11:14
   NCCL의 한계와 roofline 격차 NCCL 계열 라이브러리는 큰 연속 데이터 전송에는 강하지만, 세밀한 통신이나 복합 collective를 fuse해야 할 때는 유연성이 떨어진다. 단순한 PyTorch+NCCL 기준선은 대부분의 문제에서 communication-aware roofline의 절반도 못 미쳐, 커스텀 통신 커널의 여지가 크다.
8. 13:04
   컴파일러, DSL, 그리고 손으로 튜닝한 연산자 Triton distributed나 tilelang 같은 DSL은 빠르게 변하는 네트워킹 하드웨어를 따라가기 어렵다. 반대로 hand-tuned operator는 최고 성능을 내지만, precision이나 아키텍처가 바뀔 때 이식성이 낮고 재작성 비용이 크다.
9. 14:10
   연구 질문과 ParallelKittens 핵심 질문은 멀티 GPU 커널을 지배하는 작은 원리 집합이 있는지, 그리고 그것을 이용해 개발을 단순화할 수 있는지였다. 이를 위해 ParallelKittens라는 최소 프리미티브·패턴 집합을 만들고, 그 위에서 trade-off를 정리한 뒤 모델의 일반화 가능성도 시험했다.
10. 16:06
    전송 메커니즘: copy engine, TMA, register 수준 copy engine은 host 또는 CPU 주도로 큰 메시지 전송에 유리하고, TMA와 register level instruction은 GPU 주도로 더 작은 메시지에서 NVLink를 효율적으로 채운다. TMA는 register 소모가 적고 fine-grain overlap에 좋지만, NVSwitch가 제공하는 in-network computation까지는 충분히 활용하지 못한다.
11. 18:28
    SM 간 계산과 통신의 겹치기 intraSM 스케줄은 한 SM 내부에서 warp를 분리해 계산과 통신을 동시에 처리하고, interSM 스케줄은 다른 SM들에 역할을 나눠 더 큰 범위의 overlap을 만든다. 입력 데이터 정렬이 잘 맞으면 intraSM이 유리하고, NVSwitch의 in-network reductions를 쓰는 패턴에서는 interSM이 더 효과적이다.
12. 20:52
    Together AI와 Cursor의 실전 적용 ParallelKittens는 Together AI와 Cursor 같은 실제 제품 환경에서 쓰이고 있으며, 단일 GPU 커널에 약간의 코드만 더해 멀티 GPU 프리미티브를 삽입할 수 있다. data, sequence, expert parallelism에서 강한 기준선보다 좋은 결과를 보이며 생산성 있는 패턴으로 검증됐다.
13. 22:10
    ParallelKernelBench의 87개 문제 ParallelKernelBench는 PyTorch+torch distributed+NCCL 기준 구현과 시스템 토폴로지를 주고, 이를 unified virtual addressing을 사용하는 CUDA 커널로 다시 작성하게 한다. data, sequence, tensor, context, layer, pipeline, expert 등 실제 워크로드의 병렬화 축을 포괄하도록 87개 문제를 구성했다.
14. 24:32
    제로샷 결과와 스케일링 한계 frontier 모델은 제로샷에서 87개 중 28개를 풀고, 다중 샘플링으로도 정답 수는 늘지만 속도 향상은 약 31% 부근에서 정체됐다. 성공 패턴은 주로 인터넷에 많이 노출된 collective primitive나 tensor parallel gem 같은 익숙한 형태에 집중됐다.
15. 26:28
    실패 원인은 CUDA 문법이 아니다 여러 번 시도하거나 오류를 수정하게 하면 컴파일 자체는 되는 경우가 많지만, 모델은 collective ordering, data partitioning, intra/interSM 스케줄 선택, 전송 메커니즘 선택 같은 본질적 trade-off를 잘 못 다룬다. 특히 register transfer instruction이나 TMA를 적극적으로 쓰지 못하는 점이 반복적으로 드러났다.
16. 27:07
    에이전트 하네스와 아직 부족한 것 mini_sui_agent multi-turn harness와 Gemini 3 Pro, 로컬 bash 환경을 묶은 간단한 coding agent 구성은 87개 중 24개에서 35개로 성능을 끌어올렸다. 하지만 시간이 더 주어져도 성능은 곧 정체됐고, 더 나은 추론·실험·검증 체계가 추가로 필요하다는 점이 확인됐다.

실무 인사이트

• 멀티 GPU 성능 최적화는 이제 연산량보다 통신 패턴과 전송 메커니즘 선택이 더 중요하므로, 코드 생성보다 먼저 데이터 이동 경로를 설계해야 한다.
• 모델에게 커널을 맡길 때는 CUDA 문법보다 collective 순서, partitioning, overlap 전략 같은 설계 원리를 컨텍스트로 주는 편이 효과적이다.
• 실무에서는 NCCL 같은 범용 라이브러리로 시작하되, 병목 구간만 TMA나 register-level 전송으로 국소 최적화하는 방식이 가장 현실적이다.