Gausst松鼠会

这个屌丝很懒，什么也没留下！

热门标签

一、WebGPU 基础理论开始篇

作者：Gausst松鼠会 | 2024-06-16 16:10:47

踩

webgpu

一、WebGPU 基础理论开始篇

开始

在某种程度上，WebGPU是一个非常简单的系统。它所做的只是在 GPU 上运行 3 种类型的函数：顶点着色器，片段着色器、计算着色器。

顶点着色器计算顶点。着色器返回顶点位置。对于每组 3 个顶点，它返回在这 3 个位置之间绘制的三角形。

片段着色器计算颜色。绘制三角形时，对于每个像素要绘制，GPU 会调用您的片段着色器。然后，片段着色器返回颜色。

计算着色器更通用。它实际上只是您调用的函数和说“执行此函数 N 次”。GPU 每次传递迭代次数不断调用你的函数，这样你就可以使用这个数字在每次迭代时来做一些独特的事情。

这类似于函数的功能传递给 array.forEach 或者 array.map。你在GPU上运行的函数只是函数，就像JavaScript一样。不同的部分是它们在GPU上运行，因此要运行它们，您可以将您希望他们访问的所有数据以以下形式复制到GPU 缓冲区和纹理，它们仅输出到这些缓冲区和纹理。您需要在函数中指定函数的绑定或位置以便查找对应数据。而且，回到JavaScript，你需要将绑定了缓冲区和纹理的数据保存到绑定或位置。完成此操作后，您可以告诉GPU执行该函数/方法。

也许图片会有所帮助。这是绘制三角形的WebGPU设置的简化图：通过使用顶点着色器和片段着色器。

在这里插入图片描述

关于此图的注意事项

有一个管道。它包含 GPU 将运行的顶点着色器和片段着色器。还可以使用带有计算着色器的管道。
着色器通过绑定组间接引用资源（缓冲区、纹理、采样器）
管道通过内部状态（internal state）间接引用缓冲区来定义属性
属性从缓冲区中提取数据，并将数据馈送到顶点着色器中。
顶点着色器可能会将数据馈送到片段着色器
片段着色器通过渲染通道间接写入纹理描述

要在 GPU 上执行着色器，您需要创建所有这些资源和设置此状态。资源的创建相对简单。一个有趣的事情是，大多数 WebGPU 资源在创建后无法更改。你可以更改其内容，但不能更改其大小，用途，格式等…如果需要帮助，要更改任何这些内容，您需要创建一个新资源并销毁旧资源。

某些状态是通过创建然后执行命令缓冲区来设置的。命令缓冲区就是它们的名字所暗示的。它们是缓冲区命令。你创建编码器。编码器将命令编码到命令缓冲区中。然后完成编码器，它为您提供命令缓冲区创建。然后，您可以提交该命令缓冲区，让 WebGPU 执行命令。

下面是编码命令缓冲区的一些伪代码，后面是所创建的命令缓冲区的表示。

在这里插入图片描述

创建命令缓冲区后，您可以提交以执行

device.queue.submit([commandBuffer]);
1

上图表示命令缓冲区中某个绘制命令的状态。执行命令将设置内部状态（internal state），然后绘制命令将告诉GPU执行顶点着色器(间接地执行片段着色器)。dispatchWorkgroup命令将告诉GPU执行计算着色器。

将三角形绘制到纹理

在WebGPU中，我们可以要求画布（canvas）提供纹理然后渲染到该纹理。

要使用 WebGPU 绘制三角形，我们必须提供 2 个“着色器”。再次，着色器是在 GPU 上运行的函数。这两个着色器是

顶点着色器

顶点着色器是计算绘制三角形/线/点的顶点位置的函数
片段着色器

片段着色器是在绘制三角形/线/点时计算要绘制/光栅化的每个像素的颜色(或其他数据)的函数

让我们从一个非常小的 WebGPU 程序开始绘制一个三角形。

我们需要一个画布来显示我们的三角形：

<canvas></canvas>
1

然后撰写 JS 代码。

WebGPU 是一个异步 API，因此在异步函数中最容易使用。我们首先请求适配器，然后从适配器请求设备。

async function main() {
  const adapter = await navigator.gpu?.requestAdapter();
  const device = await adapter?.requestDevice();
  if (!device) {
    fail('need a browser that supports WebGPU');
    return;
  }
}
main();
1
2
3
4
5
6
7
8
9

通过 navigator.gpu 获取 adapter。一个 adapter 代表一个特定的 GPU。有些设备可能有多个 GPU。

