ARTICLE DETAIL

资讯详情

深耕网站视觉设计与运营推广的一线实战洞察。

wgpu 计算着色器 Workgroup 完全指南:深入理解 @workgroup_size、内置变量与并发执行模型

wgpu 计算着色器 Workgroup 完全指南:深入理解 @workgroup_size、内置变量与并发执行模型 wgpu 计算着色器 Workgroup 完全指南深入理解 workgroup_size、内置变量与并发执行模型【免费下载链接】wgpuA cross-platform, safe, pure-Rust graphics API.项目地址: https://gitcode.com/GitHub_Trending/wg/wgpuhello_workgroups是 wgpu 官方示例集中用于讲解计算着色器compute shader核心概念的入门示例它通过一段略显刻意的代码把workgroup_size属性、dispatch_workgroups派发、workgroup_id/local_invocation_id/global_invocation_id三个内置变量以及 workgroup 的内存与并发语义一次讲透。读完本文你将理解 workgroup 在 WGSL 执行模型中的真实含义、workgroup 之间与 workgroup 内部的并发保证差异并能独立读懂、改造 wgpu 中的任意计算着色器示例。示例概览一段刻意但极具教学价值的代码该示例的入口源码位于 mod.rs着色器位于 shader.wgsl配套讲解文档即本篇文章所依据的 README.md。正如文档开头所说现在你终于明白那个小小的workgroup_size(1)到底是什么意思了这个示例故意设计得非常简单粗暴原文用词是extremely bare-bones and arguably somewhat unreasonable目的是排除其他干扰、聚焦 workgroup 这一个概念。它做的事情如下准备两个长度 100 的i32数组a[i] ib[i] 2i见 mod.rs将两个数组分别绑定到着色器的两个 storage bufferbinding 0 和 binding 1派发 100 个 workgroupdispatch_workgroups(100, 1, 1)每个 workgroup 对应一个数组下标每个 workgroup 内有 2 个 invocation由workgroup_size(2, 1, 1)决定第一个 invocation 给a[wid.x]加 1第二个 invocation 给b[wid.x]加 1最后把两个 storage buffer 拷贝到 staging buffer、映射回 CPU打印输出结果。由于每个 workgroup 只覆盖一个下标却同时操作两个数组源码中特意注释道因为每个 workgroup 会同时覆盖两个数组我们只需要覆盖其中一个数组的长度即可mod.rs。运行方式cargo run --bin wgpu-examples hello_workgroups该示例在 main.rs 中被注册为名为hello_workgroups的可运行示例其webgpu: true, webgl: false标记说明它可在原生端与 WebGPUwasm上运行但不支持 WebGL 后端WebGL 2 没有计算着色器。你还可以在 lib.rs 看到它的模块注册。Workgroup 是什么TLDR / 核心要点先给出一份速览这些是文档反复强调、也是理解后续所有内容的地基Workgroup 位于一个三维网格中一次 dispatch 派发的所有 workgroup 构成一个三维网格维度由dispatch_workgroups(x, y, z)的x、y、z参数决定。同一个 workgroup 内的所有 invocation 保证并发执行这是 workgroup 最核心的语义WGSL 规范保证 workgroup 内所有 invocation 会同时运行。workgroup 之间没有任何并发保证具体表现为三条不能保证任意两个 workgroup 并行执行也不能保证任意两个 workgroup不并行执行任何一组 workgroup 之间都不保证存在任何可预测、可靠的执行顺序。workgroup 的大小由workgroup_size属性定义它修饰计算着色器的入口函数。builtin(local_invocation_id)获取当前 invocation 在其 workgroup 网格内的位置。builtin(global_invocation_id)获取当前 invocation 在整个计算着色器网格内的位置。builtin(workgroup_id)获取当前 invocation 所属的 workgroup 在派发网格中的位置。workgroup 共享workgroup地址空间的内存workgroup 内存与私有private内存类似但在整个 workgroup 内共享——同一 workgroup 内的 invocation 看到的是同一份内存而不同 workgroup 的 invocation 访问的是不同的内存。派发网格Dispatch Griddispatch_workgroups 的三维语义当你调用ComputePass::dispatch_workgroups时实际上是按x、y、z三个参数派发了一个三维的 workgroup 网格。以dispatch_workgroups(5, 2, 1)为例它会产生如下二维可见的派发网格z 1时退化为平面WWWWWWWWWW这里每个W代表一个 workgroup共 5 × 2 10 个。该 API 的定义位于 compute_pass.rs签名为pub fn dispatch_workgroups(mut self, x: u32, y: u32, z: u32)。如果希望着色器知道当前 invocation 位于整个派发网格中的哪个 workgroup就在计算着色器主函数中声明一个vec3u32类型的参数并加上builtin(workgroup_id)属性compute workgroup_size(2, 1, 1) fn main(builtin(workgroup_id) wid: vec3u32) { // wid.x / wid.y / wid.z 即当前 workgroup 在派发网格中的坐标 }一个重要的术语澄清文档特别提醒派发网格dispatch grid并不是 WGSL 的正式术语它只是本文档为了讲解方便而采用的表述。WGSL 中的正式术语有两个workgroup gridworkgroup 网格指单个 workgroup 内部的 invocation 网格compute shader grid计算着色器网格指整个派发中所有invocation 构成的总网格它被划分成一个个 workgroup 区块。Workgroup 内部workgroup_size 与 local_invocation_id在 hello-compute、repeated-compute 这些更早的示例中workgroup 大小是(1, 1, 1)于是每个 workgroup 恰好产生一个 invocation——但这只是特例并非一般情况。实际上每个 workgroup 代表其内部一个独立的小型 invocation 网格这个网格可以只有一个 invocation也可以是任意三维尺寸。网格大小即每个 workgroup 产生的 invocation 数量由workgroup_size属性决定例如workgroup_size(2, 1, 1)表示每个 workgroup 有 2 个 invocation。要获取当前 invocation 在其 workgroup 内的位置声明一个vec3u32参数并加上builtin(local_invocation_id)属性。下面用一个具体例子把三个概念拼在一起假设一次dispatch_workgroups(2, 2, 1)的派发配合workgroup_size(2, 2, 1)令w代表workgroup_id、i代表local_invocation_id则整个计算着色器网格中每个 invocation 的坐标如下表w(0, 0, 0), i(0, 0, 0)w(0, 0, 0), i(1, 0, 0)w(1, 0, 0), i(0, 0, 0)w(1, 0, 0), i(1, 0, 0)w(0, 0, 0), i(0, 1, 0)w(0, 0, 0), i(1, 1, 0)w(1, 0, 0), i(0, 1, 0)w(1, 0, 0), i(1, 1, 0)w(0, 1, 0), i(0, 0, 0)w(0, 1, 0), i(1, 0, 0)w(1, 1, 0), i(0, 0, 0)w(1, 1, 0), i(1, 0, 0)w(0, 1, 0), i(0, 1, 0)w(0, 1, 0), i(1, 1, 0)w(1, 1, 0), i(0, 1, 0)w(1, 1, 0), i(1, 1, 0)可以看到每个 workgroup 内部有 2 × 2 4 个 invocation4 个 workgroup 共 16 个 invocation。w和i的组合可以唯一确定一个 invocation 的位置。示例中的实际用法回到 hello_workgroups 的着色器shader.wgsl它同时使用了local_invocation_id和workgroup_id来让同一个 workgroup 内的两个 invocation 各司其职group(0) binding(0) varstorage, read_write a: arrayi32; group(0) binding(1) varstorage, read_write b: arrayi32; compute workgroup_size(2, 1, 1) fn main(builtin(local_invocation_id) lid: vec3u32, builtin(workgroup_id) wid: vec3u32) { if lid.x 0u { // Do computation (use your imagination) a[wid.x] 1; } else if lid.x 1u { // Do computation b[wid.x] 1; } }着色器注释解释了为何采用两个独立数组而非更常规的vec2arrayi32之类结构arrayT是 unsized 类型不能直接嵌套平时一般会交错存储或用结构体数组这里纯粹是为了演示方便。运行后a变成[1, 2, 3, ..., 100]b变成[1, 3, 5, ..., 199]每个元素各自 1。Workgroup 的执行你能保证什么不能保证什么workgroup 是 invocation 的集合。workgroup 内部的 invocation 永远保证并行执行——这是硬性保证。但除此之外保证基本到此为止你无法保证任何给定 workgroup 何时执行包括它与其他 workgroup 的相对时间关系你不能保证任意两个 workgroup 一定一起执行也不能保证它们不一起执行对于没有一起执行的 workgroup你同样无法保证它们以任何特定顺序执行。换句话说当一个函数在某个 invocation 中运行时你能确定的是它正与自己的workgroup 同伴协同工作——仅此而已。这种无保证正是为硬件实现从单核 CPU 上的串行执行到 GPU 上的大规模并行留出最大自由度。从 wgpu 的 API 层面看dispatch_workgroups只是把参数记录进计算通道的命令流中见 compute_pass.rs真正的并行调度完全交由底层后端Vulkan / Metal / DX12 / WebGPU与 GPU 硬件决定这正对应了workgroup 之间无执行保证的语义。全局视角global_invocation_id 与三个坐标的关系如前所述每个 invocation 同时存在于两个坐标系中它位于workgroup gridworkgroup 内部网格中对应builtin(local_invocation_id)它也位于compute shader grid整个派发网格按 workgroup 划分区块中对应builtin(global_invocation_id)。一个小彩蛋你可能会以为global_invocation_id是由local_invocation_id和workgroup_id计算出来的但实际上正好相反——一切运算都发生在计算着色器网格上workgroup 只是这个网格中被想象出来的区块local_invocation_id和workgroup_id反而是根据 global id 与已知的 workgroup 尺寸反推出来的。正如文档调侃的那样没错我们生活在一个……计算着色器 invocation 的矩阵里。这个冷知识虽不常用却有助于把整个模型装进一个更大的图景。屏障Barrier与 Workgroup并发语义的落点可以说workgroup 最大的价值在于与屏障barrier配合使用。由于屏障机制在 hello-synchronization 示例中已有更完整的讲解这里只做简短说明尽管各同步函数作用于不同的内存地址空间但所有同步函数都作用于 workgroup 层面即同步的是整个 workgroup。更详细的讲解见 hello-synchronization/README.md。在 hello-synchronization 的着色器shaders.wgsl中可以看到一个教科书式的配合用法patient_main先让每个 invocation 对 workgroup 级原子计数器atomicAdd(count, 1u)随后调用workgroupBarrier()确保所有 invocation 都完成累加后才由local_id.x 0u的那个 invocation 读取atomicLoad(count)写入输出。而hasty_main省略了屏障结果就可能读到未完成累加的值——这正是workgroup 内并发 屏障同步语义的直观演示。在 WGSL 中目前有三种同步函数分别对应不同地址空间storageBarrier作用于 storage 地址空间是一个简单屏障workgroupBarrier作用于 workgroup 地址空间是一个简单屏障workgroupUniformLoad同样作用于 workgroup 地址空间但它不只是屏障还承担了更多职责详见 WGSL 规范。进一步阅读对于技术控读者这份长文可能仍留有不少疑问。WebGPU 与 WGSL 的规范都很长而一个不太直观的事实是关于 workgroup 与计算着色器的大多数规范内容都写在 WGSL 规范中而非 WebGPU 规范。以下是指向规范关键章节的入口WGSL 规范中关于 workgroup 的主章节以技术语言定义了重要术语建议所有想深入该主题的读者先读这一节WebGPU 规范中从 API 视角而非 WGSL 视角介绍计算着色器的章节目前仍是 stub但未来有望充实别忘了builtin()内置输入输出builtin inputs/outputs的完整清单。小结通过 hello_workgroups 示例可以归纳出理解 wgpu / WGSL 计算着色器必须掌握的模型一次dispatch_workgroups(x, y, z)派发一个三维的 workgroup 网格每个 workgroup 由workgroup_size决定内部 invocation 网格的尺寸workgroup_id、local_invocation_id、global_invocation_id三个内置变量分别描述 invocation 在派发网格、workgroup 内部、整个计算着色器网格中的位置其中后两者存在先有 global、再反推 local 与 workgroup的推导关系并发保证只存在于 workgroup 内部跨 workgroup 的执行方式与顺序完全不可预期因此跨 workgroup 的数据依赖必须显式处理例如通过独立的 dispatch 或 queue 级同步需要 workgroup 内部协作时配合workgroupBarrier等同步函数使用这是 workgroup 高并发能力发挥作用的关键场景。掌握了这套模型再回头看 hello-compute 中那个每个 workgroup 只有一个 invocation的workgroup_size(1)相信你已经能理解它背后的完整语义了。【免费下载链接】wgpuA cross-platform, safe, pure-Rust graphics API.项目地址: https://gitcode.com/GitHub_Trending/wg/wgpu创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考
返回列表