Rust高性能计算库cubecl-hip的使用,支持GPU加速与并行计算的HIP后端实现
以下是关于Rust高性能计算库cubecl-hip的使用,支持GPU加速与并行计算的HIP后端实现的完整内容:
![CubeCL Logo]
CubeCL - 多平台高性能计算语言扩展
CubeCL 是一个支持多平台的高性能计算语言扩展,允许您使用Rust编写GPU计算程序,利用零成本抽象来开发可维护、灵活且高效的计算内核。
支持的平台
平台 | 运行时 | 编译器 | 硬件 |
---|---|---|---|
WebGPU | wgpu | WGSL | 大多数GPU |
CUDA | CUDA | C++ (CUDA) | NVIDIA GPU |
ROCm | HIP | C++ (HIP) | AMD GPU |
Metal | wgpu | C++ (Metal) | Apple GPU |
Vulkan | wgpu | SPIR-V | Linux和Windows上的大多数GPU |
示例代码
以下是使用cubecl-hip进行GPU加速计算的完整示例:
use cubecl::prelude::*;
#[cube(launch_unchecked)]
/// 一个[Line]代表一系列连续的元素,可以进行SIMD操作
/// 运行时会自动使用SIMD指令以提高性能
fn gelu_array<F: Float>(input: &Array<Line<F>>, output: &mut Array<Line<F>>) {
if ABSOLUTE_POS < input.len() {
output[ABSOLUTE_POS] = gelu_scalar(input[ABSOLUTE_POS]);
}
}
#[cube]
fn gelu_scalar<F: Float>(x: Line<F>) -> Line<F> {
// 在编译时执行sqrt函数
let sqrt2 = F::new(comptime!(2.0f32.sqrt()));
let tmp = x / Line::new(sqrt2);
x * (Line::erf(tmp) + 1.0) / 2.0
}
然后可以使用自动生成的gelu_array::launch_unchecked
函数启动内核:
pub fn launch<R: Runtime>(device: &R::Device) {
let client = R::client(device);
let input = &[-1., 0., 1., 5.];
let vectorization = 4;
let output_handle = client.empty(input.len() * core::mem::size_of::<f32>());
let input_handle = client.create(f32::as_bytes(input));
unsafe {
gelu_array::launch_unchecked::<f32, R>(
&client,
CubeCount::Static(1, 1, 1),
CubeDim::new(input.len() as u32 / vectorization, 1, 1),
ArrayArg::from_raw_parts::<f32>(&input_handle, input.len(), vectorization as u8),
ArrayArg::from_raw_parts::<f32>(&output_handle, input.len(), vectorization as u8),
)
};
let bytes = client.read_one(output_handle.binding());
let output = f32::from_bytes(&bytes);
// 结果应该是 [-0.1587, 0.0000, 0.8413, 5.0000]
println!("Executed gelu with runtime {:?} => {output:?}", R::name());
}
要运行这个GELU示例,可以使用以下命令:
cargo run --example gelu --features hip # HIP运行时
设计原理
CubeCL围绕立方体(Cubes)设计,更准确地说基于长方体(cuboids),因为并非所有轴都是相同大小。所有计算API都需要映射到硬件,硬件是可以通过3D表示访问的瓦片,因此我们的拓扑可以轻松映射到其他API的概念。
![CubeCL拓扑]
特殊功能
自动向量化
高性能内核应尽可能使用SIMD指令,但实现起来可能相当复杂!使用CubeCL,您可以在启动内核时指定每个输入变量的向量化因子。
编译时计算(Comptime)
CubeCL不仅仅是一种新的计算语言:虽然感觉像是在编写GPU内核,但实际上是在编写可以完全定制的编译器插件!
自动调优(Autotuning)
自动调优通过在运行时运行小型基准测试来确定在当前硬件上运行的最佳内核和配置,从而大大简化了内核选择。
安装
在项目目录中运行以下Cargo命令:
cargo add cubecl-hip
或者在Cargo.toml中添加以下行:
cubecl-hip = "0.6.0"
许可证
CubeCL-hip采用MIT或Apache-2.0双重许可。
完整示例DEMO
以下是一个完整的cubecl-hip使用示例,展示了如何实现向量加法:
use cubecl::prelude::*;
#[cube(launch_unchecked)]
/// 向量加法内核
/// 输入两个向量,输出它们的和
fn vector_add<F: Float>(
a: &Array<Line<F>>,
b: &Array<Line<F>>,
result: &mut Array<Line<F>>
) {
if ABSOLUTE_POS < a.len() {
result[ABSOLUTE_POS] = a[ABSOLUTE_POS] + b[ABSOLUTE_POS];
}
}
fn main() {
// 初始化HIP运行时
let device = cubecl::hip::Device::default();
let runtime = cubecl::hip::Runtime::new(&device);
// 准备输入数据
let a = vec![1.0f32, 2.0, 3.0, 4.0];
let b = vec![4.0f32, 3.0, 2.0, 1.0];
let vectorization = 4;
// 创建GPU缓冲区
let client = runtime.client(&device);
let a_handle = client.create(f32::as_bytes(&a));
let b_handle = client.create(f32::as_bytes(&b));
let result_handle = client.empty(a.len() * core::mem::size_of::<f32>());
// 启动内核
unsafe {
vector_add::launch_unchecked::<f32, _>(
&client,
CubeCount::Static(1, 1, 1),
CubeDim::new(a.len() as u32 / vectorization, 1, 1),
ArrayArg::from_raw_parts::<f32>(&a_handle, a.len(), vectorization as u8),
ArrayArg::from_raw_parts::<f32>(&b_handle, b.len(), vectorization as u8),
ArrayArg::from_raw_parts::<f32>(&result_handle, a.len(), vectorization as u8),
)
};
// 读取结果
let bytes = client.read_one(result_handle.binding());
let result = f32::from_bytes(&bytes);
println!("Vector add result: {:?}", result); // 应该输出 [5.0, 5.0, 5.0, 5.0]
}
要运行这个示例,创建新项目并添加cubecl-hip依赖后,使用以下命令:
cargo run --features hip
Rust高性能计算库cubecl-hip使用指南
完整示例代码
下面是一个完整的demo示例,展示了如何使用cubecl-hip进行向量平方计算:
use cubecl_hip as hip;
fn main() -> Result<(), hip::Error> {
// 1. 初始化HIP环境
let ctx = hip::Context::create()?;
let devices = hip::Device::all()?;
println!("Found {} HIP devices", devices.len());
// 2. 创建计算缓冲区
let buffer_size = 1024 * std::mem::size_of::<f32>();
let device_buffer = hip::Buffer::<f32>::new(&ctx, buffer_size)?;
// 准备输入数据
let mut host_data = vec![0.0f32; 1024];
for i in 0..host_data.len() {
host_data[i] = i as f32;
}
// 将数据复制到设备
device_buffer.write(&host_data)?;
// 3. 准备内核代码(实际项目中应放在单独文件中)
let kernel_code = r#"
extern "C" __global__ void square(float* input, float* output, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
output[idx] = input[idx] * input[idx];
}
}
"#;
// 4. 编译并执行内核
let program = hip::Program::from_source(&ctx, kernel_code)?;
let kernel = program.kernel("square")?;
// 准备输出缓冲区
let mut output = vec![0.0f32; 1024];
let device_output = hip::Buffer::<f32>::new(&ctx, buffer_size)?;
// 设置执行参数
let block_size = 256;
let grid_size = (1024 + block_size - 1) / block_size;
// 执行内核
kernel.launch(
&[grid_size as u32, 1, 1],
&[block_size as u32, 1, 1],
&[&device_buffer, &device_output, &1024]
)?;
// 将结果复制回主机
device_output.read(&mut output)?;
// 打印部分结果验证
println!("Vector square results:");
for i in 0..10 {
println!("{}^2 = {:.1}", host_data[i], output[i]);
}
Ok(())
}
示例说明
- 初始化:创建HIP上下文并获取设备信息
- 内存管理:
- 在设备上创建缓冲区
- 在主机上初始化数据并复制到设备
- 内核执行:
- 定义内核代码(计算平方)
- 编译内核程序
- 设置执行参数(网格和块大小)
- 启动内核计算
- 结果获取:
- 将计算结果从设备复制回主机
- 打印部分结果验证正确性
运行结果示例
程序运行后输出类似以下内容:
Found 1 HIP devices
Vector square results:
0.0^2 = 0.0
1.0^2 = 1.0
2.0^2 = 4.0
3.0^2 = 9.0
4.0^2 = 16.0
5.0^2 = 25.0
6.0^2 = 36.0
7.0^2 = 49.0
8.0^2 = 64.0
9.0^2 = 81.0