接下来，我们查找画布并为其创建上下文。这将让我们得到一个要渲染的纹理。该纹理将用于在网页。

  // Get a WebGPU context from the canvas and configure it
  const canvas = document.querySelector('canvas');
  const context = canvas.getContext('webgpu');
  const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
  context.configure({
    device,
    format: presentationFormat,
  });
1
2
3
4
5
6
7
8

首先从 canvas 画布中获取 webgpu 的上下文 context。

然后确定 canvas 的格式是什么，这里的格式可以是 “rgba8unorm” 或者 “bgra8unorm”。

接下来，我们创建一个着色器模块。着色器模块包含一个或多个着色器功能。在我们的例子中，我们将制作 1 个顶点着色器函数和 1 个片段着色器功能。

const module = device.createShaderModule({
    label: 'our hardcoded red triangle shaders',
    code: `
      @vertex fn vs(
        @builtin(vertex_index) vertexIndex : u32
      ) -> @builtin(position) vec4f {
        let pos = array(
          vec2f( 0.0,  0.5),  // top center
          vec2f(-0.5, -0.5),  // bottom left
          vec2f( 0.5, -0.5)   // bottom right
        );
 
        return vec4f(pos[vertexIndex], 0.0, 1.0);
      }
 
      @fragment fn fs() -> @location(0) vec4f {
        return vec4f(1.0, 0.0, 0.0, 1.0);
      }
    `,
  });
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

着色器是用 WGSL 语言编写的，WGSL 通常发音为 wig-sil。

