2个版本
0.1.1 | 2024年7月19日 |
---|---|
0.1.0 | 2024年4月19日 |
#277 in 算法
187 每月下载量
600KB
15K SLoC
TL;DR
使用CubeCL,您可以利用Rust编程GPU,利用零成本抽象来开发可维护、灵活且高效的计算内核。CubeCL目前完全支持函数、泛型和结构体,对特性和类型推断提供部分支持。随着项目的演进,我们预计将对Rust语言原语提供更广泛的支持,同时保持最佳性能。
示例
只需使用cube
属性标注函数,以指示它们应在GPU上运行。
use cubecl::prelude::*;
#[cube(launch)]
fn gelu_array<F: Float>(input: &Array<F>, output: &mut Array<F>) {
if ABSOLUTE_POS < input.len() {
output[ABSOLUTE_POS] = gelu_scalar::<F>(input[ABSOLUTE_POS]);
}
}
#[cube]
fn gelu_scalar<F: Float>(x: F) -> F {
x * (F::erf(x / F::sqrt(2.0.into())) + 1.0) / 2.0
}
然后,您可以使用自动生成的gelu_array::launch
函数启动内核。
fn launch<R: Runtime>(device: &R::Device) {
let client = R::client(device);
let input = &[-1., 0., 1., 5.];
let output_handle = client.empty(input.len() * core::mem::size_of::<f32>());
gelu_array::launch::<F32, R>(
client.clone(),
CubeCount::Static(1, 1, 1),
CubeDim::new(input.len() as u32, 1, 1),
ArrayArg::new(&client.create(f32::as_bytes(input)), input.len()),
ArrayArg::new(&output_handle, input.len()),
);
let bytes = client.read(output_handle.binding());
let output = f32::from_bytes(&bytes);
// Should be [-0.1587, 0.0000, 0.8413, 5.0000]
println!("Executed gelu with runtime {:?} => {output:?}", R::name());
}
fn main() {
launch::<cubecl::cuda::CudaRuntime>(&Default::default());
launch::<cubecl::wgpu::WgpuRuntime>(&Default::default());
}
要查看其实际运行,请使用以下命令运行可工作的GELU示例
cargo run --example gelu --features cuda # cuda runtime
cargo run --example gelu --features wgpu # wgpu runtime
运行时
目前唯一支持的运行时是CUDA和WebGPU。很容易添加更多,我们打算支持许多;欢迎贡献!
- CUDA
- WebGPU
- Metal(尽管WebGPU原生编译为Metal)
- ROCm
- Vulkan(尽管WebGPU原生编译为Vulkan)
我们还计划开发一个使用SIMD指令优化的JIT CPU运行时,利用Cranelift。
动机
CubeCL的目的是简化编写高度优化的跨硬件可移植的计算内核的痛苦。当您希望在多平台的同时获得最佳性能时,目前还没有合适的解决方案。您必须为不同的硬件编写自定义内核,通常使用不同的语言,如CUDA、Metal或ROCm。为了解决这个问题,我们创建了一个即时编译器,具有三个核心特性:自动向量化、comptime和autotune!
这些功能对于编写高性能内核的人来说非常有用,即使不考虑可移植性也是如此。它们提高了代码的组成性、可重用性、可测试性和可维护性,同时保持最优。CubeCL 还附带了一种针对吞吐量优化的内存管理策略,通过大量重复使用缓冲区来避免分配。
我们的目标不仅在于提供一种优化的计算语言;我们旨在开发一个基于 Rust 的高性能和科学计算生态系统。为了实现这一点,我们正在开发线性代数组件,您可以将其集成到自己的内核中。我们目前有一个高度优化的矩阵乘法模块,在可用的情况下利用 NVIDIA 硬件的 Tensor 核心,同时在其他平台上优雅地回退到基本指令。尽管还有改进的空间,特别是在使用较新 NVIDIA GPU 的自定义指令方面,但我们的实现已经提供了令人印象深刻的表现。
这只是个开始。我们计划包括更多实用程序,例如卷积、随机数生成、快速傅里叶变换和其他基本算法。我们是一个小型团队,也在构建 Burn,所以请不要犹豫,贡献和移植算法;这可以帮助比你想象的还要多!
工作原理
CubeCL 利用 Rust 的 proc 宏系统,采用独特的两步过程
- 解析:proc 宏使用 syn 包解析 GPU 内核代码。
- 扩展:宏不是立即生成中间表示(IR),而是生成一个新的 Rust 函数。
生成的函数,与原始函数在语义上相似,负责在调用时创建 IR。这种方法与传统编译器不同,传统编译器通常在解析后直接生成 IR。我们的方法实现了几个关键特性
- 编译时:由于不转换原始代码,因此可以非常容易地集成编译时优化。
- 自动向量化:通过简单地向量化 CubeCL 函数的输入,我们可以在扩展过程中确定每个中间变量的向量化因子。
- Rust 集成:生成的代码仍然是有效的 Rust 代码,允许它捆绑而没有对特定运行时的任何依赖。
设计
CubeCL 是围绕——没错——立方体设计的!更具体地说,它是基于长方体,因为所有轴的尺寸并不相同。由于所有计算 API 都需要映射到硬件,这些硬件是可以通过三维表示访问的瓦片,因此我们的拓扑可以轻松地映射到其他 API 的概念。
CubeCL - 拓扑
立方体由单元组成,因此一个 3x3x3 的立方体有 27 个单元,可以通过其沿 x、y 和 z 轴的位置来访问。同样,超立方体由立方体组成,正如立方体由单元组成一样。超立方体中的每个立方体都可以通过其相对于超立方体的 x、y 和 z 轴的位置来访问。因此,一个 3x3x3 的超立方体将会有 27 个立方体。在这个例子中,工作单元的总数将是 27 x 27 = 729。
拓扑等价 👇
由于所有拓扑变量都在内核入口点内是常量,因此我们选择使用 Rust 常量语法(大写字母)。通常在创建内核时,我们并不总是关心单元在立方体中沿每个轴的相对位置,但通常我们只关心其在整体中的位置。因此,每种类型的变量也有其自己的轴独立变量,这在其他语言中通常是不存在的,除了 WebGPU 中的 local_invocation_index
。
CubeCL | CUDA | WebGPU |
---|---|---|
CUBE_COUNT | N/A | N/A |
CUBE_COUNT_X | gridDim.x | num_workgroups.x |
CUBE_COUNT_Y | gridDim.y | num_workgroups.y |
CUBE_COUNT_Z | gridDim.z | num_workgroups.z |
CUBE_POS | N/A | N/A |
CUBE_POS_X | blockIdx.x | workgroup.x |
CUBE_POS_Y | blockIdx.y | workgroup.y |
CUBE_POS_Z | blockIdx.z | workgroup.z |
CUBE_DIM | N/A | N/A |
CUBE_DIM_X | blockDim.x | workgroup_size.x |
CUBE_DIM_Y | blockDim.y | workgroup_size.y |
CUBE_DIM_Z | blockDim.z | workgroup_size.z |
UNIT_POS | N/A | local_invocation_index |
UNIT_POS_X | threadIdx.x | local_invocation_id.x |
UNIT_POS_Y | threadIdx.y | local_invocation_id.y |
UNIT_POS_Z | threadIdx.z | local_invocation_id.z |
SUBCUBE_DIM | warpSize | subgroup_size |
ABSOLUTE_POS | N/A | N/A |
ABSOLUTE_POS_X | N/A | global_id.x |
ABSOLUTE_POS_Y | N/A | global_id.y |
ABSOLUTE_POS_Z | N/A | global_id.z |
特殊功能
自动向量化
高性能内核应尽可能依赖于SIMD指令,但这样做可能会很快变得相当复杂!使用CubeCL,您可以在启动内核时指定每个输入变量的向量化因子。在内核代码内部,您仍然只使用一种类型,它是动态向量化的,并支持自动广播。运行时能够编译内核并拥有使用最佳指令所需的所有必要信息!然而,由于算法行为可能取决于向量化因子,CubeCL允许您在需要时直接在内核中访问它,而不会损失任何性能,使用comptime系统!
Comptime
CubeCL不仅仅是一种新的计算语言:尽管感觉您正在编写GPU内核,但实际上您正在编写可以完全自定义的编译器插件!Comptime是在编译内核时第一次编译时修改编译器IR的一种方法。
这可以实现许多优化和灵活性,而无需编写许多相同内核的不同变体以确保最大性能。
特性 | 描述 |
---|---|
指令特殊化 | 并非所有指令都在所有硬件上可用,但当存在专用指令时,应通过简单的if语句启用。 |
自动向量化 | 当您可以使用SIMD指令时,应该使用!但由于并非所有硬件都支持相同的向量化因子,它可以在运行时注入! |
循环展开 | 您可能想要具有相同内核的不同版本,其中只有特定范围内的值才使用循环展开。这可以使用Comptime轻松配置。 |
形状特殊化 | 对于深度学习内核,通常至关重要的是依赖于不同大小的输入的不同内核;您可以通过将形状信息作为Comptime值传递来实现。 |
编译时计算 | 通常,您可以使用Rust运行时属性计算一个常量,并在编译内核时将其注入,以避免在每次执行时重新计算。 |
自动调整
自动调整通过在运行时运行小型基准测试来大大简化内核选择,以确定最佳配置的最佳内核,这对于可移植性是必不可少的。此功能与comptime优雅地结合,以测试不同comptime值对性能的影响;有时这可能令人惊讶!
即使基准测试可能在第一次运行应用程序时增加一些开销,但信息将被缓存在设备上并重用。对于深度学习模型等以吞吐量为导向的程序,这通常是一个不言而喻的权衡。您甚至可以将自动调整缓存与您的程序一起分发,以减少在您对部署目标有更多控制时启动时间。
资源
目前我们没有很多资源可以学习,但您可以通过查看线性代数库来了解CubeCL如何使用。如果您有任何问题或想做出贡献,请不要犹豫,加入Discord。
免责声明 & 历史
CubeCL目前处于alpha阶段。
虽然CubeCL被用于Burn,但仍然存在许多粗糙的边缘;它还没有经过完善。该项目最初是作为Burn的WebGPU专用后端开始的。在我们对其进行优化时,我们意识到我们需要一个中间表示(IR),它可以被优化然后编译到WGSL。拥有一个IR使得支持另一个编译目标变得容易,因此我们创建了一个CUDA运行时。然而,直接在那种IR中编写内核并不容易,因此我们使用syn crate创建了一个Rust前端。在利用两个平台的同时,导航CUDA和WebGPU之间的差异,迫使我们想出了一般概念,这些概念可以在任何地方工作。因此,CubeCL诞生了!
依赖项
~30–61MB
~1M SLoC