Vortex 算子开发

Vortex 算子开发一、概述1.1 什么是 Vortex1.2 四种开发路径总览二、工具链与运行管线2.1 工具链组成2.2 编译 / 运行管线(以 vxcc 为例)2.3 驱动与后端(VORTEX_DRIVER)三、环境准备3.1 前置条件3.2 通用环境变量(CUDA C / CUTLASS / Rust 三路径)3.3 Triton 路径的额外环境四、方案演示4.1 CUDA C 开发4.1.1 原理4.1.2 创建 Sample4.1.3 编译4.1.4 运行与预期输出4.2 用 CUTLASS 开发4.2.1 原理4.2.2 最小 demo:SIMT GEMM4.2.3 编译4.2.4 运行与预期输出4.3 用 Triton 开发4.3.1 原理4.3.2 最小 demo:vecadd4.3.3 运行与预期输出4.4 用 Rust 开发4.4.1 原理4.4.2 创建 Sample4.4.3 编译4.4.4 运行与预期输出五、常见问题(FAQ)六、方案选型建议七、参考资料文档定位本文以「向量加法」「GEMM」等最小可运行算子为例,演示在 Vortex 上用四条主流路径开发 Kernel 的完整流程:CUDA C、CUTLASS、Triton、Rust。适用基线Vortex 3.0 开发分支;交付树llvm_vortex_release(含vxcc与工具链共享库);运行时build/sw/runtime(含libvortex.so与simx/rtlsim两个后端)。一、概述1.1 什么是 VortexVortex 是一个基于RISC‑V的开源 GPGPU 平台,提供类 CUDA / OpenCL 的编程模型:host 侧用标准语言(C / C / Python / Rust)编写启动逻辑,device 侧编写并行kernel;编译器把二者拆分为host ELF与内嵌 kernel 镜像,运行时(libvortex.so)把 kernel 分发到功能级模拟器(simx)或 RTL 级仿真器(rtlsim)上执行。对算子开发者而言,最核心的特性是CUDA 源码零改动:__global__kernel chevron 启动语法(kernelgrid, block(args))直接可用——CUDA 兼容运行时与编译器的 kargs-unpack pass 负责设备端 ABI 与参数打包,源码与写给 NVIDIA GPU的完全一致。1.2 四种开发路径总览路径语言编译器 / 驱动编程模型Vortex 专有代码典型场景CUDA CC / Cvxcc__global__grid, block无通用算子、迁移既有 CUDA 代码CUTLASSC / C(模板)vxccCUTLASS 模板化 GEMM无(仅依赖头文件)矩阵乘等高性能 kernelTritonPythontriton.jit Vortex Triton 后端块级(block-level)SIMT无快速原型、数据并行算子、PyTorch 集成RustRustvxrustc#[no_mangle] extern C chevron 启动宏少量(cfg门控 启动宏)需要强类型 / 内存安全的 hostdevice 工程四条路径最终都产出单个可执行 ELF(kernel 镜像内嵌,VXSYMTAB尾标),运行时统一经libvortex.so启动,与具体后端(simx/rtlsim)解耦。二、工具链与运行管线2.1 工具链组成组件位置作用vxcc$LLVM_VORTEX/bin/vxcc(C 原生二进制)nvcc 风格的单命令驱动。一个源文件同时含 host 与 device 代码,编译器按__kernel/__global__自动拆分:device 函数编入内嵌 kernel 镜像并从 host ELF 剔除,其余函数编入 host ELF。chevron 语法被前端改写为vx_enqueue_launch()调用,链接 host 运行时后产出单 ELF,运行时无需额外的.vxbin文件。vxrustc$VORTEX_HOME/ci/vxrustc(bash 脚本)Rust 版驱动。流水线:rustc产出设备 IR →vxmark.py打 kernel 标记 →vxcc --device-only走既有.ll设备流 → 生成内嵌 blob(带 C mangled__vx_detail符号)→ host 侧rustc链接设备内建 shim 与 chevron 启动 proc-macro,产出单 ELF。Vortex Triton 后端triton/(由build_triton.sh构建)triton.jitkernel 经TTIR → TTGIR → LLVM IR(.ll 文本)→ vxbin编译(TTGIR → LLVM IR由 triton_vortex 插件的 emitter 完成,vxcc --device-only -x ir出 vxbin),经libvortex.so启动;torch_vortex扩展注册vortextorch 设备(privateuse1后端)。libvortex.sobuild/sw/runtime/libvortex.sohost 运行时。按$VORTEX_DRIVER加载后端库(libsimx.so/librtlsim.so),负责设备初始化、host↔device 内存拷贝、kernel 启动与同步。2.2 编译 / 运行管线(以vxcc为例)源文件(host 与 kernel 混写,无需 #ifdef 分区) │ vxcc -o out main.cu [libvortex.so] ├─ host pass : -fvortex-launch 前端,chevron → vx_enqueue_launch(),编入 ELF └─ device pass : __global__/__kernel 函数 → 机器码 → kernel 镜像(内嵌,VXSYMTAB) ▼ 单一可执行 ELF(运行时不再依赖外部 .vxbin) │ VORTEX_DRIVERsimx ./out ▼ libvortex.so → 按 $VORTEX_DRIVER 加载后端(simx | rtlsim)→ 执行2.3 驱动与后端(VORTEX_DRIVER)取值后端库说明simx(默认)libsimx.soC 功能级模拟器,速度快,用于功能验证与日常开发。rtlsimlibrtlsim.soVerilator RTL 级仿真,逐周期精确,用于时序 / 数值精确比对。不设置VORTEX_DRIVER时默认simx。两条后端对同一程序的数值应逐位一致(rtlsim速度显著慢于simx,属预期行为)。三、环境准备3.1 前置条件Vortex 工具链与运行时已构建:即./one_fast.sh至少跑过一次。若需使用 Triton 路径:build_triton.sh已构建 Triton 后端与torch_vortex扩展(套件入口为./test_triton.sh)。下文路径约定:VORTEX_HOME/data/vortex。3.2 通用环境变量(CUDA C / CUTLASS / Rust 三路径)exportVORTEX_HOME/data/vortexexportLLVM_VORTEX$VORTEX_HOME/llvm_vortex_releaseexportPATH$LLVM_VORTEX/bin:$VORTEX_HOME/ci:$PATHexportLD_LIBRARY_PATH$LLVM_VORTEX/lib:$VORTEX_HOME/build/sw/runtime:$LD_LIBRARY_PATH变量值作用VORTEX_HOME/data/vortex仓库根目录(工具链 / 模拟器 / 设备库所在)。LLVM_VORTEX$VORTEX_HOME/llvm_vortex_release交付树:bin/下有vxcc,lib/下有工具链共享库。PATH前置$LLVM_VORTEX/bin:$VORTEX_HOME/ci使vxcc(来自bin/)与vxrustc(来自ci/)可直接调用。LD_LIBRARY_PATH前置$LLVM_VORTEX/lib:$VORTEX_HOME/build/sw/runtime运行时定位工具链共享库与libvortex.so/ 后端库;同时前置可避免旧LD_LIBRARY_PATH遮蔽新版 clang 共享库。VORTEX_DRIVER是运行时变量,逐次启动时指定即可,无需预置(缺省simx)。3.3 Triton 路径的额外环境Triton 路径需在通用环境之上再source一份专用环境脚本:source$VORTEX_HOME/triton.env该脚本(仅做export,不依赖调用方 cwd)额外设置:PYTHONPATH前置triton/python(真正的 triton 包,带__init__.py;刻意不把仓库根加进PYTHONPATH,避免根目录下无__init__.py的triton/子目录以 namespace package 遮蔽真包);TRITON_DEFAULT_BACKENDvortex(显式选择 Vortex 后端);TRITON_BACKENDS_IN_TREE1(triton 未pip install,需 in-tree 后端发现);TRITON_SKIP_AMD_CODEGEN1、TRITON_BUILD_PROTONOFF(规避本机无关依赖)。四、方案演示四个示例彼此独立,可按需任选其一验证。每个示例结尾都做自检:成功打印PASS: ...并以0退出,失败打印FAIL: ...并以非零码退出。4.1 CUDA C 开发4.1.1 原理最小的纯 CUDA C 向量加法out[i] a[i] b[i]。kernel 是标准 CUDA C——普通多参__global__函数,chevron 语法启动,没有任何 Vortex 专有代码。CUDA 兼容运行时 编译器的 kargs-unpack pass 负责设备端 ABI,源码与写给NVIDIA GPU 的完全一致。4.1.2 创建 Samplecat cudac_main.cu EOF #include cuda_runtime.h #include cstdio #include stdlib.h __global__ void vecAdd(const float* a, const float* b, float* out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) out[i] a[i] b[i]; } #define CHECK_CU(expr) \ do { \ cudaError_t _e (expr); \ if (_e ! cudaSuccess) { \ printf(FAIL %s:%d: %s - %s\n, __FILE__, __LINE__, \ #expr, cudaGetErrorString(_e)); \ exit(1); \ } \ } while (0) int main() { const int n 4096; float* h_a new float[n]; float* h_b new float[n]; float* h_out new float[n]; for (int i 0; i n; i) { h_a[i] (float)(i % 97); h_b[i] (float)((i * 3) % 89); } float *d_a, *d_b, *d_out; CHECK_CU(cudaMalloc((void**)d_a, n * sizeof(float))); CHECK_CU(cudaMalloc((void**)d_b, n * sizeof(float))); CHECK_CU(cudaMalloc((void**)d_out, n * sizeof(float))); CHECK_CU(cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice)); CHECK_CU(cudaMemcpy(d_b, h_b, n * sizeof(float), cudaMemcpyHostToDevice)); const int block 128; const int grid (n block - 1) / block; vecAddgrid, block(d_a, d_b, d_out, n); CHECK_CU(cudaDeviceSynchronize()); CHECK_CU(cudaMemcpy(h_out, d_out, n * sizeof(float), cudaMemcpyDeviceToHost)); int bad 0; for (int i 0; i n; i) if (h_out[i] ! h_a[i] h_b[i]) bad; if (bad) { std::printf(FAIL: %d/%d elements wrong\n, bad, n); return 1; } std::printf(PASS: cuda-vecadd %d elements\n, n); CHECK_CU(cudaFree(d_a)); CHECK_CU(cudaFree(d_b)); CHECK_CU(cudaFree(d_out)); delete[] h_a; delete[] h_b; delete[] h_out; return 0; } EOF4.1.3 编译# vxcc 单命令编译;末尾的 *.so(host 运行时库)为可选参数——# 省略时 vxcc 默认链接交付树里的 libvortex.a。vxcc-ocudac_main cudac_main.cu4.1.4 运行与预期输出VORTEX_DRIVERsimx ./cudac_main[VXDRV] Device ctor: ALLOC_BASE_ADDR0x10000, GLOBAL_MEM_SIZE0x200000000, VM_ENABLED compile0 [VXDRV] pinned_size_0x0 PASS: cuda-vecadd 4096 elements4.2 用 CUTLASS 开发4.2.1 原理标准 CUTLASS 模板化 GEMMD alpha·A·B beta·C,零 Vortex 专有代码,仅额外依赖 CUTLASS 头文件。本例使用 SIMT(CUDA-core)GEMM,f32 输入 / f32 输出,threadblock tile32×64×8,以 CPU 参考值逐元素比对(容差1e-4)。4.2.2 最小 demo:SIMT GEMMcat cutlass_main.cu EOF #include cmath #include cstdio #include vector #include cuda_runtime.h #include cutlass/cutlass.h #include cutlass/layout/matrix.h #include cutlass/gemm/device/gemm.h #define CUDA_CHECK(call) \ do { \ cudaError_t e_ (call); \ if (e_ ! cudaSuccess) { \ std::printf(CUDA error: %s (line %d)\n, cudaGetErrorString(e_), \ __LINE__); \ std::exit(1); \ } \ } while (0) // SIMT (CUDA-core) GEMM, f32 in / f32 out, threadblock tile 32x64x8 using Gemm cutlass::gemm::device::Gemm float, cutlass::layout::RowMajor, // A [M,K] row-major float, cutlass::layout::ColumnMajor, // B [K,N] column-major float, cutlass::layout::RowMajor, // C/D [M,N] row-major float, // ElementAccumulator cutlass::arch::OpClassSimt, cutlass::arch::Sm80, cutlass::gemm::GemmShape32, 64, 8; int main() { const int M 8, N 8, K 8; const float alpha 1.0f, beta 0.0f; std::vectorfloat hA(size_t(M) * K), hB(size_t(N) * K), hC(size_t(M) * N, 0.0f); for (int i 0; i M * K; i) hA[i] 0.1f * float(i); for (int i 0; i N * K; i) hB[i] 0.05f * float(i); // CPU reference: D[m][n] sum_k A[m][k]*B[k][n] std::vectorfloat ref(size_t(M) * N); for (int m 0; m M; m) for (int n 0; n N; n) { float acc 0.0f; for (int k 0; k K; k) acc hA[m * K k] * hB[n * K k]; ref[m * N n] acc; } float *dA, *dB, *dC; CUDA_CHECK(cudaMalloc(dA, sizeof(float) * M * K)); CUDA_CHECK(cudaMalloc(dB, sizeof(float) * N * K)); CUDA_CHECK(cudaMalloc(dC, sizeof(float) * M * N)); CUDA_CHECK(cudaMemcpy(dA, hA.data(), sizeof(float) * M * K, cudaMemcpyHostToDevice)); CUDA_CHECK(cudaMemcpy(dB, hB.data(), sizeof(float) * N * K, cudaMemcpyHostToDevice)); CUDA_CHECK(cudaMemcpy(dC, hC.data(), sizeof(float) * M * N, cudaMemcpyHostToDevice)); Gemm op; Gemm::Arguments args{ {M, N, K}, {dA, Gemm::LayoutA(K)}, {dB, Gemm::LayoutB(K)}, {dC, Gemm::LayoutC(N)}, {dC, Gemm::LayoutC(N)}, {alpha, beta}, 1}; // split_k size_t ws Gemm::get_workspace_size(args); void* dws nullptr; if (ws) CUDA_CHECK(cudaMalloc(dws, ws)); if (op.initialize(args, dws) ! cutlass::Status::kSuccess) { std::printf(FAIL: initialize\n); return 1; } if (op() ! cutlass::Status::kSuccess) { std::printf(FAIL: run (CUDA: %s)\n, cudaGetErrorString(cudaGetLastError())); return 1; } CUDA_CHECK(cudaDeviceSynchronize()); std::vectorfloat got(size_t(M) * N); CUDA_CHECK(cudaMemcpy(got.data(), dC, sizeof(float) * M * N, cudaMemcpyDeviceToHost)); bool ok true; for (int i 0; i M * N; i) if (std::fabs(ref[i] - got[i]) 1e-4f) ok false; if (!ok) { std::printf(FAIL: D[0][0] expected %.6f got %.6f\n, ref[0], got[0]); return 1; } std::printf(PASS: gemm %dx%dx%d D[0][0]%.4f\n, M, N, K, got[0]); return 0; } EOF4.2.3 编译# vxcc 编译;-isystem 指定 cutlass 核心头目录 tools/utilvxcc -isystem/data/vortex/cutlass/include\-isystem/data/vortex/cutlass/tools/util/include\-ocutlass_main cutlass_main.cu4.2.4 运行与预期输出VORTEX_DRIVERsimx ./cutlass_main[VXDRV] Device ctor: ALLOC_BASE_ADDR0x10000, GLOBAL_MEM_SIZE0x200000000, VM_ENABLED compile0 [VXDRV] pinned_size_0x0 PASS: gemm 8x8x8 D[0][0]0.70004.3 用 Triton 开发4.3.1 原理张量放在 Vortex 设备(torchprivateuse1后端,设备名vortex)上;triton.jitkernel 由Vortex Triton 后端编译(TTIR → TTGIR → LLVM IR(.ll 文本)→ vxbin),经libvortex.so启动。相比 CUDA C,Triton 以块(block)为并行单位书写,边界用 mask 处理,适合快速原型与数据并行算子,且天然嵌入 PyTorch 生态。4.3.2 最小 demo:vecaddcattriton_main.pyEOFimportos os.environ.setdefault(TRITON_DEFAULT_BACKEND,vortex)# IMPORTANT: import triton BEFORE adding /data/vortex to sys.path,# otherwise the repos triton/ dir shadows the installed package.importtritonimporttriton.languageastlimportsys VHos.environ.get(VORTEX_HOME,/data/vortex)sys.path.insert(0,VH)importtorchimporttorch_vortex# noqa: F401 (registers the vortex torch device)triton.jitdefvecadd_kernel(x_ptr,y_ptr,out_ptr,n,BLOCK:tl.constexpr):pidtl.program_id(0)offspid*BLOCKtl.arange(0,BLOCK)maskoffsn xtl.load(x_ptroffs,maskmask)ytl.load(y_ptroffs,maskmask)tl.store(out_ptroffs,xy,maskmask)defmain():N4096BLOCK128xtorch.arange(N,dtypetorch.float32)*0.01ytorch.arange(N,dtypetorch.float32)*-0.005xvx.to(vortex)yvy.to(vortex)outtorch.empty(N,devicevortex)grid(triton.cdiv(N,BLOCK),)vecadd_kernel[grid](xv,yv,out,N,BLOCKBLOCK)torch.vortex.synchronize()refxy gotout.cpu()err(got-ref).abs().max().item()print(fvecadd: max|err| {err:.3e}(N{N}))asserterr1e-5,fvecadd mismatch (max|err|{err:.3e})print(PASS: vecadd triton-on-vortex)if__name____main__:main()EOF4.3.3 运行与预期输出# 需先 source $VORTEX_HOME/triton.env(见 3.3)VORTEX_DRIVERsimx python3 triton_main.pyvecadd: max|err| 0.000e00 (N4096) PASS: vecadd triton-on-vortexTriton kernel 首次运行会触发后端编译并写入~/.triton缓存;后续运行命中缓存、显著更快。改 kernel 源码后建议清理缓存再复测。4.4 用 Rust 开发4.4.1 原理Rust 路径同样把kernel 与 host 写在同一文件,用cfg做分区:#[cfg(vx_kernel)]标记的设备模块在「设备构建」下编译,#[cfg(not(vx_kernel))]标记的 host 部分(main、启动逻辑)仅在 host 构建下编译——与 C 流在 AST层用-fvortex-kernel-only拆分等价。kernel:直接标注#[no_mangle]的pub extern C fn name即被视为kernel(vxrustc据此自动发现,可用--kernels覆盖);启动:host 侧经 chevron 启动 proc-macrovxk!(name)生成同名启动宏,调用处写成name!(grid, block(args),与 CUDA 语法观感一致;公共部分:host/device 共用的设备内建(vxr_*)与 host API 封装集中在common.rs,经mod common;复用(设备内建由sw/rust/vx_rust_shim.c的设备节提供)。4.4.2 创建 Samplecatrust_main.rsEOF#![allow(dead_code)]#[macro_use]modcommon;#[cfg(not(vx_kernel))]usevxk_macro::vxk;#[cfg(not(vx_kernel))]vxk!(vecadd);#[cfg(vx_kernel)]pubmodkx{usecrate::common::device::*;/// d[i] a[i] b[i]#[no_mangle]pubexternCfnvecadd(d:*mutf32,a:*constf32,b:*constf32,n:u32){unsafe{letidxvxr_block_id_x()*vxr_block_dim_x()vxr_thread_id_x();ifidxn{*d.add(idxasusize)*a.add(idxasusize)*b.add(idxasusize);}}}}#[cfg(not(vx_kernel))]fnmain(){usecommon::ffi;usecommon::host;usecommon::host::bytemsg;letn:u321024;leth_a:Vecf32(0..n).map(|i|(iasf32)*0.25).collect();leth_b:Vecf32(0..n).map(|i|(iasf32)*0.50).collect();leth_ref:Vecf32h_a.iter().zip(h_b).map(|(x,y)|xy).collect();letahost::buffer(ffi::VX_MEM_READ,nasu64*4);letbhost::buffer(ffi::VX_MEM_READ,nasu64*4);letdhost::buffer(ffi::VX_MEM_WRITE,nasu64*4);let(a_addr,b_addr,d_addr)(host::addr(a),host::addr(b),host::addr(d));host::wait(host::upload(a,bytemsg(h_a)));host::wait(host::upload(b,bytemsg(h_b)));let(grid,block)host::max_occupancy(1,[n]);letevvecadd!(grid,block(d_addras*mutf32,a_addras*constf32,b_addras*constf32,n));letmuth_dvec![0f32;nasusize];host::wait(host::download(d,muth_d,Some(ev)));foriin0..nasusize{// Dyadic values → the sum is exact; compare bitwise.ifh_d[i]!h_ref[i]{eprintln!(vecadd mismatch at {i}: got {} want {},h_d[i],h_ref[i]);std::process::exit(1);}}host::release(a);host::release(b);host::release(d);println!(vecadd: PASSED (n {n}));}EOF4.4.3 编译# common.rs 不在本目录,用 -I 指向其所在目录(vxrustc 会把该目录下的# *.rs 与源文件组装进一个临时 module root 再分别做设备 / host 两次 rustc)。vxrustc-orust_main rust_main.rs-I$VORTEX_HOME/tests/rust_compat4.4.4 运行与预期输出VORTEX_DRIVERsimx ./rust_main[VXDRV] Device ctor: ALLOC_BASE_ADDR0x10000, GLOBAL_MEM_SIZE0x200000000, VM_ENABLED compile0 [VXDRV] pinned_size_0x0 vecadd: PASSED (n 1024)五、常见问题(FAQ)现象原因 / 处理运行时报error while loading shared libraries: libvortex.so/libsimx.soLD_LIBRARY_PATH未包含build/sw/runtime(见 3.2)。command not found: vxcc/vxrustcPATH未前置$LLVM_VORTEX/bin与$VORTEX_HOME/ci(见 3.2)。rtlsim比simx慢很多预期行为:RTL 级仿真逐周期建模,仅用于精确比对;日常开发用simx。Triton:ModuleNotFoundError: No module named triton未source triton.env,PYTHONPATH缺triton/python。Triton:import 到的不是 Vortex 的 tritonimport triton必须先于把/data/vortex加入sys.path(见 4.3.2 注释);并确认TRITON_BACKENDS_IN_TREE1。Rust:file not found for module common-I未指向common.rs所在目录(见 4.4.3)。改 kernel 后结果未变化(Triton)命中~/.triton编译缓存,清理缓存或改TRITON_CACHE_PATH后复测。六、方案选型建议迁移既有 CUDA kernel / 通用逐元素、归约类算子→CUDA C(vxcc),源码零改动,上手成本最低。矩阵乘、卷积等需要模板化 tile 的高性能 kernel→CUTLASS(vxccCUTLASS 头文件),直接复用 CUTLASS 的算子库与调优经验。快速原型、块级数据并行算子、需与 PyTorch 张量无缝集成→Triton,Python 书写、块级抽象,迭代速度快。追求强类型、内存安全、可维护的 hostdevice 工程→Rust(vxrustc),以cfg分区 chevron 宏保持 CUDA 观感,同时获得 Rust 的类型系统。四条路径共享同一套工具链、运行时与后端,产出形态一致(单 ELF 内嵌 kernel镜像),可在同一工程内按需混用。七、参考资料docs/vortex-cuda-demo.md— CUDA C 路径详解docs/vortex-triton-demo.md— Triton 路径详解vortex-cutlass-adaptation-technical-report.md— CUTLASS 适配技术报告triple_chevron_syntax.md— chevron 启动语法说明tests/rust_compat/README.md— Rust 路径(vxrustc)与common.rs约定README.md— Vortex 部署与端到端验证总览