可以看到上面我们定义了一个 vs 函数，它通过 @vertex 属性来声明。这意味着它是一个顶点着色器函数。

     @vertex fn vs(
        @builtin(vertex_index) vertexIndex : u32
      ) -> @builtin(position) vec4f {
         ...
1
2
3
4

它接受一个我们命名为vertexIndex的参数。vertexIndex是一个u32，意思是一个32位无符号整数。它从名为vertex_index的内置函数中获取其值。vertex_index类似于迭代号，类似于JavaScript数组中的索引。Map (function(value, index){…})。如果我们告诉GPU通过调用draw来执行这个函数10次，第一次vertex_index将是0，第二次它将是1，第三次它将是2，等等…

我们的vs函数被声明为返回一个vec4f，它是由4个32位浮点值组成的向量。可以把它想象成一个包含4个值的数组，或者一个具有4个属性的对象，比如{x: 0, y: 0, z: 0, w: 0}。这个返回值将被赋值给内置位置。在“triangle-list”模式下，每执行3次顶点着色器，就会绘制一个连接我们返回的3个位置值的三角形。

WebGPU中的位置需要在剪切空间中返回，其中X从左边的-1.0到右边的+1.0,Y从底部的-1.0到顶部的+1.0。无论我们绘制的纹理大小如何，都是如此。

在这里插入图片描述

vs函数声明了一个包含3个vec2fs的数组。每个vec2f由两个32位浮点值组成。

        let pos = array(
          vec2f( 0.0,  0.5),  // top center
          vec2f(-0.5, -0.5),  // bottom left
          vec2f( 0.5, -0.5)   // bottom right
        );
1
2
3
4
5

最后，它使用vertexIndex从数组的3个值中返回一个。由于函数需要4个浮点数作为返回类型，而pos是vec2f的数组，因此代码为剩下的2个值提供了0.0和1.0。

 return vec4f(pos[vertexIndex], 0.0, 1.0);
1

着色器模块还声明了一个名为fs的函数，该函数使用@fragment属性声明，使其成为一个片段着色器函数。

  @fragment fn fs() -> @location(0) vec4f {
1

此函数不接受任何参数，并返回location(0)处的vec4f。这意味着它将写入第一个渲染目标。我们稍后将使第一个渲染目标成为我们的画布canvas纹理。

  return vec4f(1, 0, 0, 1);
1

代码返回1,0,0,1，表示是红色的。在WebGPU中，颜色通常被指定为从0.0到1.0的浮点值，其中以上4个值分别对应红色、绿色、蓝色和alpha。

当GPU栅格化三角形(使用像素绘制它)时，它将调用片段着色器来找出每个像素的颜色。在我们的例子中，我们只是返回红色。

还有一点需要注意的是标签。几乎你用WebGPU创建的每个对象都可以带一个标签。标签是完全可选的，但最好的做法是给你做的所有东西都贴上标签。原因是，当你得到一个错误时，大多数WebGPU实现将打印一个错误消息，其中包括与错误相关的东西的标签。

在一个普通的应用程序中，你会有100或1000个缓冲区、纹理、着色器模块、管道等……如果你遇到类似"WGSL syntax error in shaderModule at line 10"这样的错误，如果你有100个着色器模块，哪一个出错了?如果你标记了这个模块，那么你会看到类似于"WGSL syntax error in shaderModule(‘our hardcoded red triangle shaders’) at line 10"这样的错误信息，这是一个更有用的错误信息，将为你节省大量的时间来跟踪这个问题。

现在我们已经创建了一个着色器模块，接下来我们需要制作一个渲染管道

  const pipeline = device.createRenderPipeline({
    label: 'our hardcoded red triangle pipeline',
    layout: 'auto',
    vertex: {
      module,
      entryPoint: 'vs',
    },
    fragment: {
      module,
      entryPoint: 'fs',
      targets: [{ format: presentationFormat }],
    },
  });
1
2
3
4
5
6
7
8
9
10
11
12
13

在这种情况下，没有什么可看的。我们将layout设置为auto，这意味着要求WebGPU从着色器中获取数据的布局。但我们没有使用任何数据。

然后我们告诉渲染管道使用来自着色器模块的vs函数作为顶点着色器，使用fs函数作为片段着色器。否则，我们告诉它第一个渲染目标的格式。“render target”表示我们将要渲染的纹理。我们创建了一个管道，我们必须指定纹理的格式，我们将使用这个管道最终渲染。

target数组的元素0对应于我们为片段着色器返回值指定的位置0。稍后，我们将目标设置为画布的纹理。接下来我们准备一个GPURenderPassDescriptor，它描述了我们想要绘制哪些纹理以及如何使用它们。

  const renderPassDescriptor = {
    label: 'our basic canvas renderPass',
    colorAttachments: [
      {
        // view: <- to be filled out when we render
        clearValue: [0.3, 0.3, 0.3, 1],
        loadOp: 'clear',
        storeOp: 'store',
      },
    ],
  }; 
1
2
3
4
5
6
7
8
9
10
11

一个GPURenderPassDescriptor有一个用于颜色附件colorAttachments的数组，它列出了我们将渲染的纹理以及如何处理它们。我们将等待填充我们真正想要渲染的纹理。现在，我们设置了一个明确的值semi-dark gray，并设置了loadOp和storeOp。loadOp: clear指定在绘制之前将纹理清除到clear值。另一个选项是load，这意味着将纹理的现有内容加载到GPU中，这样我们就可以在已经存在的内容上进行绘制。storeOp: store的意思是存储我们绘制的结果。我们也可以传入discard，它会丢弃我们绘制的内容。我们将在另一篇文章中介绍为什么要这样做。

现在是时候渲染了。

  function render() {
    // Get the current texture from the canvas context and
    // set it as the texture to render to.
    renderPassDescriptor.colorAttachments[0].view =
        context.getCurrentTexture().createView();
 
    // make a command encoder to start encoding commands
    const encoder = device.createCommandEncoder({ label: 'our encoder' });
 
    // make a render pass encoder to encode render specific commands
    const pass = encoder.beginRenderPass(renderPassDescriptor);
    pass.setPipeline(pipeline);
    pass.draw(3);  // call our vertex shader 3 times
    pass.end();
 
    const commandBuffer = encoder.finish();
    device.queue.submit([commandBuffer]);
  }
 
  render();
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

首先，我们调用context. getcurrentexture()来获取一个将出现在画布上的纹理。

调用createView得到一个视图到纹理的特定部分，但没有参数，它将返回默认部分，这就是我们在这个例子中想要的。

目前，我们唯一的颜色附件是画布上的纹理视图，它是通过我们一开始创建的上下文获得的。同样，数组colorAttachments的元素0对应@location(0)，这是我们为片段着色器的返回值指定的。

接下来我们创建一个命令编码器。命令编码器用于创建命令缓冲区。我们使用它对命令进行编码，然后“提交”它创建的命令缓冲区以执行命令。

然后，我们使用命令encoder通过调用beginRenderPass来创建渲染通道编码器。渲染通道编码器是用于创建与渲染相关的命令的特定编码器。

我们传入renderPassDescriptor来告诉它我们想渲染哪个纹理。我们对命令setPipeline进行编码，设置管道，然后通过调用draw with 3来告诉它执行我们的顶点着色器3次。

默认情况下，我们的顶点着色器每执行3次，就会通过连接刚刚从顶点着色器返回的3个值来绘制一个三角形。

我们结束渲染过程，然后完成编码器。这为我们提供了一个表示我们刚刚指定的步骤的命令缓冲区。最后，我们提交要执行的命令缓冲区。

当draw命令被执行时，这将是我们的状态：

在这里插入图片描述

上述例子中没有纹理，没有缓冲区，没有绑定组，但我们有一个管道，一个顶点和片段着色器，以及一个渲染传递描述符，它告诉我们的着色器渲染画布纹理。

以下就是我们的代码及其运行结果，这里的代码进行了些微调，并且使用了 TypeScript 来进行编写的：

HTML：

<!--
 * @Description: 
 * @Author: tianyw
 * @Date: 2022-11-11 12:50:23
 * @LastEditTime: 2023-04-09 15:58:11
 * @LastEditors: tianyw
-->
<!DOCTYPE html>
<html lang="en">

<head>
    <meta charset="UTF-8" />
    <meta name="viewport" content="width=device-width, initial-scale=1.0" />
    <title>001hello-triangle</title>
    <style>
        html,
        body {
            margin: 0;
            width: 100%;
            height: 100%;
            background: #000;
            color: #fff;
            display: flex;
            text-align: center;
            flex-direction: column;
            justify-content: center;
        }

        div,
        canvas {
            height: 100%;
            width: 100%;
        }
    </style>
</head>

<body>
    <div id="001hello-triangle">
        <canvas id="gpucanvas"></canvas>
    </div>
    <script type="module" src="./001hello-triangle.ts"></script>

</body>

</html>
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45

TS:

/*
 * @Description:
 * @Author: tianyw
 * @Date: 2023-04-08 20:03:35
 * @LastEditTime: 2023-09-16 01:07:44
 * @LastEditors: tianyw
 */
export type SampleInit = (params: {
  canvas: HTMLCanvasElement;
}) => void | Promise<void>;

import triangleVertWGSL from "./shaders/triangle.vert.wgsl?raw";
import redFragWGSL from "./shaders/red.frag.wgsl?raw";
const init: SampleInit = async ({ canvas }) => {
  const adapter = await navigator.gpu?.requestAdapter();
  if (!adapter) return;
  const device = await adapter?.requestDevice();
  if(!device) {
    console.error("need a browser that supports WebGPU");
    return;
  }
  const context = canvas.getContext("webgpu");
  if (!context) return;
  const devicePixelRatio = window.devicePixelRatio || 1;
  canvas.width = canvas.clientWidth * devicePixelRatio;
  canvas.height = canvas.clientHeight * devicePixelRatio;
  const presentationFormat = navigator.gpu.getPreferredCanvasFormat();

  context.configure({
    device,
    format: presentationFormat,
    alphaMode: "premultiplied"
  });

  const pipeline = device.createRenderPipeline({
    layout: "auto",
    vertex: {
      module: device.createShaderModule({
        code: triangleVertWGSL
      }),
      entryPoint: "main"
    },
    fragment: {
      module: device.createShaderModule({
        code: redFragWGSL
      }),
      entryPoint: "main",
      targets: [
        {
          format: presentationFormat
        }
      ]
    },
    primitive: {
      // topology: "line-list"
      // topology: "line-strip"
      //  topology: "point-list"
      topology: "triangle-list"
      // topology: "triangle-strip"
    }
  });

  function frame() {
    const commandEncoder = device.createCommandEncoder();
    if (!context) return;
    const textureView = context.getCurrentTexture().createView();
    const renderPassDescriptor: GPURenderPassDescriptor = {
      colorAttachments: [
        {
          view: textureView,
          clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 },
          loadOp: "clear",
          storeOp: "store"
        }
      ]
    };

    const passEncoder = commandEncoder.beginRenderPass(renderPassDescriptor);
    passEncoder.setPipeline(pipeline);
    passEncoder.draw(3, 1, 0, 0);
    passEncoder.end();

    device.queue.submit([commandEncoder.finish()]);
    requestAnimationFrame(frame);
  }

  requestAnimationFrame(frame);
};

const canvas = document.getElementById("gpucanvas") as HTMLCanvasElement;
init({ canvas: canvas });

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92

Shader：

片元着色器：

@fragment
fn main() -> @location(0) vec4<f32> {
  return vec4(1.0, 0.0, 0.0, 1.0);
}

1
2
3
4
5

顶点着色器：

@vertex
fn main(
  @builtin(vertex_index) VertexIndex : u32
) -> @builtin(position) vec4<f32> {
  var pos = array<vec2<f32>, 3>(
    vec2(0.0, 0.5),
    vec2(-0.5, -0.5),
    vec2(0.5, -0.5)
  );

  return vec4<f32>(pos[VertexIndex], 0.0, 1.0);
}

1
2
3
4
5
6
7
8
9
10
11
12
13

目录结构：

在这里插入图片描述

运行结果：

在这里插入图片描述

需要强调的是，我们调用的所有这些函数，如setPipeline和draw，只是将命令添加到命令缓冲区。它们实际上并不执行这些命令。当我们将命令缓冲区提交到设备队列时（device queue），命令就会被执行。

WebGPU使用我们从顶点着色器返回的每3个顶点来栅格化三角形。它通过确定哪些像素的中心在三角形内来实现这一点。然后，它为每个像素调用我们的片段着色器，以询问将其设置为什么颜色。

想象一下我们渲染的纹理是15x11像素。这些是要绘制的像素

在这里插入图片描述

所以，现在我们已经看到了一个非常小的WebGPU示例。很明显，在着色器中硬编码一个三角形是不灵活的。我们需要提供数据的方法，我们将在以下文章中介绍这些方法。从上面的代码中可以看出:

WebGPU只是运行着色器。它由你用代码填充它们来做有用的事情
着色器在着色器模块中指定，然后转换为管道
WebGPU可以绘制三角形
WebGPU绘制纹理(我们碰巧从画布中得到一个纹理)
WebGPU通过编码命令工作，然后提交它们。

在GPU上运行计算

让我们写一个在GPU上进行一些计算的基本示例。

我们从获得WebGPU设备的相同代码开始：

async function main() {
  const adapter = await gpu?.requestAdapter();
  const device = await adapter?.requestDevice();
  if (!device) {
    fail('need a browser that supports WebGPU');
    return;
  }
1
2
3
4
5
6
7

创建着色器：

  const module = device.createShaderModule({
    label: 'doubling compute module',
    code: `
      @group(0) @binding(0) var<storage, read_write> data: array<f32>;
 
      @compute @workgroup_size(1) fn computeSomething(
        @builtin(global_invocation_id) id: vec3<u32>
      ) {
        let i = id.x;
        data[i] = data[i] * 2.0;
      }
    `,
  });
1
2
3
4
5
6
7
8
9
10
11
12
13

首先声明一个类型为storage的变量data，我们希望它既能读又能写。

 @group(0) @binding(0) var<storage, read_write> data: array<f32>;
1

我们将其类型声明为array<f32>这意味着一个32位浮点值的数组。我们告诉它，我们将在bindGroup 0 (@group(0))中的绑定位置0(binding(0))上指定这个数组。然后我们声明一个名为computeSomething的函数，它带有@compute属性，使其成为一个compute着色器。

    @compute @workgroup_size(1) fn computeSomething(
        @builtin(global_invocation_id) id: vec3u
      ) {
        ...
1
2
3
4

计算着色器需要声明工作组的大小，我们将在后面介绍。现在，我们只需通过属性@workgroup_size(1)将其设置为1。我们声明它有一个使用vec3u的参数id。vec3u是三个无符号32整数值。就像上面的顶点着色器一样，这是迭代次数。不同的是，计算着色器迭代次数是3维的(有3个值)。我们声明id来从内置的global_invocation_id中获取值。

你可以把计算着色器想象成这样运行。这是一个过于简化的例子，但现在就可以了。

// pseudo code
function dispatchWorkgroups(width, height, depth) {
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const workgroup_id = {x, y, z};
        dispatchWorkgroup(workgroup_id)
      }
    }
  }
}
 
