Rust-GPU/Rust-CUDA 核心知识点：Kernel ABI 详解
项目地址: https://gitcode.com/gh_mirrors/ru/Rust-CUDA 引言
在 GPU 编程中，理解如何正确地向内核函数传递参数至关重要。本文将深入探讨 Rust-GPU/Rust-CUDA 项目中的 Kernel ABI（应用二进制接口），帮助开发者掌握参数传递的正确方式。
基本概念
Kernel ABI 定义了 CPU 端如何向 GPU 内核传递参数。在 Rust-GPU/Rust-CUDA 中，所有内核函数都必须使用 extern "C" 调用约定，这是由 #[kernel] 宏强制执行的。
重要提示：当前文档仅适用于非 Rust 调用约定（特别是 "C" 约定）。虽然项目目前覆盖了除 Rust 外的所有 ABI，但开发者应仅使用 "C" 约定。
结构体传递
结构体作为值传递时，会以字节数组的形式直接传递，这与 CUDA/PTX ABI 的预期一致。
示例：
#[derive(Clone, Copy)]
#[repr(C)]
pub struct ExampleStruct {
pub field1: u16,
pub field2: u64,
pub field3: u128,
}
#[kernel]
pub unsafe fn process_struct(data: ExampleStruct) {
// 处理逻辑
}
对应的 PTX 代码会显示为字节数组参数：
.visible .entry process_struct(
.param .align 16 .b8 process_struct_param_0[32]
)
调用时应直接传递结构体实例，而非引用：
let instance = ExampleStruct {
field1: 1,
field2: 2,
field3: 3
};
unsafe {
launch!(
module.process_struct<<<1, 1, 0, stream>>>(instance)
)?;
}
数组处理
数组的处理方式与结构体类似，总是以值形式作为字节数组传递。
切片传递机制
切片在传递时会被拆分为两个参数：
- 数据起始指针（32位或64位，取决于架构）
- 切片长度
示例内核：
#[kernel]
pub unsafe fn process_slice(data: &[u8]) {
// 处理逻辑
}
对应的 PTX 代码（nvptx64架构）：
.visible .entry process_slice(
.param .u64 process_slice_param_0,
.param .u64 process_slice_param_1
)
调用时需要分别传递指针和长度：
let buffer = [5u8; 10].as_dbuf()?;
unsafe {
launch!(
module.process_slice<<<1, 1, 0, stream>>>(buffer.as_device_ptr(), buffer.len())
)?;
}
注意：虽然编译器可能警告切片不是标准的 C 类型，但这些警告可以安全忽略。
重要限制：不能传递可变切片，因为这会导致别名规则冲突。替代方案包括：
- 使用
&[UnsafeCell<T>]并在确认访问不冲突后转换为可变引用 - 更常见的做法是使用裸指针
特殊类型处理
零大小类型（ZST）
零大小类型会被完全忽略，不会出现在最终 PTX 代码中。
基本数据类型
基本类型直接按值传递，对应 PTX 的特殊类型：
- 有符号整数：.s8, .s16, .s32, .s64
- 无符号整数：.u8, .u16, .u32, .u64
- 浮点数：.f32, .f64
例外情况：
- u128 和 i128 会作为字节数组传递（但这不影响 CPU 端的传递方式）
引用和指针
引用和指针都作为指针传递，调用时需要传递设备内存指针：
#[kernel]
pub unsafe fn process_reference(data: &u8) {
// 处理逻辑
}
调用示例：
let value = DeviceBox::new(&5)?;
unsafe {
launch!(
module.process_reference<<<1, 1, 0, stream>>>(value.as_device_ptr())
)?;
}
重要注意事项
关于 repr(Rust) 类型
虽然不禁止在内核参数中使用 repr(Rust) 类型，但强烈不建议这样做。原因是 rustc 可能在不同编译会话中改变这些类型的表示方式，导致难以追踪的错误。
建议做法：
- 内核参数应只使用 repr(C) 类型
- 切片是个例外，因为它的参数布局是有保证的
最佳实践总结
- 始终使用
#[kernel]宏声明内核函数 - 结构体和数组直接传递值，不要传递引用
- 切片传递指针和长度两个参数
- 避免使用可变切片
- 基本类型直接传递
- 引用和指针传递设备内存地址
- 优先使用 repr(C) 类型作为参数
通过遵循这些准则，可以确保内核参数的正确传递，避免常见的 GPU 编程陷阱。
创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
