새 GPU가 나올 때마다, 우리는 커널을 처음부터 다시 짜야 했다
ICML 2026. 하드웨어 인식 GPU 커널 최적화에 관한 발표.
이 글은 ICML 2026에서 들은 발표를 발표자 1인칭 시점으로 재구성한 것입니다.
좋은 커널을 쓰려면 하드웨어를 알아야 한다. 여기서 제가 말하는 하드웨어 인식이란 단순히 “이 GPU가 빠르다”는 걸 아는 수준이 아니라, 그 칩에만 있는 최적화 기법이 무엇인지, 그리고 그걸 실제 커널 코드로 어떻게 구현하는지를 정확히 아는 것이다. 그런데 GPU는 세대가 바뀔 때마다 이 답이 완전히 달라진다. 텐서 코어 명령어 하나만 봐도 그렇다. MMA에서 WGMMA로, 다시 TCGEN05.MMA로 — 계보가 이어질 때마다 명령어 자체가 통째로 바뀐다. 그리고 성능을 제대로 뽑으려면 이 하드웨어 전용 명령어를 반드시 써야 한다.
문제는 이 지점에서 그동안 쌓아온 추상화가 전부 무너진다는 것이다. 특히 최신 칩으로 갈수록 더 심하다. CUDA 지원 자체도 새 GPU에서는 아직 부족한데, 하드웨어 전용 커널 지원은 그보다 훨씬 더 희소하다. 그러다 보니 모델들은 Blackwell이나 Rubin 같은 최신 칩을 위한 커널을 잘 쓰지 못한다. AMD, Google, Amazon처럼 다른 벤더로 넘어가면 상황은 더 복잡해진다. 트위터를 보면 아시겠지만 요즘 새로운 칩들이 쏟아져 나오고 있고, 그 칩들 모두가 자기만의 방식으로 프로그래밍되기를 원한다.
행과 열로 하드웨어 인식을 구조화하다
그래서 이 프로젝트에서는 사람의 입력을 어떻게 구조화해야 하드웨어 인식 커널 최적화가 확장 가능해지는지를 물었다. 답을 찾다 보니 흥미로운 사실 하나를 발견했다. 원시 수준으로 내려가 보면, 사실 모든 커널이 공유하는 핵심 전략의 집합은 그렇게 많지 않다. ThunderKittens 같은 DSL을 보면, 어떤 ML 워크로드로도 확장하는 데 필요한 요소가 생각보다 적다. 밖에 존재하는 흥미로운 커널들을 뜯어보면 결국 어디서나 반복되는 공통 프리미티브들로 환원되고, 이건 유닛 테스트로 정리하고 프로파일링할 수 있다.
이걸 바탕으로 저희는 택소노미를 하나 만들었다. 세부 구조는 잠시 넘어가더라도, 핵심은 열에는 아키텍처를, 행에는 최적화 전략을 두었다는 것이다. 각 셀은 사람이 채워 넣어야 한다 — 그 최적화가 무엇인지 설명하고, 테스트할 수 있는 유닛 테스트를 주고, 조건화할 수 있는 솔루션 커널을 보여주는 식으로. 새로운 최적화 전략을 추가하고 싶으면 행을 하나 더 그리면 된다. 한 번만 비용을 치르면 그다음부터는 계속 확장된다. 새 칩으로 넓히고 싶을 때도 마찬가지다. 엔비디아가 곧 루빈을 낼 거고, 구글은 TPU를 공개용으로 풀고 있고, Etched 같은 칩도 좋은 성능을 보여주며 등장하고 있다. 걱정할 것 없다 — 열을 하나 추가하면 그만이다.
이 택소노미를 커널 최적화 태스크에 그대로 심어 넣을 수 있다는 게 핵심이었다. OpenHands 같은 에이전트에게 이 택소노미와 유닛 테스트에 대한 읽기 권한을 준다. 에이전트는 이 택소노미에 스스로를 조건화하고 필요한 걸 검증한다. 실제로 관심 있는 워크로드를 구현하려 할 때는 이 택소노미를 활용해 최적화를 시도하고, 그게 통했는지 프로파일링한다. 여기서 재밌는 지점은, 특정 최적화가 통했는지를 아주 명확하게 알려주는 프로파일링 지표가 있다는 것이다. 그리고 한번 성공한 하드웨어 인식 커널이 나오면, 그다음 최적화는 이미 잘 작동하는 커널을 조건으로 삼아 이어갈 수 있다.
토치를 이기는 것만으론 부족하다
커널이라는 태스크 자체가 워낙 넓다 보니, 저희는 포팅의 여러 축에 걸쳐 평가하고 싶었다. 포팅이란 새로운 아키텍처, 새로운 벤더, 새로운 정밀도를 대상으로 커널을 새로 쓰는 일을 말한다. ATLAS로 실험해보니 GEMM이나 어텐션 같은 표준 워크로드는 물론이고, 새로 떠오르는 어텐션 변형들에서도 흥미로운 결과가 나왔다. 아키텍처도 Ampere, Hopper 같은 예전 세대부터 Blackwell, AMD의 MI350까지 다뤘고, 정밀도 쪽에서도 FP8, MXFP4, MXFP6처럼 새 아키텍처가 들고나오는 저정밀도 포맷까지 다뤘다 — TFLOPs를 제대로 뽑으려면 이게 꼭 필요하다.
여기서 하나 짚고 싶은 게 있다. 새로 떠오르는 어텐션 변형들을 다룰 때 보니, torch는 사실 그렇게 좋은 비교 대상이 아니었다. 커널 최적화 연구 상당수가 torch와 비교해서 이겼다고 보여주지만, 이미 세상에는 torch를 10배에서 20배까지 이기는 훌륭한 Triton 커널들이 나와 있다. 저희는 그런 커널들과 비교했고, 그래도 꽤 경쟁력 있는 결과를 얻었다 — 특히 Triton 컴파일러가 아직 최적화가 덜 된 Blackwell에서 그랬다. 그리고 코딩 에이전트라는 맥락에서 저희가 정말 신경 쓰는 건 테스트 타임 컴퓨트를 어떻게 확장하느냐다. 턴을 더 많이 줬을 때 얼마나 더 좋은 성능을 뽑아낼 수 있는가. ATLAS를 표준 코딩 에이전트에 얹어보니, 베이스라인보다 훨씬 높은 지점에서 플래토가 형성됐다. 좋은 커널을 계속 써낼 수 있다는 뜻이다.
이제 조금 멀리서 보고 싶다. 저는 칩 설계를 이렇게 생각한다. 우리에겐 항상 관심 있는 워크로드가 있다 — 서브쿼드러틱 어텐션일 수도 있고, 디코드 토큰이 많은 추론일 수도 있다. 이런 워크로드는 시스템을 설계하는 방식 자체를 바꾼다. 칩을 설계하려는 사람은 먼저 아키텍처가 어떤 모습이어야 할지 이해하고, 그걸 테이프아웃하고, 스펙을 정하고, 칩을 검증한다. 이 과정은 예전엔 몇 년씩 걸렸는데 지금은 훨씬 빨라지고 있다. 그런데 그다음이 진짜 문제다. 그 새 하드웨어에 커널을 포팅해야 하고, 거기서 파생되는 모든 결정을 내려야 한다. 이 작업은 종종 몇 달에서 몇 년이 걸리고, 반드시 하드웨어와 함께 공동 설계되어야 한다. 이 전체 과정을 다 거쳐야만 비로소 워크로드를 가속할 수 있다.
그래서 저는 이걸 하나의 루프로 본다. 새로운 워크로드는 계속 나온다 — 여기 ICML에 나온 몇몇 논문도 살아남아서 그걸 위한 새 칩이 설계되길 바란다. 그런데 정작 이 루프에서 병목은 칩이 아니라 소프트웨어다. 제 지도교수님이 즐겨 보여주시는 그림이 하나 있는데, 칩을 실제로 설계하고 배포할 때 칩 자체의 비용은 그렇게 크지 않다. 진짜 비용은 소프트웨어에서 나온다. 그리고 더 많은 칩을 더 빠른 주기로 설계하게 될수록, AI 연구가 이렇게 말도 안 되는 속도로 진행될수록, 자율적인 커널 최적화에 대한 필요는 점점 커진다 — 그것도 그냥 자율적인 게 아니라, 쏟아져 나오는 이 수많은 칩들과 함께 공동 설계되는 자율성이어야 한다.
지금은 프로그래머블 AI 가속기의 르네상스다. 특정 워크로드를 위한 흥미로운 칩을 설계하는 회사들이 정말 많고, 심지어 그 회사들 상당수가 자기네 소프트웨어 스택의 많은 부분이 이미 Claude Code나 Codex로 돌아간다고 공개적으로 말하고 있다. 이건 코딩 에이전트가 이 수많은 가속기를 위한 하드웨어 인식 커널 최적화라는 영역에서 정말 독특한 위치에 서 있다는 뜻이다. 이 작업이 앞으로 나올 여러 연구 중 하나가 되어, 이 많은 가속기를 위해 하드웨어를 제대로 인식하는 커널을 어떻게 효율적으로 쓸지 깊이 고민하는 흐름에 보탬이 되었으면 한다.
질문 하나를 받았다. 이런 공동 설계 절차 때문에 다음 세대 칩의 설계 자체가 어떻게 바뀔 것 같으냐고. 제 생각은 이렇다. 칩을 설계하려면 결국 그 칩이 좋은지 나쁜지 확인하기 위해 커널을 써봐야 한다. 그런데 자율적인 커널 최적화가 가능해지면, 칩을 설계하는 과정 자체가 더 쉽고 더 좋아진다. 그러니까 이건 아주 밀접하게 맞물린 과정이다. 그래서 이 타임라인이 이렇게 빨라지고 있는 거고, 그게 정말 흥미로운 지점이다. 칩과 커널은 더 이상 순서대로 오는 두 단계가 아니다. 이제는 같이 태어난다.