GPU 프로그래밍은 오랫동안 CUDA C++나 벤더 고유 도구에 묶여 있었다. 러스트로 GPU를 쓰려는 시도가 없지는 않았지만, 대개는 GPU 전용 언어를 흉내 낸 별도 하위 집합이거나 특수한 어노테이션을 요구했다. 이런 배경에서 GPU-네이티브 소프트웨어 회사를 표방하는 VectorWare가 러스트의 포터블 SIMD(core::simd)를 GPU 위에서 그대로 실행하는 데 성공했다고 발표한 것은 눈여겨볼 만하다. 핵심은 새로운 벡터 타입이나 문법을 만든 것이 아니라, 이미 존재하는 러스트 표준 추상화를 GPU의 실행 모델에 대응시켰다는 점이다.
이 회사는 앞서 std::thread를 GPU의 워프(warp)에 매핑해 러스트 스레드를 GPU에서 돌리는 작업을 해왔다. 다만 그때는 워프 단위의 동시성만 활용했을 뿐, 하나의 워프 안에 존재하는 병렬 레인은 놀리고 있었다. CPU에서 스레드 내부 병렬성의 추상화가 SIMD인 것처럼, GPU에서도 그 빈자리를 채울 도구가 필요했고 그것이 이번 포터블 SIMD 매핑이다.
워프는 곧 벡터 유닛이다
기존에 러스트로 SIMD를 쓰려면 x86-64의 _mm256_add_ps나 Arm의 vaddq_f32 같은 아키텍처별 인트린식을 직접 불러야 했고, 여러 아키텍처를 지원하려면 각각 따로 구현해야 했다. 포터블 SIMD는 그 위에 Simd이라는 제네릭 타입을 두어, 산술·비교·리덕션·레인 셔플을 한 번만 작성하면 컴파일러가 대상 CPU의 벡터 명령으로 낮춰주는 구조다. 게다가 이 타입은 std가 아니라 core에 있어 GPU에서 std 지원조차 필요하지 않다.
VectorWare의 통찰은 GPU를 포터블 SIMD가 겨냥할 또 하나의 벡터 하드웨어로 본 데 있다. 엔비디아가 SIMT(Single Instruction, Multiple Thread)라 부르는 GPU의 실행 모델에서 워프는 하나의 명령을 발행하고 32개 레인이 각자의 데이터에 그 명령을 수행한다. 이는 곧 하나의 명령이 여러 데이터를 처리한다는 SIMD의 정의 그 자체이며, 레인별 주소 지정이 추가될 뿐 본질은 다르지 않다. 예컨대 Simd는 32개 레인 각각에 i16 원소 하나씩을 배정하고, 두 벡터를 더하면 모든 레인이 동시에 덧셈을 수행하는 단일 워프 명령으로 컴파일된다.
연산별로 보면 대응은 자연스럽다. 덧셈·곱셈·비교 같은 원소별 연산은 Simd에 구현된 러스트 트레이트에서 나와 GPU가 네이티브로 실행한다. reduce_sum이나 reduce_max 같은 리덕션은 GPU의 워프 셔플 명령으로 레인 간 값을 교환·결합한다. simd_swizzle!이나 회전 같은 크로스 레인 셔플 역시 같은 워프 셔플 프리미티브에 매핑되고, Mask의 any·all 같은 질의는 GPU의 vote/ballot 명령을 쓴다. 루프 카운터나 상수처럼 모든 레인이 동일하게 계산하는 스칼라 값은 CUDA에서처럼 워프 전역에 복제되는데, 이는 ISPC 같은 언어가 명시하는 uniform 대 varying 구분이 러스트 타입에서 저절로 드러나는 셈이다.
어긋나는 지점과 타입 기반 IR
추상화와 하드웨어가 맞지 않는 유일한 지점은 레인 수다. CPU에서 Simd은 N이 1부터 64까지 자유롭지만 GPU는 폭이 고정돼 엔비디아는 32, AMD는 32 또는 64다. N이 이 폭과 정확히 일치할 때만 일대일 매핑이 성립하며, 더 작으면 일부 레인이 놀고 더 크면 레인마다 여러 원소를 처리해 명령이 늘어난다. 워프를 '레인 간 데이터 이동·결합 프리미티브와 활성 레인 규칙을 갖춘 작은 기계'로 보고, 그 위에 작업을 배치하는 것이 프로그래밍이 된다. VectorWare는 이 기계에 별도 자료구조가 아니라 러스트의 타입·제네릭·const 제네릭·트레이트 바운드로 인코딩한 IR을 부여했다. 연산이 형태를 타입에 담고 있어 잘못된 프로그램 상당수는 아예 구성조차 되지 않으며, IR은 GPU에서 인터프리터 없이 곧바로 명령으로 낮춰져 손으로 쓴 PTX 대비 추가 비용이 없다. 동시에 같은 타입으로 CPU에서 결정론적으로 실행하는 참조 인터프리터, 이른바 워프 레인 프로그래밍용 Miri를 만들어 차등 테스트에 활용한다.
실무자가 짚어야 할 한계
실무 관점에서 매력은 분명하다. 동일한 소스가 CPU와 GPU에서 돌아가므로, 이미 포터블 SIMD를 쓰는 코드와 라이브러리는 재작성 없이 GPU 실행 후보가 된다. Simd은 평범한 소유 값이라 빌림 검사기·라이프타임·타입 검사가 CPU와 똑같이 적용되고, GPU 전용 어노테이션도 없다. 다만 전제 조건이 만만치 않다. 포터블 SIMD 자체가 아직 러스트에서 불안정 기능이라 nightly의 #![feature(portable_simd)]가 필요하고 API가 바뀔 수 있다. 제로 코스트는 벡터 폭이 워프 레인 수와 일치할 때만 성립하며, 임의의 순열 셔플은 여러 명령이나 공유 메모리 경유를 요구할 수 있다. 리덕션이나 all/any 같은 수평 연산은 워프 내 동기화 지점으로 작동해 스케줄러의 작업 중첩을 제약한다. VectorWare 스스로도 다른 러스트 기능과의 상호작용을 위해 컴파일러를 손봐야 했고 모든 경우를 다뤘다고 확신하지 못한다고 밝힌다.
지금은 엔비디아를 겨냥하지만 IR 자체는 아키텍처 중립적이어서 AMD 웨이브프론트나 벌칸 서브그룹으로 확장할 여지가 있다고 한다. 이 회사는 스레드·SIMD·async를 GPU에 각각 매핑한 데 이어 세 축을 조합하는 방향, 그리고 행렬형 SIMD를 텐서 코어로 낮추거나 평범한 스칼라 루프를 자동 벡터화하는 과제를 다음 목표로 제시한다. 러스트를 쓰지 않는 개발자를 위해 향후 제품은 여러 언어와 런타임을 지원하겠다는 단서도 달았다. 결국 이 발표의 의미는 성능 수치가 아니라, GPU를 '특수한 이기종 장치'가 아니라 러스트가 겨냥하는 평범한 플랫폼의 하나로 취급할 수 있음을 보여준 데 있다. 그 방향이 실제 프로덕션에 닿으려면 nightly 의존과 컴파일러 개편이라는 불확실성을 먼저 통과해야 한다.