VectorWare đưa Rust Portable SIMD lên GPU: Viết một lần, chạy mọi nơi

Bản tin

10 tháng 8, 2026 | 12 phút đọc

Chế độ chi tiết: Tắt

Mã GPU giờ đây đã có thể sử dụng SIMD di động (portable SIMD) của Rust. Chúng tôi chia sẻ phương pháp triển khai và những tiềm năng mà nó mở ra cho lập trình GPU.

Tại VectorWare, chúng tôi đang xây dựng công ty phần mềm thuần GPU (GPU-native) đầu tiên. Hôm nay, chúng tôi rất vui mừng thông báo rằng chúng tôi đã có thể sử dụng thành công SIMD di động của Rust (core::simd) trên GPU. Cột mốc này đánh dấu một bước tiến quan trọng trong tầm nhìn của chúng tôi, nhằm giúp các nhà phát triển viết những ứng dụng phức tạp, hiệu năng cao, tận dụng toàn bộ sức mạnh phần cứng GPU bằng cách sử dụng các trừu tượng (abstractions) quen thuộc của Rust.

Tính song song dưới cấp độ luồng (thread)

Khi chúng tôi đưa các luồng Rust lên GPU, chúng tôi đã ánh xạ mỗi std::thread thành một warp trên GPU. Điều này cho phép chúng tôi chạy nhiều luồng đồng thời trên GPU nhưng lại không sử dụng được các làn (lane) song song bên trong mỗi luồng/warp đó.

Trên CPU, trừu tượng cho tính song song bên trong một luồng chính là SIMD. Một lệnh duy nhất sẽ thao tác trên nhiều phần tử dữ liệu được đóng gói vào một đơn vị vector: trong khi mã vô hướng (scalar code) cộng hai số, thì một lệnh cộng SIMD sẽ lấy hai vector, ví dụ gồm tám giá trị f32, và tạo ra tám tổng cùng một lúc. Tính song song dữ liệu này nằm bên trong một luồng đơn lẻ, ở cấp độ thấp hơn cả nơi hệ điều hành lập lịch.

Luồng CPU | Lệnh SIMD | 0 1 2 N ⋯ | Các làn SIMD | Luồng CPU

SIMD di động của Rust

Trước đây, việc viết SIMD trong Rust đồng nghĩa với việc phải sử dụng các intrinsic đặc thù cho từng kiến trúc trong core::arch, chẳng hạn như _mm256_add_ps trên x86-64 hoặc vaddq_f32 trên Arm. Các intrinsic này chỉ dành riêng cho một tập lệnh nhất định, vì vậy một chương trình chạy trên nhiều kiến trúc cần phải có các bản triển khai riêng biệt cho từng loại.

Thay vào đó, SIMD di động của Rust thêm một lớp trừu tượng phía trên các intrinsic này. Nó cung cấp một kiểu generic duy nhất là Simd<T, N> đại diện cho một vector gồm N phần tử kiểu T. Chương trình sẽ viết các phép toán số học, so sánh, rút gọn (reduction) và hoán đổi làn (lane shuffle) một lần duy nhất với Simd, và trình biên dịch sẽ hạ cấp (lower) chúng xuống bất kỳ lệnh vector nào mà CPU mục tiêu hỗ trợ.

Tại VectorWare, chúng tôi nhận ra rằng GPU chỉ là một phần cứng vector khác mà SIMD di động có thể nhắm tới. Điểm cộng là SIMD di động nằm trong core thay vì std, và nó thậm chí không cần đến sự hỗ trợ của std mà chúng tôi đã mang lên GPU.

SIMT chính là SIMD

GPU thực thi theo mô hình mà NVIDIA gọi là SIMT, hay Single Instruction, Multiple Thread (Một lệnh, nhiều luồng). Một warp phát ra một lệnh, và mỗi làn trong số 32 làn của nó sẽ chạy lệnh đó trên dữ liệu riêng của mình. Một lệnh thao tác trên nhiều phần tử dữ liệu chính xác là ý nghĩa của SIMD, và việc đánh địa chỉ theo từng làn mà SIMT thêm vào không làm thay đổi điều đó. Một warp là một đơn vị vector rộng và một vector SIMD di động ánh xạ trực tiếp lên đơn vị đó.

