GPU에서 Rust의 portable

고성능 애플리케이션 개발자가 소스 코드 수정 없이 Rust의 portable SIMD(core::simd, 여러 데이터를 한 번에 처리하는 명령어 집합)를 GPU에서 그대로 실행할 수 있게 됐다. VectorWare가 GPU 환경에서 이 기능을 성공적으로 구현하며 하드웨어 성능을 최대한 끌어올리는 기반을 마련했다. 개발자는 평소 사용하던 Rust 추상화 도구를 그대로 활용해 복잡한 고성능 프로그램을 작성할 수 있다. 이는 GPU 네이티브 소프트웨어 기업을 지향하는 VectorWare의 비전을 실현하는 중요한 단계다.

CPU와 GPU는 한 번에 처리하는 데이터 통로인 레인 수(lane count)가 달라 서로 호환되지 않는 문제가 있었다. CPU의 Simd<T, N>은 레인 수를 1에서 64까지 자유롭게 정할 수 있지만, GPU 하드웨어는 처리 폭이 고정되어 있다. NVIDIA GPU는 32개, AMD GPU는 32개 또는 64개만 지원한다.

VectorWare는 이 간극을 메우기 위해 IR(중간 표현, 하드웨어 간의 차이를 조율하는 중간 단계 언어)을 도입했다. 별도의 데이터 구조를 설계하는 대신 제네릭, const generics, trait bounds 같은 Rust 타입 시스템에 이 설정을 인코딩해 관리한다. 하드웨어마다 다른 레인 수 제약을 Rust의 타입 체계 안에서 직접 해결한 결과다.

portable SIMD는 아키텍처별 벤더 고유

`_mm256_add_ps` 같은 x86-64 전용 명령어와 `vaddq_f32` 같은 Arm 전용 명령어를 각각 따로 짜야 했던 제약이 사라졌다. 기존에는 `core::arch`(하드웨어 제조사가 제공하는 특수 명령어 집합)를 통해 특정 아키텍처에 종속된 코드를 직접 구현해야 했다. 이제는 `Simd<T, N>`이라는 단일 제네릭 타입을 통해 T 타입의 요소 N개를 가진 벡터를 정의하는 추상화 계층을 사용한다. 이렇게 작성한 코드는 컴파일러가 대상 CPU의 벡터 명령어로 자동 변환하므로 하드웨어마다 다른 구현체를 만들 필요가 없다.

NVIDIA GPU의 SIMT(하나의 명령어로 여러 스레드를 동시에 처리하는 방식) 모델은 이 벡터 구조를 그대로 받아들인다. GPU의 워프(warp)는 하나의 명령어를 발행하고 32개의 레인이 각자의 데이터를 처리하는 단위다. `Simd<i16, 32>`와 같은 타입은 워프의 32개 레인에 요소를 하나씩 할당해 단일 워프 명령어로 컴파일된다. CPU를 위해 설계된 벡터 유닛이 GPU의 넓은 벡터 유닛에 직접 매핑되어 작동하는 구조다.

데이터의 순서를 바꾸거나 합치는 복잡한 연산도 GPU 하드웨어 명령어와 직접 연결했다. `reduce_sum`이나 `reduce_max` 같은 리덕션(여러 값을 하나로 합치는 연산)은 GPU의 워프 셔플(warp shuffle) 명령어를 사용한다. `simd_swizzle!` 같은 셔플이나 회전 연산 역시 동일한 워프 셔플 프리미티브(기본 연산 단위)를 활용해 구현했다. `Mask<T, N>`에서 제공하는 `any`나 `all` 같은 수평 쿼리는 GPU의 보트(vote) 및 발롯(ballot) 명령어로 변환되어 처리된다.

Rust Nightly 버전에서 `#![feature(portable_simd)]` 플래그를 사용하면 기존에 portable SIMD로 작성한 CPU 라이브러리를 별도 리라이팅 없이 NVIDIA GPU에서 가속할 수 있다. 다만 컴파일러 수정 사항에 대한 모든 케이스의 안정성은 아직 완전히 검증되지 않은 상태다.