WebGPU 的新变化 (Chrome 123)

François Beaufort
François Beaufort
GitHub

WGSL 中对 DP4a 内置函数的支持

DP4a（4 个元素的点积和累加）是指深度学习推理中常用于量化的一组 GPU 指令。它可以高效地执行 8 位整数点积运算，从而加快此类 int8 量化模型的计算速度。与 f32 版本相比，它可以节省高达 75% 的内存和网络带宽，并提升任何机器学习模型在推理方面的性能。因此，很多热门 AI 框架现在都广泛使用它。

当 navigator.gpu.wgslLanguageFeatures 中存在 "packed_4x8_integer_dot_product" WGSL 语言扩展时，您现在可以使用 32 位整数标量来封装 8 位整数的 4 分量矢量，并将其用作 WGSL 着色器代码中点积指令的输入，同时使用 dot4U8Packed 和 dot4I8Packed 内置函数。您还可以将打包和解包指令与 8 位整数打包的 4 分量矢量与 pack4xI8、pack4xU8、pack4xI8Clamp、pack4xU8Clamp、unpack4xI8 和 unpack4xU8 WGSL 内置函数结合使用。

建议在 WGSL 着色器代码顶部使用 requires-directive 来指明 requires packed_4x8_integer_dot_product; 可能无法移植。请参阅以下示例和问题 tint:1497。

if (!navigator.gpu.wgslLanguageFeatures.has("packed_4x8_integer_dot_product")) {
  throw new Error(`DP4a built-in functions are not available`);
}

const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();

const shaderModule = device.createShaderModule({ code: `
  requires packed_4x8_integer_dot_product;

  fn main() {
    const result: u32 = dot4U8Packed(0x01020304u, 0x02040405u); // 42
  }`,
});

特别感谢 Intel 位于上海的 Web 图形团队推动这一规范和实现顺利完成！

WGSL 中的无限制指针参数

"unrestricted_pointer_parameters" WGSL 语言扩展放宽了对可传递到 WGSL 函数的指针的限制：

  • storage、uniform 和 workgroup 地址空间的参数指针指向用户声明的函数。

  • 将指向结构成员和数组元素的指针传递给用户声明的函数。

如需详细了解，请参阅将指针用作函数参数 | WGSL 导览。

此功能可以使用 navigator.gpu.wgslLanguageFeatures 进行特征检测。建议您始终在 WGSL 着色器代码顶部使用 requires-directive 来指明可能无法移植到 requires unrestricted_pointer_parameters;。请参阅以下示例、WGSL 规范变更和问题 tint:2053。

if (!navigator.gpu.wgslLanguageFeatures.has("unrestricted_pointer_parameters")) {
  throw new Error(`Unrestricted pointer parameters are not available`);
}

const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();

const shaderModule = device.createShaderModule({ code: `
  requires unrestricted_pointer_parameters;

  @group(0) @binding(0) var<storage, read_write> S : i32;

  fn func(pointer : ptr<storage, i32, read_write>) {
    *pointer = 42;
  }

  @compute @workgroup_size(1)
  fn main() {
    func(&S);
  }`
});

WGSL 中用于解引用复合体的语法糖

当 navigator.gpu.wgslLanguageFeatures 中存在 "pointer_composite_access" WGSL 语言扩展时，您的 WGSL 着色器代码现在支持使用同一点 (.) 语法访问复杂数据类型的组件，无论您是直接处理数据还是使用指向这些数据的指针。具体方法如下：

  • 如果 foo 是指针，使用 foo.bar 编写 (*foo).bar 会更方便。通常需要使用星号 (*) 将指针转换为可解引用的“引用”，但现在指针和引用都更相似，并且几乎可以互换。

  • 如果 foo 不是指针：点 (.) 运算符的运作方式与您直接访问成员的方式完全相同。

同样，如果 pa 是存储数组起始地址的指针，则使用 pa[i] 可让您直接访问存储该数组的第 'i 个元素的内存位置。

建议使用 requires-directive，通过 WGSL 着色器代码顶部的 requires pointer_composite_access; 表明可能不可移植。请参阅以下示例和问题 tint:2113。

if (!navigator.gpu.wgslLanguageFeatures.has("pointer_composite_access")) {
  throw new Error(`Pointer composite access is not available`);
}

const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();

const shaderModule = device.createShaderModule({ code: `
  requires pointer_composite_access;

  fn main() {
    var a = vec3f();
    let p : ptr<function, vec3f> = &a;
    let r1 = (*p).x; // always valid.
    let r2 = p.x; // requires pointer composite access.
  }`
});

为模板和深度方面提供单独的只读状态

以前，渲染阶段中的只读深度-模板附件要求深度和模板两个方面都为只读。此限制现已解除。现在，您可以以只读方式使用深度方面，例如用于接触阴影跟踪，同时写入模板缓冲区以识别要进一步处理的像素。请参阅问题 dawn:2146。

Dawn 更新

现在，当错误发生时，系统会立即调用通过 wgpuDeviceSetUncapturedErrorCallback() 设置的未捕获错误回调。这是开发者在调试时始终期望和希望的。请参阅 change dawn:173620。

已实现 webgpu.h API 中的 wgpuSurfaceGetPreferredFormat() 方法。请参阅问题 dawn:1362。

本文仅介绍了一些主要亮点。查看详尽的提交内容列表。

WebGPU 中的新变化

WebGPU 新变化系列中涵盖的所有内容的列表。

Chrome 131

Chrome 130

Chrome 129

Chrome 128

Chrome 127

Chrome 126

Chrome 125

Chrome 124

Chrome 123

Chrome 122

Chrome 121

Chrome 120

Chrome 119

Chrome 118

Chrome 117

Chrome 116

Chrome 115

Chrome 114

Chrome 113