Codex 자동 연구로 기준보다 232배 빠른 GPU 커널 만들기

4 hours ago 3

GPU Mode의 qr_v2 대회에서 Codex 기반 반복 최적화로 torch.geqrf 기준 약 419,000µs를 1,805µs까지 줄여 183명 중 12위를 기록함 순차 의존성이 강한 Householder QR을 블록 Householder·WY 표현으로 재구성해 좁은 패널만 직렬 처리하고, 후행 행렬 갱신은 텐서 코어에 적합한 GEMM으로 전환함 14일 동안 1,500회 넘게 제출하며 /goal, 프로파일링, 로그, 서브 에이전트를 활용했고, 커널 융합·고정 크기 특수화·CUDA 그래프로 실행 및 패널 오버헤드를 줄임 3,000µs 아래에서는 단일 후보 조정이 한계에 부딪혀 3~5개 후보 빔, 고위험 구조 변경, 조언자 모델, NCU·Modal 교차 프로파일링으로 아이디어의 다양성을 확보함 도메인 지식은 더 나은 질문과 실험 설계를 가능하게 했으며, 입력 분포별 특화·라이브러리 함수 제거·FP16 상주·tcgen05 활용은 남은 개선 기회였음 대회 과제와 결과 GPU Mode와 Core Automation이 개최한 자동 연구 대회의 과제는 배치 정사각 행렬에 대한 compact-Householder QR 분해 구현이었음 입력은 batch x n x n 형태의 FP32 CUDA 행렬 A이며, 출력은 torch.geqrf(A)와 동일한 (H, tau) 형식이어야 했음 H의 상삼각 부분에는 R, 하삼각 부분에는 Householder 벡터가 저장됨 tau에는 각 반사 변환의 계수가 저장됨 검사기는 torch.linalg.householder_product(H, tau)로 Q를 복원하고 R = triu(H)를 취해 A ≈ QR과 직교성을 확인함 내부 계산에는 FP16, FP8, NVFP4를 사용할 수 있었지만, 반환 결과는 FP32 수준의 QR 검사를 통과해야 했음 주요 크기는 512 x 512였고 1024, 2048, 4096 사례도 포함됐으며, 올바른 제출은 크기와 조건별 실행 시간의 기하평균으로 순위가 결정됨 최종 결과는 183명 중 12위, 기준 약 419,000µs 대비 1,805µs로 약 232배 가속이었음 추적 가능한 제출 이력은 108,803µs에서 시작해 1,805µs까지 98.34% 감소함 14일 동안 1,500회 넘게 제출함 전체 문제와 순위는 Problem Link and Leaderboard에서 확인할 수 있으며, 대회는 Linear Algebra Kernels in the Age of Re...

Read Entire Article