Vala プログラミング

WebGPU プログラミング

おなが@京都先端科学大

WebGPU Triangle and Instancing (WebGPU, WGSL update)

06/28/2022
chrome canary version 105
WebGPU, WGSL 06/17/2022

chrome canary について
(WebGPU を利用するには、chrome://flags/#enable-unsafe-webgpu
を設定する。)
 この時点では GPUDevice の読み込みが不安定になっています。
 requestDevice の read error となる場合は、再読み込みをします。
 これを何回か繰り返すと、エラーは無いが表示が出来なくなる場合があります。
 この場合は、chrome を再起動します。

[実行結果]
Triangle

Instancing

変更箇所
Triangle
webGPU
error
1 canvas.getContext("gpupresent") -> canvas.getContext("webgpu")
2 context.configureSwapChain() -> context.configure()
3 context.getPreferredFormat(adapter) ->
  navigator.gpu.getPreferredCanvasFormat()
4 swapChain.getCurrentTexture().createView() ->
  context.getCurrentTexture().createView()
5 render pass: endPass() -> end()

warning
1 CanvasConfiguration
 alphaMode default change ->
  alphaMode を使用するときは、”premultiplied" を設定
2 render pipeline lauout
 layout の設定が必要 -> implicit pipeline layout の場合は、’auto' を設定
3 renderPass ColorAttachments
 attachment -> view
 loadValue -> clearValue
 loadOp -> 要設定
 storeOp -> 要設定

WGSL
attributes の変更

1 [[stage(vertex)]] -> @vertex
2 [[builtin(vertex_index)]] -> @builtin(vertex_index)
   [[builtin(position)]] -> @builtin(position)
3 [[stage(fragment)]] -> @fragment
4 [[location(0)]] -> @location(0)

Instancing
上記以外の変更箇所
WebGPU
1 compute pass dispatch() -> dispatchWorkgroups()
2 compute pass endPass() -> end()

WGSL
1 structure

 [[block]] struct Uniforms {} -> [[block]] 不要
 struct Uniforms {
    proj_view : mat4x4<f32>
 }
 メンバー区切りの変更 ";" -> ","

2 attributes

 [[binding(0), group(0)]] -> @binding(0) @group(0)
 [[stage(compute)]] -> @compute @workgroup_size(64)

プログラム
triangle.html

<!DOCTYPE html>
<html>
<head>
    <meta charset="utf-8">
    <title>WebGPU Triangle WGSL</title>
</head>
<body>
<canvas id="webgpu-canvas" width="400" height="400"></canvas>
<script>

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

    const canvas = document.getElementById("webgpu-canvas");
    const context = canvas.getContext("webgpu");

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

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

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

    const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
    context.configure({
        device,
        format: presentationFormat,
        alphaMode: "premultiplied",
    });

    const pipeline = device.createRenderPipeline({
        layout: "auto",
        vertex: {
            module: device.createShaderModule({
                code: wgslShaders.vertex
            }),
            entryPoint: "main",
        },
        fragment: {
            module: device.createShaderModule({
                code: wgslShaders.fragment
            }),
            entryPoint: "main",
            targets: [{
                format: "bgra8unorm",
            }],
        },
        primitive: {
            topology: "triangle-list"
        },
    });

    // render
    requestAnimationFrame(function frame() {
        const commandEncoder = device.createCommandEncoder();
        
        const textureView = context.getCurrentTexture().createView();
        const renderPassDesc = {
            colorAttachments: [{
                view: textureView,
                clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 1.0 },
                loadOp: 'clear',
                storeOp: 'store',
            }]
        };

        // render pass
        const renderPass = commandEncoder.beginRenderPass(renderPassDesc);
        renderPass.setPipeline(pipeline);
        renderPass.draw(3, 1, 0, 0);
        renderPass.end();

        device.queue.submit([commandEncoder.finish()]);

        requestAnimationFrame(frame);
    });
})();
</script>
</body>
</html>

instancing.html

<!DOCTYPE html>
<html>
<head>
    <meta charset="utf-8">
    <title>WebGPU Instancing</title>
    <script src="../libs/gl-matrix-min.js"></script>
    <script src="../libs/webgl-util.js"></script>
    <script src="../libs/utils.js"></script>
