최신 GPU는 최대 성능에 도달하기 위해 SIMT 스타일의 병렬성뿐만 아니라 컴퓨팅과 데이터 이동 간 소프트웨어 관리 컨커런시(concurrency)도 요구합니다. 성능 엔지니어들은 작업 분할을 계산 리소스 계층 구조(스레드, 워프, 워프그룹, 블록, 클러스터)로 추론해야 하며, 많은 경우 메모리 계층 구조(레지스터, 텐서 코어 누산기, 공유 메모리, 전역 메모리)의 서로 다른 레벨에서 비동기 텐서 코어 및 memcpy 명령어를 사용해야 합니다. 아웃-오브-오더 실행이 하드웨어에 의해 관리되고 프로그래머로부터 숨겨지는 CPU와 달리, GPU는 이러한 비동기 명령어들을 통해 명시적인 명령어 재배열(instruction reordering)을 소프트웨어에 노출합니다. 잘 알려진 GPU 프로그래밍 언어들은 일반적으로 안전성 보장 없이 직접적인 저수준 제어를 제공하거나(예: CUDA C++ 인라인 어셈블리 또는 인트린식), 비동기 명령어를 컴파일러 백엔드에서 숨겨서 분석하기 더 쉬운 고수준 추상화(예: Triton의 타일 기반 모델)를 제공합니다. 이 방식은 성능 엔지니어가 중요한 세부 사항을 튜닝하여 성능을 극대화하는 것을 방해할 수 있습니다. 우리는 CUDA에 대한 최소한의 추상화를 생성하는 명령형(imperative), 저수준 언어인 Exo-GPU를 제안합니다. 우리의 핵심 아이디어는 병렬성과 동기화를 근본적인 제어 흐름 원시 요소(control flow primitives)가 아니라 순차 코드에 대한 단순한 주석(annotation)으로 취급하여, 이러한 구조들이 프로그램 의미론(semantics)을 변경하지 않음을 검증할 수 있게 하는 것입니다. 그 이점은 두 가지입니다: 프로그래머는 숨겨진 제어 흐름이나 변형 없이 코드를 추론할 수 있으며, 동시에 Exo-GPU 컴파일러가 순차-병렬 동등성(sequential-parallel equivalence)을 검증하도록 허용합니다. 이는 병렬 실행이 그 순차적 해석과 기능적으로 동등함을 보장합니다. 우리는 Exo-GPU를 사용하여 H100 GPU의 GEMM 커널을 wgmma, TMA, 그리고 split-k를 이용해 작성했습니다. 우리의 커널은 큰 문제 크기에서 이론적 최대 성능의 80% 이상을 달성했으며, 일부 사례에서는 공급업체 제공 CUBLAS 라이브러리보다 우수한 성능을 보였습니다.