wgpu 光线追踪实战：用计算着色器对加速结构中的 AABB 做硬件光线查询（ray-aabb 示例深度解析）

wgpu 光线追踪实战用计算着色器对加速结构中的 AABB 做硬件光线查询ray-aabb 示例深度解析【免费下载链接】wgpuA cross-platform, safe, pure-Rust graphics API.项目地址: https://gitcode.com/GitHub_Trending/wg/wgpu导读本篇文章围绕 wgpu 仓库中的ray_aabb_compute示例位于 examples/features/src/ray_aabb_compute深入讲解如何在 wgpu 中搭建「BLAS/TLAS 加速结构 WGSL 硬件光线查询」的完整数据通路先用 AABB 几何构建底层加速结构再在计算着色器里通过ray_query对盒子做硬件求交并把每个命中盒转换为一个内切球体进行解析着色。读完本文你将掌握 wgpu 实验性光线追踪 APIEXPERIMENTAL_RAY_QUERY的加速结构创建、光线查询循环、自定义命中处理与计算着色器渲染全流程并能独立跑通和复刻这一示例。示例概览从 AABB 命中到球面着色按 examples/features/src/ray_aabb_compute/README.md 的描述这个示例做的事可以概括为三步从加速结构中对 AABB 几何做光线追踪ray-trace把每一次命中hit转换成球体计算着色器使用**硬件光线查询hardware ray queries**对 AABBs 求交在每个盒子里内切一个球再用**光线-球体求交ray-sphere intersection**解析命中并着色。换句话说示例用「盒子 解析球」的组合演示了 wgpu 中加速结构加速结构内的叶子并不直接存储三角形而是 AABB与 WGSL 内置光线查询函数配合的典型用法——这正是程序化几何procedural geometry光线追踪的核心思路。示例还特别注明了一个平台限制注意Metal 后端目前不支持光线查询中的 AABB 求交因此该示例在 Metal 上无法运行。这一限制也体现在示例的注册表与测试配置中详见下文「平台兼容性与测试」一节。运行方式与示例注册README 给出的运行命令在仓库根目录执行cargo run --bin wgpu-examples ray_aabb_compute示例二进制来自 examples/features/Cargo.toml 中定义的[[bin]] wgpu-examples入口为src/main.rsharness false即该 crate 同时以库和二进制形式存在。示例的派发逻辑在 examples/features/src/main.rs 的EXAMPLES常量表中ray_aabb_compute被注册为wgpu_examples::ray_aabb_compute::main且webgl: false, webgpu: false——即它只面向原生native目标运行不会出现在 Web 平台WebGL/WebGPU的示例列表中与ray_cube_compute、ray_traced_triangle等其他光线追踪示例一致源码注释注明 Web 端尚无光线追踪扩展。示例主入口 examples/features/src/ray_aabb_compute/mod.rs 最后一行调用crate::framework::run::Example(ray-aabb)通过 examples/features/src/framework.rs 的Exampletrait 约定init / resize / update / render生命周期。功能与限制声明示例通过 trait 的required_*系列方法向框架声明自己的硬性需求mod.rs需求取值说明required_features()TEXTURE_BINDING_ARRAY \| VERTEX_WRITABLE_STORAGE \| EXPERIMENTAL_RAY_QUERY必须启用实验性光线查询特性同时需要纹理绑定数组与可写顶点存储框架/测试基础设施所需required_downlevel_capabilities()COMPUTE_SHADERS至少需要计算着色器能力required_limits()Limits::default().using_minimum_supported_acceleration_structure_values()使用驱动支持的最小加速结构相关 limit保证在尽可能多的硬件上可运行加速结构本身不单独对应一个 Feature正如 wgpu/src/documentation/extensions/ray_tracing.rs 明确说明的Acceleration structures do not have a separate feature, instead they are enabled byFeatures::EXPERIMENTAL_RAY_QUERY, unlike vulkan.该文档同时强调这套 API 仍是实验性的、预期会有破坏性变更。数据与资源准备AABB 数据的宿主结构三个盒子在 CPU 侧被描述为GpuAabbmod.rs内存布局与 WGSL 端 shader.wgsl 中的GpuAabb严格对应#[repr(C)] #[derive(Clone, Copy, Pod, Zeroable)] struct GpuAabb { min: [f32; 3], // min_x, min_y, min_z max: [f32; 3], // max_x, max_y, max_z _pad: [f32; 2], // _pad0, _pad1对齐到 32 字节 }初始化代码mod.rs沿 X 轴均匀摆放了三个盒子尺寸均为 2×1×1左侧min(-3.5, -0.5, -0.5)max(-1.5, 0.5, 0.5)中间min(-0.5, -0.5, -0.5)max(0.5, 0.5, 0.5)右侧min(1.5, -0.5, -0.5)max(3.5, 0.5, 0.5)。这个 AABB 缓冲区同时被两个角色使用mod.rslet aabb_buf device.create_buffer_init(wgpu::util::BufferInitDescriptor { label: Some(AABB primitives), contents: bytemuck::cast_slice(aabb_data), usage: wgpu::BufferUsages::BLAS_INPUT | wgpu::BufferUsages::STORAGE, });BLAS_INPUT作为构建 BLAS底层加速结构的几何输入STORAGE供计算着色器以arrayGpuAabb, 3只读访问用于把命中索引primitive_index映射回盒子的几何数据。相机逆矩阵 Uniform示例并不真正“看”三角形网格而是直接从像素反推出世界空间的光线因此 uniform 里放的是视图矩阵与投影矩阵的逆矩阵mod.rs相机位于(0.0, 0.5, 5.0)看向原点使用Mat4::look_at_rh投影使用Mat4::perspective_rh(59°, 宽高比, near0.001, far1000.0)存储view_inverse与proj_inverse两个逆矩阵。离屏渲染目标storage texture光线追踪的结果不是直接画到屏幕而是先写入一张离屏纹理mod.rs尺寸与 surface 相同config.width × config.height格式Rgba8Unorm单 mip、单层、无 MSAAUsage 同时包含TEXTURE_BINDING | STORAGE_BINDING计算着色器把它当 storage texture 写入后续 blit 阶段再以普通纹理采样上屏。搭建加速结构BLAS TLAS创建 BLASAABB 几何先按 AABB 几何声明 BLAS 的尺寸需求mod.rslet aabb_size_desc wgpu::BlasAABBGeometrySizeDescriptor { primitive_count: aabb_data.len() as u32, // 3 个盒子 flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE, }; let blas device.create_blas( wgpu::CreateBlasDescriptor { label: None, flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE, update_mode: wgpu::AccelerationStructureUpdateMode::Build, }, wgpu::BlasGeometrySizeDescriptors::AABBs { descriptors: vec![aabb_size_desc.clone()], }, );几点值得注意PREFER_FAST_TRACE本示例的几何是静态的、每帧只有 TLAS 变换在动因此优先追踪性能ray_tracing.rs 文档建议对移动几何使用PREFER_FAST_BUILD以加快重建AABB 的 BLAS 与三角形 BLAS 不同文档指出对于 AABB 几何加速结构不保证存储 AABB 本身避免引入第三种求交类型但包围它们的“分支”节点会紧密贴合 AABB 以保证追踪性能——这正是后续着色器中rayQueryGetCandidateIntersection返回RAY_QUERY_INTERSECTION_AABB命中、而最终由用户自行提交命中generate/commit的原因。创建 TLAS 并装入实例TLAS 用create_tlas创建max_instances: 1mod.rs。随后通过TlasInstance把 BLAS 挂进 TLASmod.rstlas[0] Some(wgpu::TlasInstance::new( blas, affine_to_rows(Affine3A::from_rotation_translation( Quat::IDENTITY, Vec3::new(0.0, 0.0, 0.0), )), 0, // custom_data 0xff, // mask ));TlasInstance的语义定义在 wgpu/src/api/blas.rs包含 BLAS 引用、行主序 3×4 仿射变换矩阵transform: [f32; 12]、只允许使用低 24 位的custom_data以及用于光线裁剪的 8 位mask——光线查询时命中条件为(shader_cull_mask tlas_instance.mask) ! 0u本例 mask 取0xff表示不裁剪任何实例。affine_to_rowsmod.rs把 glam 的Affine3A按“3×3 旋转 平移”展开成 12 个浮点的行主序数组。构建与每帧重建构建发生在命令编码阶段mod.rsencoder.build_acceleration_structures( iter::once(wgpu::BlasBuildEntry { blas: blas, geometry: wgpu::BlasGeometries::AabbGeometries(vec![wgpu::BlasAabbGeometry { size: aabb_size_desc, stride: mem::size_of::GpuAabb() as wgpu::BufferAddress, // 32 字节步长 aabb_buffer: aabb_buf, primitive_offset: 0, }]), }), iter::once(tlas), );BLAS 与 TLAS 在同一次build_acceleration_structures调用中一起构建。文档强调BLAS 必须在与引用它的 TLAS 相同的构建或更早的构建中完成TLAS 在被着色器使用前必须已构建且其引用的所有 BLAS 也必须在同一次或更早的构建中完成。每帧渲染时mod.rs示例只更新 TLAS 的变换再单独重建 TLASBLAS 迭代器传空self.tlas[0].as_mut().unwrap().transform affine_to_rows(Affine3A::from_rotation_translation( Quat::from_rotation_y(anim_time * 0.4), // 绕 Y 轴旋转 Vec3::new(0.0, 0.0, 0.0), )); let mut encoder device.create_command_encoder(wgpu::CommandEncoderDescriptor { label: None }); encoder.build_acceleration_structures(iter::empty(), iter::once(self.tlas));于是三个盒子会以约 0.4 rad/s 的角速度绕 Y 轴转动但球体形状不变——因为球定义在对象空间旋转由object_to_world矩阵自动带到世界空间。注意 TLAS 引用的是 BLAS而非拷贝文档提醒只要其引用的 BLAS 被修改TLAS 就必须重建本例 BLAS 静态因此只需重建 TLAS。计算着色器硬件光线查询的完整流程着色器源码在 examples/features/src/ray_aabb_compute/shader.wgsl第一行即声明启用扩展enable wgpu_ray_query;wgpu_ray_query是 naga 的 WGSL 扩展开启后可用的内置光线查询函数命名大多沿用 Vulkan详见 wgpu/src/documentation/extensions/ray_tracing.rs 的完整签名与约束。绑定布局group(0) binding(0) var output: texture_storage_2drgba8unorm, write; // 离屏渲染目标 group(0) binding(1) varuniform uniforms: Uniforms; // view_inv / proj_inv group(0) binding(2) var acc_struct: acceleration_structure; // TLAS group(0) binding(3) varstorage, read gpu_aabbs: arrayGpuAabb, 3; // 原始盒子数据对应的 Rust 绑定组在 mod.rs 中按相同顺序填充纹理视图 → uniform 缓冲 →wgpu::BindingResource::AccelerationStructure(tlas)→ AABB 存储缓冲。由像素反推光线计算着色器工作组大小为workgroup_size(8, 8)shader.wgsl每个线程对应一个像素dispatch 尺寸见下一节let target_size textureDimensions(output); let pixel_center vec2f32(global_id.xy) vec2f32(0.5); let in_uv pixel_center / vec2f32(target_size.xy); let d in_uv * 2.0 - 1.0; // 归一化设备坐标 NDC let origin (uniforms.view_inv * vec4f32(0.0, 0.0, 0.0, 1.0)).xyz; let temp uniforms.proj_inv * vec4f32(d.x, d.y, 1.0, 1.0); let dir normalize((uniforms.view_inv * vec4f32(normalize(temp.xyz), 0.0)).xyz);即像素中心 → UV → NDC → 用proj_inv反投影到相机空间方向 → 再用view_inv转到世界空间得到世界空间的光线原点与方向。初始化与遍历循环var rq: ray_query; rayQueryInitialize(rq, acc_struct, RayDesc(0u, 0xFFu, 0.01, 200.0, origin, dir)); while (rayQueryProceed(rq)) { let c rayQueryGetCandidateIntersection(rq); if (c.kind RAY_QUERY_INTERSECTION_AABB) { // ... 自定义解析命中 } }RayDesc的六个字段按文档定义如下字段本例取值含义flags0u光线标志例如FORCE_OPAQUE(0x1)、TERMINATE_ON_FIRST_HIT(0x4)、CULL_OPAQUE(0x40)、SKIP_TRIANGLES(0x100)、SKIP_AABBS(0x200)等cull_mask0xFFu与每个TlasInstance::mask按位与非零才可能命中t_min0.01最近命中距离避免自交/浮点误差t_max200.0最远命中距离origin世界空间相机位置光线起点dir归一化方向t以dir长度为单位计量遍历语义rayQueryProceed返回true表示遇到候选Candidate命中对非透明几何/程序化 AABB 而言随后可以用rayQueryGetCandidateIntersection取得候选信息调用rayQueryProceed会继续之前的遍历直到返回false表示遍历完成。把 AABB 命中解析成球ray-sphere 求交AABB 命中是“体积命中”候选交点的t字段并不可用文档明确t仅对候选三角形命中返回因此示例用球体解析替代if (c.kind RAY_QUERY_INTERSECTION_AABB) { let aabb gpu_aabbs[c.primitive_index]; let min_v vec3f32(aabb.min_x, aabb.min_y, aabb.min_z); let max_v vec3f32(aabb.max_x, aabb.max_y, aabb.max_z); let extent max_v - min_v; let radius min(extent.x, min(extent.y, extent.z)) * 0.45; // 内切球半径略小于内切球 let center_o (min_v max_v) * 0.5; // 对象空间球心 盒子中心 let center_w (c.object_to_world * vec4f32(center_o, 1.0)).xyz; // 变换到世界空间 let t ray_sphere(origin, dir, center_w, radius); if (t 0.0) { rayQueryGenerateIntersection(rq, t); // 提交自定义命中距离 } }这里的几何要点盒子三轴长度为extent取三个方向的最小值再乘0.45作为球半径保证球被盒子完全包裹球心取盒子中点object_to_worldmat4x3f32负责把对象空间的球心变换到世界空间——这正是旋转 TLAS 时球体随之转动的原理。ray_sphereshader.wgsl是标准的二次方程解析解法oc origin - centerb dot(oc, dir)c dot(oc, oc) - r²判别式disc b² - c无实根返回-1否则取最近正根小于0.001的近距离根被跳过防止自相交噪声。提交命中与最终着色rayQueryGenerateIntersection(rq, t)会把距离t的自定义命中提交为程序化几何命中。遍历结束后读取最终结果let hit rayQueryGetCommittedIntersection(rq); if (hit.kind RAY_QUERY_INTERSECTION_GENERATED) { // 重新取出盒子数据计算球心 let p origin dir * hit.t; // 世界空间命中点 let n normalize(p - center_w); // 球面法线 color vec4f32(n * 0.5 vec3f32(0.5), 1.0); // 法线可视化 } textureStore(output, global_id.xy, color);命中类型的三个常量文档RAY_QUERY_INTERSECTION_NONE 0、RAY_QUERY_INTERSECTION_TRIANGLE 1、RAY_QUERY_INTERSECTION_GENERATED 2程序化几何的最终命中只能由GenerateIntersection/ConfirmIntersection产生、RAY_QUERY_INTERSECTION_AABB 3仅在候选阶段出现。示例把球面法线映射到 RGB 输出因此三个“盒子”在截图中呈现为各自独立的球体渐变——未命中的像素保持初始化的背景渐变颜色vec4(uv, 0.0, 1.0)即左上红、右下绿的渐变与截图完全吻合。上屏compute 结果到窗口的 blit 通道计算通道把结果写入离屏 storage texture 后渲染通道用一个全屏三角形把它采样到交换链表面mod.rs// 1) 计算通道每 8×8 像素一个工作组 let mut cpass encoder.begin_compute_pass(wgpu::ComputePassDescriptor { label: None, timestamp_writes: None }); cpass.set_pipeline(self.compute_pipeline); cpass.set_bind_group(0, Some(self.compute_bind_group), []); cpass.dispatch_workgroups(self.rt_target.width() / 8, self.rt_target.height() / 8, 1); // 2) 渲染通道blit 上屏 let mut rpass encoder.begin_render_pass(/* 背景 LoadOp::Clear(GREEN) */); rpass.set_pipeline(self.blit_pipeline); rpass.set_bind_group(0, Some(self.blit_bind_group), []); rpass.draw(0..3, 0..1);blit 顶点着色器blit.wgsl不读顶点缓冲只用builtin(vertex_index)0..3生成覆盖屏幕的三角形并把纹理坐标传给片元着色器采样r_color采样器为线性过滤、ClampToEdge寻址mod.rs。这样“计算着色器追踪 → storage texture → blit 全屏三角形”的流程就完成了。平台兼容性与测试Metal 不可用README 明确指出 Metal 后端尚不支持光线查询中的 AABB 求交因此该示例在 Metal 上无法运行这也解释了为什么注册表中webgl: false, webgpu: false该示例为纯原生、且依赖后端能力。自动化快照测试模块末尾mod.rs注册了一个wgpu_test::gpu_test!测试ray_aabb_compute以 1024×768 渲染并与 examples/features/src/ray_aabb_compute/screenshot.png 做Mean(0.02)的像素均值比较同时disable_mtl_shader_validation()处理 Metal 着色器校验问题对应 gfx-rs/wgpu#9100。这也说明仓库以截图作为该示例的行为基准。延伸阅读完整的 wgpu 光线追踪 API 与 WGSL 光线查询函数签名、RayDesc合法性约束、RayIntersection字段语义 wgpu/src/documentation/extensions/ray_tracing.rsTlasInstance的构造、变换矩阵与 mask 语义 wgpu/src/api/blas.rs加速结构构建命令的 Feature 前置条件 wgpu/src/api/command_encoder.rs同族的其他光线追踪示例三角形、场景、阴影等位于 examples/features/src可对比 AABB 程序化几何与常规三角形几何在构建和着色上的差异。【免费下载链接】wgpuA cross-platform, safe, pure-Rust graphics API.项目地址: https://gitcode.com/GitHub_Trending/wg/wgpu创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考