Luồng CPU | 0 1 2 N ⋯ | Các làn SIMD | ≈ | Warp GPU | 0 1 2 N ⋯ | Các làn warp

Ví dụ, một Simd<i16, 32> cung cấp một phần tử i16 cho mỗi làn trong số 32 làn của warp, và việc cộng hai vector như vậy sẽ được biên dịch thành một lệnh warp duy nhất, trong đó mỗi làn cộng phần tử của nó cùng một lúc.

CPU | let a: Simd<i16, 32> = [1, 1, 1,..., 1]; let b: Simd<i16, 32> = [2, 2, 2,..., 2]; let c = a + b; | biên dịch thành vpaddw %zmm2, %zmm1, %zmm0 | a0+b0 (làn 0), a1+b1 (làn 1), a2+b2 (làn 2), a31+b31 (làn 31) ⋯ | println!("{c:?}");

GPU | let a: Simd<i16, 32> = [1, 1, 1,..., 1]; let b: Simd<i16, 32> = [2, 2, 2,..., 2]; let c = a + b; | biên dịch thành add.s16 %rs3, %rs1, %rs2; | a0+b0 (làn 0), a1+b1 (làn 1), a2+b2 (làn 2), a31+b31 (làn 31) ⋯ | println!("{c:?}");

Sự ánh xạ mới này hoàn thiện hệ thống phân cấp song song từ công trình trước đây của chúng tôi. Trên CPU, một luồng chứa các làn SIMD, và trên GPU, std::thread của chúng tôi là một warp mà các làn phần cứng của nó đóng vai trò tương tự. Trong cả hai trường hợp, core::simd đều điều khiển các làn đó.

CPU | ⋯ luồng 0 (0 1 2 N ⋯) ⋯ luồng 1 (0 1 2 N ⋯) ⋯ luồng N (0 1 2 N ⋯) | Các làn SIMD | ≈ | GPU | ⋯ warp 0 (0 1 2 N ⋯) ⋯ warp 1 (0 1 2 N ⋯) ⋯ warp N (0 1 2 N ⋯) | Các làn warp

Lần đầu tiên trên thế giới: core::simd trên GPU

Giống như các bài viết trước, điều này rất khó để hiển thị trực quan vì mã nguồn là Rust thông thường. Các kiểu core::simd giống nhau, vốn hạ cấp xuống SIMD x86-64 trên máy tính xách tay, nay hạ cấp xuống các thao tác warp trên GPU mà không cần thay đổi mã nguồn.

Tại đây, chúng tôi định nghĩa một quy trình SIMD di động nhỏ và gọi nó từ main. Nó thực hiện các tính năng cốt lõi của mô hình: số học theo từng phần tử, một phép so sánh tạo ra mặt nạ làn (lane mask), một phép chọn (select) được điều khiển bởi mặt nạ đó, và một phép rút gọn ngang (horizontal reduction) trên các làn.

Điểm bắt đầu là một hàm fn main bình thường không có chú thích đặc thù cho GPU. Chuỗi công cụ (toolchain) của chúng tôi biên dịch nó thành một GPU kernel, và kết quả được in ra từ thiết bị bằng cách sử dụng sự hỗ trợ std của chúng tôi.

Dưới đây là bản ghi chương trình đang chạy trên GPU, tạo ra kết quả chính xác như khi chạy trên CPU.

Triển khai

Như đã đề cập trước đó, sự ánh xạ này dựa trên một quan sát đơn giản: một warp là một đơn vị vector mà các làn của nó có thể được đánh địa chỉ riêng lẻ. Khi Simd<T, N> được bố trí theo từng làn, mỗi nhóm thao tác sẽ có một đối tác tương ứng ở cấp độ warp.

Các thao tác SIMD theo từng phần tử là trường hợp dễ nhất. Phép cộng, nhân, so sánh và các toán tử theo làn khác đến từ các triển khai trait Rust thông thường trên Simd như Add. GPU chạy chúng một cách tự nhiên.

Các phép rút gọn SIMD như reduce_sum và reduce_max kết hợp mọi làn thành một giá trị vô hướng. Chúng sử dụng các lệnh hoán đổi warp (warp shuffle) của GPU để trao đổi và kết hợp các giá trị giữa các làn, tạo ra cùng một kết quả vô hướng trong mọi làn.

