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 个元素的内存位置。

建议在 WGSL 着色器代码的顶部使用 requires-directive 来表明 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() 设置的未捕获错误回调会在发生错误时立即调用。这是开发者在调试时一直期望和需要的。请参阅更改 dawn:173620。

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

以下仅介绍了一些关键亮点。查看详尽的提交列表。

WebGPU 最新动态

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

Chrome 153-154

Chrome 151-152

Chrome 149-150

Chrome 147-148

Chrome 146

Chrome 145

Chrome 144

Chrome 143

Chrome 142

Chrome 141

Chrome 140

Chrome 139

Chrome 138

Chrome 137

Chrome 136

Chrome 135

Chrome 134

Chrome 133

Chrome 132

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