</head>
<body>
<canvas id="webgpu-canvas" width="500" height="400"></canvas>
<script>

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

    const canvas = document.getElementById("webgpu-canvas");
    const context = canvas.getContext("webgpu");

    const NUM_PARTICLES = 100;
    const POINT_SIZE = 0.1;

    //// camera
    const eye = [0, 0, 4];
    const center = [0, 0, 0];
    const up = [0, 1, 0];

    //// initial particle data
    var particleData = new Float32Array(8 * NUM_PARTICLES);
    for (let i = 0; i < particleData.length; i += 8)
    {
        //// position data
        particleData[i]     = Math.random() * 2.0 - 1.0;
        particleData[i + 1] = Math.random() * 2.0 - 1.0;
        particleData[i + 2] = Math.random() * 2.0 - 1.0;
        particleData[i + 3] = 1;
        //// velocity data
        particleData[i + 4] = (Math.random() * 2.0 - 1.0) * 0.005;
        particleData[i + 5] = (Math.random() * 2.0 - 1.0) * 0.005;
        particleData[i + 6] = (Math.random() * 2.0 - 1.0) * 0.005;
        particleData[i + 7] = 1;
    }

    const particleBuffer = device.createBuffer({
        size: particleData.byteLength,
        usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE
    });
    device.queue.writeBuffer(particleBuffer, 0, particleData);

    //// WGSL Shader (vertex, fragnebt, compute)
    const wgslShaders = {
    vertex:`
    struct Uniforms {
        proj_view : mat4x4<f32>
    };

    @binding(0) @group(0) var<uniform> uniforms : Uniforms;

    @vertex
    fn main(@location(0) vertexPosition : vec3<f32>,
            @location(2) position : vec3<f32>)
        -> @builtin(position) vec4<f32>{
        var scale : f32 = ${POINT_SIZE};
        var scaleMTX : mat4x4<f32> = mat4x4<f32>(
            vec4<f32>(scale, 0.0,   0.0,   0.0),
            vec4<f32>(0.0,   scale, 0.0,   0.0),
            vec4<f32>(0.0,   0.0,   scale, 0.0),
            vec4<f32>(position, 1.0)
        );

        return uniforms.proj_view * scaleMTX * vec4<f32>(vertexPosition, 1.0);
    }
    `,

    fragment:`
    @fragment
    fn main() -> @location(0) vec4<f32> {
        return vec4<f32>(1.0, 0.0, 0.0, 1.0);
    }
    `,
    
    compute:`
    struct Particle {
        pos : vec3<f32>,
        vel : vec3<f32>,
    };
    struct Particles {
        particles : array<Particle, ${NUM_PARTICLES}>
    };

    @binding(0) @group(0) var<storage, read_write> particle : Particles;

    @compute @workgroup_size(64)
    fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
        var index: u32 = GlobalInvocationID.x;

        var position: vec3<f32> = particle.particles[index].pos;
        var velocity: vec3<f32> = particle.particles[index].vel;
        var new_position: vec3<f32> = position + velocity;

        if (abs(new_position.x) >= 1.0) {
            particle.particles[index].vel.x = - velocity.x;
        }
        if (abs(new_position.y) >= 1.0) {
            particle.particles[index].vel.y = - velocity.y;
        }
        if (abs(new_position.z) >= 1.0) {
            particle.particles[index].vel.z = - velocity.z;
        }

        particle.particles[index].pos = new_position;
    }
    `
    };

    const computePipeline = device.createComputePipeline({
        layout: "auto",
        compute: {
            module: device.createShaderModule({
                code: wgslShaders.compute}),
            entryPoint: "main"
        }
    });

    const computeBindGroup = device.createBindGroup({
        layout: computePipeline.getBindGroupLayout(0),
        entries: [
            {
                binding: 0,
                resource: {
                    buffer: particleBuffer
                }
            },
        ]
    });

    const cubeData = utils.createCube();
    const numVertices = cubeData.positions.length / 3;

    const vertexBuffer = device.createBuffer({
        size: cubeData.positions.byteLength,
        usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST
    });
    device.queue.writeBuffer(vertexBuffer, 0, cubeData.positions);

    //// Setup render outputs
    const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
    context.configure({
        device,
        format: presentationFormat,
        alphaMode: "premultiplied",
    });

    var pipeline_point = device.createRenderPipeline({
        layout: "auto",
        vertex: {
            module: device.createShaderModule({
                code: wgslShaders.vertex}),
            entryPoint: "main",
            buffers: [
            {
                arrayStride: 12,
                attributes: [
                {
                    shaderLocation: 0,
                    format: "float32x3",
                    offset: 0
                }]
            },
            {
                arrayStride: 16 * 2,
                stepMode: "instance",
                attributes: [{
                    shaderLocation: 2,
                    format: "float32x3",
                    offset: 0
                }]
            }
            ],
        },
        fragment: {
            module: device.createShaderModule({
                code: wgslShaders.fragment}),
            entryPoint: "main",
            targets: [{
                format: "bgra8unorm",
            }],
        },
        primitive: {
            topology: "triangle-list"
        },
    });

    //// Setup renderPassDesc
    var renderPassDesc = {
        colorAttachments: [{
            view: undefined,
            clearValue: { r: 0.2, g: 0.2, b: 0.2, a: 1.0 },
                loadOp: 'clear',
                storeOp: 'store',
        }],
    };

    //// Create uniform buffer
    var viewParamsBuffer = device.createBuffer({
        size: 16 * 4,
        usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST
    });

    //// Create bind group (setup uniform buffer)
    var bindGroup = device.createBindGroup({
        layout: pipeline_point.getBindGroupLayout(0),
        entries: [
            {
                binding: 0,
                resource: {
                    buffer: viewParamsBuffer
                }
            }
        ]
    });

    //// Create arcball camera and view projection matrix
    var camera = new ArcballCamera(eye, center, up,
        0.5, [canvas.width, canvas.height]);
    var projection = mat4.perspective(mat4.create(), 50 * Math.PI / 180.0,
        canvas.width / canvas.height, 0.1, 100);
    //// projection_view matrix
    var projView = mat4.create();

    //// Controller
    var controller = new Controller();
    controller.mousemove = function(prev, cur, evt) {
        if (evt.buttons == 1) {
            camera.rotate(prev, cur);
        } else if (evt.buttons == 2) {
            camera.pan([cur[0] - prev[0], prev[1] - cur[1]]);
        }
    };
    controller.wheel = function(amt) { camera.zoom(amt * 3.0); };
    controller.registerForCanvas(canvas);

    //// render
    requestAnimationFrame(function frame() {
        const commandEncoder = device.createCommandEncoder();

        //// compute pass
        var computePass = commandEncoder.beginComputePass();

        computePass.setPipeline(computePipeline);
        computePass.setBindGroup(0, computeBindGroup);
        computePass.dispatchWorkgroups(NUM_PARTICLES);

        computePass.end();

        //// render pass setting
        renderPassDesc.colorAttachments[0].view =
            context.getCurrentTexture().createView();

        projView = mat4.mul(projView, projection, camera.camera);
        device.queue.writeBuffer(viewParamsBuffer, 0, projView);

        const renderPass = commandEncoder.beginRenderPass(renderPassDesc);
        renderPass.setPipeline(pipeline_point);
        renderPass.setVertexBuffer(0, vertexBuffer);
        renderPass.setVertexBuffer(1, particleBuffer);
        renderPass.setBindGroup(0, bindGroup);
        renderPass.draw(numVertices, NUM_PARTICLES, 0, 0);

        renderPass.end();

        device.queue.submit([commandEncoder.finish()]);

        requestAnimationFrame(frame);
    });
})();
</script>
</body>
</html>

プログラムには、以下のライブラリを使用しています。
1 point(cubeで表現)を描画するためのcubeプログラム utils.js
  WebGPU Examples https://github.com/tsherif/webgpu-examples
2 cameraプログラム webgl-util.min.js
  WebGPU Experiments https://github.com/Twinklebear/webgpu-experiments
3 matrix and vector計算ライブラリ gl-matrix.js
  https://github.com/toji/gl-matrix