function dispatchWorkgroup(workgroup_id) {
  // from @workgroup_size in WGSL
  const workgroup_size = shaderCode.workgroup_size;
  const {x: width, y: height, z: depth} = workgroup.size;
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const local_invocation_id = {x, y, z};
        const global_invocation_id =
            workgroup_id * workgroup_size + local_invocation_id;
        computeShader(global_invocation_id)
      }
    }
  }
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27

由于我们设置了@workgroup_size(1)，因此上面的伪代码实际上变成了

// pseudo code
function dispatchWorkgroups(width, height, depth) {
  for (z = 0; z < depth; ++z) {
    for (y = 0; y < height; ++y) {
      for (x = 0; x < width; ++x) {
        const workgroup_id = {x, y, z};
        dispatchWorkgroup(workgroup_id)
      }
    }
  }
}
 
function dispatchWorkgroup(workgroup_id) {
  const global_invocation_id = workgroup_id;
  computeShader(global_invocation_id)
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16

最后，我们使用id的x属性来索引数据，并将每个值乘以2

   let i = id.x;
        data[i] = data[i] * 2.0;
1
2

上面，i只是3个迭代数中的第一个。现在我们已经创建了着色器，我们需要创建一个管道

const pipeline = device.createComputePipeline({
    label: 'doubling compute pipeline',
    layout: 'auto',
    compute: {
      module,
      entryPoint: 'computeSomething',
    },
  });
1
2
3
4
5
6
7
8

这里我们只是告诉它，我们正在使用创建的着色器模块中的计算阶段，并希望调用computeSomething函数。布局再次是auto，告诉WebGPU从着色器中找出布局。接下来我们需要一些数据

  const input = new Float32Array([1, 3, 5]);
1

这些数据只存在于JavaScript中。为了让WebGPU使用它，我们需要在GPU上创建一个缓冲区，并将数据复制到缓冲区中。

 // create a buffer on the GPU to hold our computation
  // input and output
  const workBuffer = device.createBuffer({
    label: 'work buffer',
    size: input.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
  });
  // Copy our input data to that buffer
  device.queue.writeBuffer(workBuffer, 0, input);
1
2
3
4
5
6
7
8
9

通过调用 device.createBuffer 来创建一个缓冲区。size是以字节为单位的大小，在本例中是12，因为3个值组成的Float32Array的size以字节为单位是12。

我们创建的每个WebGPU缓冲区都必须指定用法 usage。我们可以传递一堆标志 flags 以供使用，但不是所有标志都可以一起使用。在这里，我们通过传递GPUBufferUsage.STORAGE参数，让这个缓冲区可用作存储 storage 。这使得它与来自着色器的var<storage,…>相匹配/兼容。此外，我们希望能够将数据复制到该缓冲区，因此我们需要包含GPUBufferUsage.COPY_DST flag。最后，我们希望能够从缓冲区中复制数据，因此我们引入了GPUBufferUsage.COPY_SRC。

注意，你不能直接从JavaScript中读取WebGPU缓冲区的内容。相反，你必须“映射/map”它，这是从WebGPU请求访问缓冲区的另一种方式，因为缓冲区可能正在使用，并且它可能只存在于GPU上。

可以在JavaScript中映射的WebGPU缓冲区不能用于其他太多地方。换句话说，我们不能映射我们刚刚创建的缓冲区，如果我们尝试添加标志/flag使其可映射，我们将得到一个错误，它与usage STORAGE不兼容。

因此，为了查看计算结果，我们需要另一个缓冲区。运行计算后，我们将上面的缓冲区复制到这个结果缓冲区，并设置其标志，以便我们可以映射它。

// create a buffer on the GPU to get a copy of the results
  const resultBuffer = device.createBuffer({
    label: 'result buffer',
    size: input.byteLength,
    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
  });
1
2
3
4
5
6

MAP_READ意味着我们希望能够映射这个缓冲区以读取数据。为了告诉着色器我们想要处理的缓冲区，我们需要创建一个bindGroup

// Setup a bindGroup to tell the shader which
  // buffer to use for the computation
  const bindGroup = device.createBindGroup({
    label: 'bindGroup for work buffer',
    layout: pipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: { buffer: workBuffer } },
    ],
  });