Các phép hoán đổi SIMD giữa các làn, như simd_swizzle! và các phép xoay (rotates), di chuyển các phần tử giữa các làn. Vì một làn SIMD chính là một làn warp GPU, chúng ánh xạ lên cùng các nguyên hàm hoán đổi warp giúp các làn GPU trao đổi dữ liệu hiệu quả.

Các mặt nạ SIMD (SIMD masks) cũng ánh xạ gọn gàng tương tự. Một Mask<T, N> cung cấp một vị ngữ (predicate) cho mỗi làn SIMD. Mask::select thực hiện việc chọn trong mỗi làn warp. Các truy vấn mặt nạ ngang như any và all sử dụng các lệnh vote và ballot của GPU.

Các giá trị vô hướng trong mã xung quanh, chẳng hạn như bộ đếm vòng lặp hoặc hằng số, được tính toán giống hệt nhau bởi mọi làn và do đó chỉ đơn giản là được sao chép trên toàn bộ warp giống như trong CUDA thông thường. Đây là sự phân biệt giữa uniform (đồng nhất) và varying (biến thiên) giống như các ngôn ngữ song song dữ liệu như ISPC đã làm rõ, ngoại trừ việc ở đây nó xuất phát từ chính các kiểu dữ liệu của Rust: một f32 đơn thuần là uniform, một Simd<f32, 32> là varying.

Làm việc với các làn

Điểm duy nhất mà sự trừu tượng và phần cứng không khớp nhau là số lượng làn. Trên CPU, một Simd<T, N> cho phép bất kỳ N nào từ 1 đến 64, nhưng phần cứng GPU có độ rộng cố định: 32 làn trên NVIDIA và 32 hoặc 64 trên AMD. Sự ánh xạ chỉ là một-một khi N khớp với độ rộng đó. Một N nhỏ hơn sẽ để lại một số làn nhàn rỗi, trong khi một N lớn hơn sẽ cung cấp cho một số hoặc tất cả các làn nhiều hơn một phần tử để xử lý.

Khi khối lượng công việc lớn hơn độ rộng của warp, chúng ta cần một cách để chỉ định làn nào làm gì. Việc coi warp như một "cỗ máy" nhỏ của riêng nó sẽ rất hữu ích: một tập hợp cố định các nguyên hàm để di chuyển và kết hợp dữ liệu giữa các làn, cộng với các bất biến về việc làn nào đang hoạt động và mỗi làn chứa bao nhiêu dữ liệu. "Lập trình" cho nó nghĩa là đặt công việc vào các làn theo những quy tắc đó.

Tại VectorWare, chúng tôi cung cấp cho cỗ máy đó một IR (Intermediate Representation). Thay vì một cấu trúc dữ liệu độc lập, chúng tôi mã hóa nó trong hệ thống kiểu của Rust bằng cách sử dụng các kiểu, generics, const generics và trait bounds. Một chương trình được cấu tạo từ các thao tác có kiểu: ballots, shuffles, reductions, scans, gathers, scatters, atomics và strip mining cho các vector rộng hơn warp. Các toán hạng, hình dạng thực thi và dung lượng cũng được định kiểu. Vì các thao tác mang hình dạng của chúng trong các kiểu, nhiều chương trình không hợp lệ sẽ không thể được tạo ra ngay từ đầu.

IR không cần trình thông dịch trên GPU. Mỗi thao tác hạ cấp trực tiếp xuống các lệnh tương ứng với chi phí bằng không so với PTX viết tay. Các kiểu tương tự cho phép chúng tôi chạy nó trên cả CPU. Chúng tôi đã xây dựng một trình thông dịch tham chiếu thực thi IR một cách xác định, một dạng Miri cho lập trình làn warp. Chúng tôi sử dụng nó để mô phỏng mã GPU và để kiểm thử vi sai (differential testing).

Công việc của chúng tôi hiện nhắm vào NVIDIA, nhưng không có gì ở đây là đặc thù của CUDA. Các wavefront của AMD và subgroup của Vulkan cũng cung cấp các nguyên hàm và ngữ nghĩa tương tự. Bản thân IR là Rust bất khả tri về kiến trúc.

Lợi ích