이제 GPU 코드에서 Rust의 스레드를 사용할 수 있습니다. 구현 방식과 이것이 GPU 프로그래밍에서 무엇을 가능하게 하는지 공유합니다.
2026년 3월 24일 15분 읽기
Pedantic mode:끄기
이제 GPU 코드에서 Rust의 스레드를 사용할 수 있습니다. 구현 방식과 이것이 GPU 프로그래밍에서 무엇을 가능하게 하는지 공유합니다.
VectorWare에서는 최초의 GPU-native software company를 만들고 있습니다. 오늘 우리는 GPU에서 Rust의 std::thread를 성공적으로 사용할 수 있게 되었다는 소식을 매우 기쁘게 전합니다. 이 이정표는 개발자가 익숙한 Rust 추상화를 사용해 GPU 하드웨어의 전체 성능을 활용하는 복잡한 고성능 애플리케이션을 작성할 수 있게 하려는 우리의 비전에 있어 중요한 진전입니다.
CPU와 GPU는 근본적으로 다른 방식으로 프로그램을 실행합니다. CPU 프로그램은 하나의 스레드에서 시작해 필요에 따라 추가 스레드를 생성합니다. 각 스레드는 독립적으로 실행되며, 언제 어떤 방식으로 동시성을 도입할지는 프로그래머가 제어합니다.
GPU 프로그램은 다르게 동작합니다. GPU 프로그램은 하나 이상의 커널로 구성됩니다. 각 커널은 병렬로 실행되는 매우 많은 인스턴스로 실행됩니다. 동시성은 프로그래머가 명시적으로 도입하는 것이 아닙니다. 그것은 하드웨어가 GPU 프로그램을 실행하는 방식에 본질적으로 내재되어 있습니다.
이 모델은 행렬 곱셈, 이미지 처리, 그래픽 렌더링처럼 모든 warp가 서로 다른 데이터에 대해 같은 작업을 수행하는 균일한 워크로드에 잘 맞습니다.
GPU 프로그램이 더 정교해질수록 개발자들은 warp specialization을 사용해 같은 프로그램의 서로 다른 부분을 서로 다른 warp에서 동시에 활성화합니다.
대부분의 CPU 프로그래밍 모델은 프로그램의 진입점으로 main 함수를 사용합니다. 실행은 정확히 하나의 스레드에서 시작하므로, 프로그램을 함수로 표현하는 것은 자연스럽습니다. 함수 본문은 그 하나의 스레드가 수행하는 작업을 설명합니다.
VectorWarefn main() {
// Single threaded CPU code
}
놀랍게도 대부분의 GPU 프로그래밍 모델도 진입점으로 함수를 사용합니다. 프로그래머는 함수가 한 번 실행되는 것처럼 작성하지만, 하드웨어는 그 함수를 수천 번 병렬로 실행합니다. GPU 커널은 일반적인 CPU 함수처럼 보이는 함수이지만 동작은 매우 다릅니다.
// CUDA C kernel
__global__ void scale(float* data) {
// This code is executed thousands of times in parallel,
// each with the same function parameters, but different values for the global
// indexes such as `blockIdx` and `threadIdx`.
int i = blockIdx.x * blockDim.x + threadIdx.x;
data[i] *= 2.0f;
}
프로그래밍 모델과 실행 모델 사이의 이러한 불일치는 GPU 프로그래밍이 어려운 이유 중 하나입니다. 한 번 실행되는 함수와 수천 번 병렬로 실행되는 함수는 의미론이 매우 다르지만, 컴파일러와 프로그래머 모두 코드만 보고는 이를 쉽게 추론할 수 없습니다. 실제로는 공유 데이터에 대한 올바른 인덱싱 유지와 경쟁 상태 회피 같은 불변 조건을 수동으로 지키는 책임이 프로그래머에게 있습니다.
Rust로 작성된 GPU 프로그램도 같은 패턴을 따르며 함수로 모델링됩니다. 같은 GPU 커널을 Rust로 작성한 예를 보겠습니다.
use core::arch::nvptx::*;
pub unsafe extern "ptx-kernel" fn scale(data: *mut f32) {
// This code is executed thousands of times in parallel,
// each with the same function parameters, but different values
// for the global indexes such as `_block_idx_x` and `_thread_idx_x`.
let i = (_block_idx_x() * _block_dim_x() + _thread_idx_x()) as usize;
*data.add(i) *= 2.0;
}
이 커널은 unsafe가 필요하며 reference 대신 *mut f32 raw pointer를 받습니다. GPU는 이 함수를 수천 개 인스턴스로 동시에 실행하고 각 인스턴스는 같은 포인터를 받기 때문에, Rust의 소유권 모델을 사용하는 함수로 이를 안전하게 표현할 방법이 없습니다. Rust는 fn main이 단일 스레드에서 실행되는 CPU의 실행 모델을 중심으로 설계되었고, 언어는 그 위에서 안전성을 강제할 수 있습니다. GPU의 실행 모델은 이 언어에 낯선 것이며 커널 경계는 FFI 경계처럼 취급됩니다. 즉 raw pointer, unsafe, 그리고 컴파일러 보장이 없습니다. 이것도 작동은 하지만, 이상적으로는 Rust의 안전성 보장이 GPU까지 확장되어야 합니다.
GPU 특유의 의미론을 포착하기 위해 새로운 타입과 어노테이션을 도입할 수도 있습니다. 하지만 그렇게 하면 일반적인 Rust와 분리된 새로운 프로그래밍 모델이 생기게 됩니다. 프로그래머는 새로운 추상화를 배워야 하고 GPU 전용 코드를 작성해야 합니다. 우리는 GPU 코드가 Rust 생태계와 자연스럽게 통합되는 평범한 Rust 코드처럼 보이기를 원합니다.
우리는 Rust가 GPU 실행 모델을 이해하도록 만들고, 컴파일러가 GPU 고유의 불변 조건을 추론하도록 확장할 수도 있습니다. 그것은 세심한 설계 작업이 많이 필요한 장기 프로젝트입니다. Rust compiler team의 구성원으로서 우리는 이 노력에 기여하고 싶지만, 오늘 당장 Rust로 안전한 GPU 코드를 작성하고 싶습니다.
함수가 GPU에서 실행되는 방식과 유사하게 이를 호출하는 CPU 하네스를 만들면 일부 안전성을 되찾을 수 있습니다. 각 GPU 인스턴스는 CPU 스레드로 모델링할 수 있고, 이를 통해 miri 같은 도구가 가능한 인터리빙을 탐색하고 Rust 메모리 모델 아래에서 정의되지 않은 동작을 검사할 수 있습니다. VectorWare에서는 이런 하네스를 사용합니다.
CPU 기반 하네스는 GPU 실행의 일부 측면을 포착할 수 있을지 모르지만, 언제나 새는 추상화로 남습니다. 이상적인 최종 상태는 GPU 프로그램이 CPU 프로그램처럼 동작하여 Rust 컴파일러가 두 환경에서 동일한 불변 조건을 추론할 수 있게 되는 것입니다. 단일 함수 커널 진입점은 동시성이 암묵적이기 때문에 이를 어렵게 만듭니다. 동시성이 명시적인 모델은 프로그래머와 컴파일러 모두에게 추론하기 더 쉽습니다. 스레드는 그런 모델 중 하나입니다.
std::thread를 지원하나요?Rust 프로그램은 동시성을 위해 주로 futures와 threads라는 두 가지 모델을 사용합니다.
이전 글에서 우리는 GPU에서 최초로 실행되는 futures와 async/await를 시연했습니다. 그러나 Rust의 std를 GPU로 가져왔을 때는 스레드를 구현하지 않았습니다. 어떻게 해야 할지가 분명하지 않았고, GPU-native 앱을 작성하기 위한 인체공학적인 동시성 수단으로 이미 async/await가 있었기 때문입니다.
그럼에도 Rust 생태계의 상당 부분은 futures보다 스레드를 중심으로 구축되어 있습니다. rayon 같은 널리 사용되는 스레드 풀, tokio 같은 async 런타임, 그리고 병렬성을 위한 많은 라이브러리가 모두 std::thread에 의존합니다. 스레드를 지원하면 기존 생태계의 큰 부분이 열립니다.
std::thread를 GPU threads에 매핑하지 않나요?warp 내부에서 GPU는 많은 threads를 가집니다. 이를 lanes라고도 부릅니다. 가장 분명한 접근은 각 std::thread를 그중 하나에 매핑하는 것입니다.
하지만 GPU의 "thread"는 CPU 프로그래머가 말하는 "thread"와 다릅니다. GPU thread는 warp 내부의 단일 lane이며, 독립적인 실행 컨텍스트라기보다 CPU의 SIMD lane에 더 가깝습니다.
CPU thread는 자기만의 스택, 자기만의 프로그램 카운터를 가지며 독립적으로 스케줄될 수 있습니다. GPU lane은 이런 방식으로 동작하지 않습니다. warp 내부의 lane들은 lockstep으로 함께 전진합니다. std::thread를 GPU lane에 매핑하면 Rust가 기대하는 의미론을 깨뜨리게 됩니다.
그리고 그것은 느릴 것입니다. warp 내부의 lane들이 서로 다른 분기를 타면, GPU 하드웨어는 비활성 lane을 마스킹하고 각 경로를 순차적으로 실행할 수 있습니다. 이것을 divergence라고 합니다. thread::spawn()이 lane에 매핑된다면, 생성된 lane과 호출한 lane은 서로 다른 코드를 실행하면서 같은 warp에 있게 됩니다. 하드웨어가 이들을 직렬화할 수 있으므로 동시성의 이점이 사라집니다.
std::threadGPU에서 스레드를 실행하는 것은 코드를 시각적으로 보여 주기 어렵습니다. 코드가 평범한 Rust처럼 보이고 그렇게 실행되기 때문입니다. 의도적으로 CPU에서 쓰는 것과 같은 문법이 수정 없이 GPU에서 실행됩니다. 아래는 두 개의 스레드를 생성하고, 각 스레드에서 작업을 수행한 뒤, 이를 join하는 Rust 프로그램입니다.
use std::thread;
fn main() {
let a = thread::spawn(|| {
let mut sum = 0u64;
for i in 0..1000 {
sum += i;
}
sum
});
let b = thread::spawn(|| {
let mut product = 1u64;
for i in 1..20 {
product *= i;
}
product
});
let sum = a.join().unwrap();
let factorial = b.join().unwrap();
println!("sum: {sum}, factorial: {factorial}");
}
아래는 이 프로그램이 GPU에서 실행되는 녹화입니다. 코드는 GPU 커널로 컴파일되어 디바이스에서 실행됩니다. 두 스레드는 각각 별도의 warp에서 실행되고, 결과는 우리의 std support를 통해 GPU에서 출력됩니다.
GPU에서 std::thread를 지원할 수 있게 해 주는 핵심 관찰은 세 가지입니다.
std::thread는 이와 같은 개념을 언어가 제공하는 API 뒤에 두고, 소유권, 타입 검사, 수명 검사를 내장한 형태라고 볼 수 있습니다.동작 방식은 이렇습니다. 우리는 각 std::thread를 하나의 GPU warp에 매핑합니다. 커널이 시작되면 Warp 0만 활성화됩니다. Warp 0은 CPU의 메인 스레드처럼 main을 실행합니다. 다른 모든 warp는 잠듭니다. thread::spawn()을 호출하면 잠자고 있던 warp 하나를 깨워 생성된 클로저를 실행시킵니다. thread::join()을 호출하면 부모 warp는 자식 warp가 끝날 때까지 블록됩니다.
아래는 코드의 각 부분을 어떤 warp가 실행하는지 보여 주도록 주석을 단 같은 프로그램입니다.
warp-as-thread 모델이 마련되면 std::thread의 나머지 부분도 자연스럽게 따라옵니다. thread::current()는 현재 warp의 식별자를 반환합니다. thread::sleep()는 nanosleep을 사용해 warp를 잠재웁니다. thread::yield_now()는 warp를 스케줄러에 다시 양보합니다. thread 이름, thread IDs, 그리고 builder patterns도 모두 예상대로 동작합니다.
우리의 구현은 NVIDIA GPU를 대상으로 하지만, 이 접근 자체에 CUDA 고유인 부분은 없습니다. Vulkan에는 subgroups가 있고 HIP/ROCm에는 wavefronts가 있으며, 둘 다 같은 warp-as-thread 매핑을 구현하는 데 사용할 수 있습니다.
이 접근의 큰 장점 중 하나는 프로그래머가 warp, block, grid를 생각할 필요가 없다는 점입니다. CPU에서와 마찬가지로 생성된 스레드는 자기만의 실행 컨텍스트를 가지며, 부모와 다른 코드를 실행하고, join할 수 있습니다. 이것이 어떻게 일어나는지에 대한 세부 메커니즘은 추상화 뒤로 숨겨집니다.
또 다른 장점은 이 접근이 구조적으로 divergence를 방지한다는 점입니다. divergence는 warp 내부의 lane들이 서로 다른 분기를 탈 때 발생합니다. thread::spawn()이 하나의 클로저를 하나의 warp에 매핑하기 때문에, 그 warp의 모든 lane은 같은 코드를 실행합니다. 단일 std::thread 내부에서 divergent branching을 표현할 방법이 없으므로 divergence는 발생할 수 없습니다. 최악의 경우 워크로드가 warp당 한 개 lane만 사용하고 나머지 lane은 유휴 상태로 남을 수 있습니다. 하지만 유휴 lane은 divergent lane보다 명백히 낫습니다. 유휴 lane은 용량을 낭비하지만 divergent lane은 실행을 직렬화합니다. lane 전체에 걸쳐 넓게 실행해야 하는 코드는 여전히 스레드의 클로저 내부에서 warp_shuffle_idx 같은 warp-level intrinsics를 사용해 명시적으로 그렇게 할 수 있습니다.
가장 중요한 점은, 이 접근을 사용하면 Rust의 borrow checker와 lifetime이 그대로 동작한다는 것입니다. 이것은 Rust 프로그래머가 이미 알고 있는 의미론이며, 기존 코드는 바로 그런 의미론을 전제로 작성되어 있습니다. 우리는 Rust에 새로운 GPU 프로그래밍 모델을 도입하는 것이 아닙니다. 우리는 Rust의 프로그래밍 모델을 GPU 위에 매핑하고 있습니다. VectorWare에서는 GPU가 평범한 Rust 플랫폼처럼 동작하도록 만들고 있습니다.
이제 GPU에서 std::thread와 async/await가 모두 동작하므로, Rust 생태계의 큰 부분이 GPU 하드웨어에서 실행 가능해집니다. 병렬성을 위해 스레드를 사용하거나, I/O를 위해 async를 사용하거나, 혹은 둘을 조합하는 라이브러리들이 이제 거의 수정 없이 또는 전혀 수정 없이 GPU를 대상으로 할 수 있습니다. 앞으로의 블로그 글에서는 이것이 가능하게 하는 흥미로운 응용 사례를 소개할 예정입니다.
그렇다고 해서 우리의 최종 목표가 단지 기존의 CPU 지향 소프트웨어를 GPU에서 실행하는 것이라고 생각하지는 않습니다. 우리는 CPU-native 소프트웨어로는 불가능한 방식으로 하드웨어를 활용하는 새로운 GPU-native 애플리케이션을 작성하는 데 더 큰 기대를 걸고 있습니다.
warp는 유한한 자원입니다. 너무 많은 스레드를 생성하면 사용 가능한 warp가 고갈되어 큐잉이 필요하거나 실패하게 됩니다. 실제로는 이것이 생각보다 큰 문제는 아닙니다. 대부분의 Rust 코드는 이미 하드웨어에 적응하기 때문입니다. 표준 라이브러리는 하드웨어가 동시에 실행할 수 있는 스레드 수를 질의하는 std::thread::available_parallelism()을 제공합니다. GPU에서는 이 함수가 사용 가능한 warp 수를 반환하도록 했습니다.
GPU에서의 스레드 동기화는 CPU보다 더 비쌉니다. mutex, condition variable, 그리고 다른 블로킹 프리미티브는 점유율과 처리량을 낮출 수 있는 warp 수준의 스케줄링 결정을 필요로 합니다.
warp-as-thread 모델은 생성된 각 스레드가 warp 전체를 소비한다는 뜻입니다. 생성된 코드가 모든 lane을 사용하지 않으면 GPU 하드웨어 활용도가 낮아집니다. 이는 복잡하고 divergent한 워크로드에는 허용 가능하지만, 데이터 병렬성이 더 적합한 단순한 병렬 작업에는 낭비입니다.
아직 스레드가 할당되지 않은 warp는 유휴 상태로 남아 유용한 작업 없이 하드웨어 자원을 소모합니다. 우리는 이에 대한 완화책을 가지고 있으며, 향후 글에서 우리의 접근을 설명할 예정입니다.
마지막으로 스레드용 스택 메모리는 GPU 메모리에서 할당되어야 하는데, 이는 CPU 메모리보다 더 제약이 큽니다. 깊은 호출 스택이나 많은 동시 스레드는 사용 가능한 메모리를 고갈시킬 수 있습니다. 기본 CUDA 스택 크기는 대부분의 복잡한 애플리케이션에는 너무 작으며 cudaLimitStackSize를 통해 변경해야 합니다.
우리는 이 작업을 몇 달 전에 완료했습니다. GPU에서 이처럼 빠르게 진전을 이룰 수 있었던 것은 Rust의 추상화와 생태계가 가진 힘을 잘 보여 줍니다.
회사로서 우리는 모든 사람이 Rust를 사용하는 것은 아니라는 점을 이해합니다. 앞으로의 제품은 여러 프로그래밍 언어와 런타임을 지원할 것입니다. 하지만 우리는 Rust가 고성능이면서 신뢰할 수 있는 GPU-native 애플리케이션을 구축하는 데 특히 잘 맞는 언어라고 믿으며, 바로 그것에 가장 큰 기대를 걸고 있습니다.
우리의 진행 상황을 계속 받아 보려면 X, Bluesky, LinkedIn, 또는 우리의 blog를 구독해 주세요. 앞으로 몇 달 동안 우리의 작업에 대해 더 많은 내용을 공유할 예정입니다. 또한 hello@vectorware.com으로도 연락하실 수 있습니다.