当前位置:   article > 正文

一、WebGPU 基础理论开始篇

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 上运行的函数。这两个着色器是

  1. 顶点着色器

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

  2. 片段着色器

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

让我们从一个非常小的 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
推荐阅读
相关标签
  

闽ICP备14008679号