1
2
3
4
5
6
7
8
9

我们从管道中获取bindGroup的布局。然后我们设置bindGroup入口entries。pipeline.getBindGroupLayout(0)中的0对应着色器中的@group(0)。{binding: 0…的值对应着色器中的@group(0) @binding(0)。

现在我们可以开始编码命令了

// Encode commands to do the computation
  const encoder = device.createCommandEncoder({
    label: 'doubling encoder',
  });
  const pass = encoder.beginComputePass({
    label: 'doubling compute pass',
  });
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(input.length);
  pass.end();
1
2
3
4
5
6
7
8
9
10
11

我们创建一个命令编码器。我们启动一个compute pass。我们设置管道，然后设置bindGroup。这里是传递的0 在 setBindGroup(0, bindGroup)对应着色器中的@group(0)。然后，我们调用dispatchWorkgroups，在这种情况下，我们向它传递输入 input.length值为3，告诉WebGPU运行3次compute shader。然后我们结束 pass。

以下是dispatchWorkgroups执行时的情况:

在这里插入图片描述

计算完成后，我们请求WebGPU从workBuffer复制到resultBuffer

 // Encode a command to copy the results to a mappable buffer.
  encoder.copyBufferToBuffer(workBuffer, 0, resultBuffer, 0, resultBuffer.size);
1
2

现在我们可以 finish 编码器以获取命令缓冲区，然后提交该命令缓冲区。

  // Finish encoding and submit the commands
  const commandBuffer = encoder.finish();
  device.queue.submit([commandBuffer]);
1
2
3

然后，我们映射结果缓冲区并获得数据的一份副本

 // Read the results
  await resultBuffer.mapAsync(GPUMapMode.READ);
  const result = new Float32Array(resultBuffer.getMappedRange());
 
  console.log('input', input);
  console.log('result', result);
 
  resultBuffer.unmap();
1
2
3
4
5
6
7
8

为了映射结果缓冲区，我们调用mapAsync，并且必须等待它完成。一旦映射，我们可以调用resultBuffer.getMappedRange()，它不带参数将返回整个缓冲区的ArrayBuffer。我们把它放在Float32Array类型的数组视图中，然后我们可以查看这些值。一个重要的细节是，getMappedRange返回的ArrayBuffer只在调用unmap之前有效。在将其长度设置为0后解除映射，其数据不再可访问。
运行一下，我们可以看到我们得到了结果，所有的数字都翻了一倍。

在这里插入图片描述

我们将在其他文章中介绍如何真正使用计算着色器。现在，你希望已经对WebGPU的功能有了一些了解。其他一切由你决定!可以把WebGPU想象成类似于其他编程语言。它提供了一些基本功能，剩下的留给你自己创造。WebGPU编程的特殊之处在于这些函数，顶点着色器、片段着色器和计算着色器，都在GPU上运行。一个GPU可能有超过10000个处理器，这意味着它们可能并行执行超过10000个计算，这可能比你的CPU并行执行的能力高出3个或更多的数量级。

以下为上面的编码及其运行结果：

HTML：

<!--
 * @Description: 
 * @Author: tianyw
 * @Date: 2022-11-11 12:50:23
 * @LastEditTime: 2023-09-17 16:33:32
 * @LastEditors: tianyw
-->
<!DOCTYPE html>
<html lang="en">

<head>
    <meta charset="UTF-8" />
    <meta name="viewport" content="width=device-width, initial-scale=1.0" />
    <title>001hello-triangle</title>
    <style>
        html,
        body {
            margin: 0;
            width: 100%;
            height: 100%;
            background: #000;
            color: #fff;
            display: flex;
            text-align: center;
            flex-direction: column;
            justify-content: center;
        }

        div,
        canvas {
            height: 100%;
            width: 100%;
        }
    </style>
</head>

<body>
    <div id="002hello-compute">
        <canvas id="gpucanvas"></canvas>
    </div>
    <script type="module" src="./002hello-compute.ts"></script>

</body>

</html>
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45

TS:

/*
 * @Description:
 * @Author: tianyw
 * @Date: 2023-04-08 20:03:35
 * @LastEditTime: 2023-09-17 16:41:00
 * @LastEditors: tianyw
 */
export type SampleInit = (params: {
  canvas: HTMLCanvasElement;
}) => void | Promise<void>;

import triangleVertWGSL from "./shaders/triangle.vert.wgsl?raw";
import redFragWGSL from "./shaders/red.frag.wgsl?raw";
import dataComputeWGSL from "./shaders/data.compute.wgsl?raw";
const init: SampleInit = async ({ canvas }) => {
  const adapter = await navigator.gpu?.requestAdapter();
  if (!adapter) return;
  const device = await adapter?.requestDevice();
  if (!device) {
    console.error("need a browser that supports WebGPU");
    return;
  }
  const context = canvas.getContext("webgpu");
  if (!context) return;
  const devicePixelRatio = window.devicePixelRatio || 1;
  canvas.width = canvas.clientWidth * devicePixelRatio;
  canvas.height = canvas.clientHeight * devicePixelRatio;
  const presentationFormat = navigator.gpu.getPreferredCanvasFormat();

  context.configure({
    device,
    format: presentationFormat,
    alphaMode: "premultiplied"
  });

  const renderPipeline = device.createRenderPipeline({
    layout: "auto",
    vertex: {
      module: device.createShaderModule({
        code: triangleVertWGSL
      }),
      entryPoint: "main"
    },
    fragment: {
      module: device.createShaderModule({
        code: redFragWGSL
      }),
      entryPoint: "main",
      targets: [
        {
          format: presentationFormat
        }
      ]
    },
    primitive: {
      // topology: "line-list"
      // topology: "line-strip"
      //  topology: "point-list"
      topology: "triangle-list"
      // topology: "triangle-strip"
    }
  });
  async function createCompute() {
    const computePipeLine = device.createComputePipeline({
      label: "doubling compute module",
      layout: "auto",
      compute: {
        module: device.createShaderModule({
          label: "doubling compute module",
          code: dataComputeWGSL
        }),
        entryPoint: "computeSomething"
      }
    });
    const input = new Float32Array([1, 3, 5]);
    const commandEncoder = device.createCommandEncoder({
      label: "doubling encoder"
    });
    const computePass = commandEncoder.beginComputePass({
      label: "doubling compute pass"
    });

    const workBuffer = device.createBuffer({
      label: "work buffer",
      size: input.byteLength,
      usage:
        GPUBufferUsage.STORAGE |
        GPUBufferUsage.COPY_SRC |
        GPUBufferUsage.COPY_DST
    });
    device.queue.writeBuffer(workBuffer, 0, input);

    const resultBuffer = device.createBuffer({
      label: "result buffer",
      size: input.byteLength,
      usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST
    });

    const bindGroup = device.createBindGroup({
      label: "bindGroup for work buffer",
      layout: computePipeLine.getBindGroupLayout(0),
      entries: [{ binding: 0, resource: { buffer: workBuffer } }]
    });

    computePass.setPipeline(computePipeLine);
    computePass.setBindGroup(0, bindGroup);
    computePass.dispatchWorkgroups(input.length);
    computePass.end();

    commandEncoder.copyBufferToBuffer(
      workBuffer,
      0,
      resultBuffer,
      0,
      resultBuffer.size
    );

    const commandBuffer = commandEncoder.finish();
    device.queue.submit([commandBuffer]);

    await resultBuffer.mapAsync(GPUMapMode.READ);
    const result = new Float32Array(resultBuffer.getMappedRange().slice(0));
    resultBuffer.unmap();

    console.log("input", input);
    console.log("result", result);
  }

  function frame() {
    const renderCommandEncoder = device.createCommandEncoder({
      label: "render vert frag"
    });
    if (!context) return;

    const textureView = context.getCurrentTexture().createView();
    const renderPassDescriptor: GPURenderPassDescriptor = {
      colorAttachments: [
        {
          view: textureView,
          clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 },
          loadOp: "clear",
          storeOp: "store"
        }
      ]
    };
    const renderPass =
      renderCommandEncoder.beginRenderPass(renderPassDescriptor);
    renderPass.setPipeline(renderPipeline);
    renderPass.draw(3, 1, 0, 0);
    renderPass.end();
    const renderBuffer = renderCommandEncoder.finish();
    device.queue.submit([renderBuffer]);

    requestAnimationFrame(frame);
  }
  createCompute();
  requestAnimationFrame(frame);
};

const canvas = document.getElementById("gpucanvas") as HTMLCanvasElement;
init({ canvas: canvas });

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162

Shaders:

vert:

@vertex
fn main(
  @builtin(vertex_index) VertexIndex : u32
) -> @builtin(position) vec4<f32> {
  var pos = array<vec2<f32>, 3>(
    vec2(0.0, 0.5),
    vec2(-0.5, -0.5),
    vec2(0.5, -0.5)
  );

  return vec4<f32>(pos[VertexIndex], 0.0, 1.0);
}

1
2
3
4
5
6
7
8
9
10
11
12
13

frag:

@fragment
fn main() -> @location(0) vec4<f32> {
  return vec4(1.0, 0.0, 0.0, 1.0);
}

1
2
3
4
5

compute:

@group(0) @binding(0) var<storage, read_write> data: array<f32>;

@compute @workgroup_size(1)fn computeSomething(@builtin(global_invocation_id) id: vec3<u32>) {
    let i = id.x;
    data[i] = data[i] * 2.0;
}
1
2
3
4
5
6

在这里插入图片描述

运行结果：

在这里插入图片描述

声明：本文内容由网友自发贡献，不代表【wpsshop博客】立场，版权归原作者所有，本站不承担相应法律责任。如您发现有侵权的内容，请联系我们。转载请注明出处：https://www.wpsshop.cn/w/Gausst松鼠会/article/detail/727332