LoginSignup
2
3
@tetsutakamurata76

はじめての　WebGPU。でシンプルなWebアプリケーション。

Last updated at Posted at 2024-09-08

ショートストーリータイトル: 「WebGPU。行列の星」

WebGPUの初期化


async function initWebGPU() {
    if (!navigator.gpu) {
        console.log("WebGPUはこのブラウザでサポートされていません。");
        return;
    }

    try {
        const adapter = await navigator.gpu.requestAdapter();
        if (!adapter) {
            console.log("WebGPUアダプターの取得に失敗しました。");
            return;
        }
        
        const device = await adapter.requestDevice();
        const outputElement = document.getElementById('output');

東京の静かな夜、ビルの間をすり抜ける夜風が心地よく感じられる。しかし、下町の一角にある小さなオフィスの中では、主人公のプログラマである浩司が熱心にモニターの前に向かっていた。彼のデスクには、複数のコードが並び、手元にはコーヒーのカップが置かれている。浩司の目は、画面に映し出された行列の計算プログラムに釘付けだ。

行列の初期化


        const matrixSize = 4;
        const matrixA = new Float32Array(matrixSize * matrixSize);
        const matrixB = new Float32Array(matrixSize * matrixSize);
        const matrixC = new Float32Array(matrixSize * matrixSize);

        for (let i = 0; i < matrixSize * matrixSize; i++) {
            matrixA[i] = Math.random();
            matrixB[i] = Math.random();
        }

浩司は、最近のプロジェクトである「行列の星」に取り組んでいた。このプロジェクトは、東京の高度な技術を駆使して行列の計算を効率化し、リアルタイムで結果を表示するものだった。彼のタスクは、WebGPUという新しい技術を使って、4x4の行列の掛け算を高速に行い、その結果を美しく可視化することだった。

WebGPU バッファの作成とデータ転送


        const gpuBufferA = device.createBuffer({
            size: matrixA.byteLength,
            usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
            mappedAtCreation: true
        });
        new Float32Array(gpuBufferA.getMappedRange()).set(matrixA);
        gpuBufferA.unmap();

        const gpuBufferB = device.createBuffer({
            size: matrixB.byteLength,
            usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
            mappedAtCreation: true
        });
        new Float32Array(gpuBufferB.getMappedRange()).set(matrixB);
        gpuBufferB.unmap();

        const gpuBufferC = device.createBuffer({
            size: matrixC.byteLength,
            usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST
        });

        const gpuBufferRead = device.createBuffer({
            size: matrixC.byteLength,
            usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
        });

「よし、これで行列AとBのデータがGPUに転送できたはずだ。」浩司は、プログラムの一部が正常に動作していることを確認し、ホッと息をつく。彼は行列AとBにランダムなデータを設定し、これらのデータをGPUのバッファに送り込む作業を終えたばかりだった。

シェーダコードの作成


        const shaderCode = `
            @group(0) @binding(0) var<storage, read> matrixA : array<f32>;
            @group(0) @binding(1) var<storage, read> matrixB : array<f32>;
            @group(0) @binding(2) var<storage, read_write> matrixC : array<f32>;

            @compute @workgroup_size(1)
            fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
                let row = global_id.x;
                let col = global_id.y;
                if (row < 4u && col < 4u) {
                    let size = 4u;
                    var sum = 0.0;
                    for (var i = 0u; i < size; i = i + 1u) {
                        sum = sum + matrixA[row * size + i] * matrixB[i * size + col];
                    }
                    matrixC[row * size + col] = sum;
                }
            }
        `;

「さて、次はシェーダコードを定義しなきゃ。」浩司はキーボードを叩き、WGSL（WebGPU Shading Language）で行列の掛け算を計算するシェーダコードを記述し始めた。彼の目には、計算が正確に行われることを確認する使命感が宿っていた。コードを書き終えた彼は、シェーダをコンパイルし、コンピュートパイプラインを構築する作業に取り掛かった。

コンピュートパイプラインとバインドグループの作成


        const shaderModule = device.createShaderModule({
            code: shaderCode
        });

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

        const bindGroup = device.createBindGroup({
            layout: computePipeline.getBindGroupLayout(0),
            entries: [
                { binding: 0, resource: { buffer: gpuBufferA }},
                { binding: 1, resource: { buffer: gpuBufferB }},
                { binding: 2, resource: { buffer: gpuBufferC }}
            ]
        });

その間も、東京の夜は静かに過ぎていく。ビルの明かりが夜空に映り、星のように輝いている。浩司は、これらの星のような行列が、彼のプログラムの中で織りなす美しい計算のパターンを想像していた。彼のプログラムが正しく動作し、行列の掛け算がスムーズに計算されることを願っていた。

「さあ、計算結果が出るはずだ。」浩司は、プログラムを実行し、GPUが計算を終えるのを待った。処理が完了し、結果が画面に表示されると、彼はその美しい数値の並びに感動した。それはまるで、星空に輝く星々が、一つ一つの計算によって織りなされた、美しい夜空のようだった。

コマンドエンコーディングと実行


        const commandEncoder = device.createCommandEncoder();
        const passEncoder = commandEncoder.beginComputePass();
        passEncoder.setPipeline(computePipeline);
        passEncoder.setBindGroup(0, bindGroup);
        passEncoder.dispatchWorkgroups(matrixSize, matrixSize);
        passEncoder.end();

        commandEncoder.copyBufferToBuffer(gpuBufferC, 0, gpuBufferRead, 0, matrixC.byteLength);

        const start = performance.now();
        const commands = commandEncoder.finish();
        device.queue.submit([commands]);

        await gpuBufferRead.mapAsync(GPUMapMode.READ);
        const arrayBuffer = gpuBufferRead.getMappedRange();
        const result = new Float32Array(arrayBuffer);
        const end = performance.now();

        outputElement.textContent = `行列C（結果）:\n${result}\n処理時間: ${(end - start).toFixed(2)} ms`;

        gpuBufferRead.unmap();

「行列C、計算結果だ…これが僕の作り上げたものなんだ。」浩司は画面に表示された結果を見て、自分の仕事が形になった瞬間を噛み締めた。処理時間も、彼の期待通り、驚くほど短かった。彼の努力が実を結び、東京の夜空のように美しい結果が得られたのだ。

image.png

その夜、浩司はパソコンの前で微笑みながら、満足感に包まれていた。行列の星々が、彼のプログラムの中で静かに輝いていた。東京の繁忙な街並みの中で、彼は一人、技術と美の融合を実感していた。

そして、夜が更けると、浩司はまた明日への準備をしながら、静かに机を片付けた。彼の心には、新たな挑戦と可能性が広がっていた。

説明。

WebGPUを使って4x4の行列の掛け算を実行し、結果を表示するシンプルなWebアプリケーションです。以下に、コードの各部分について詳細に解説します。

  1. HTML 構造

  2. JavaScript コード

WebGPUの初期化

async function initWebGPU() {
    if (!navigator.gpu) {
        console.log("WebGPUはこのブラウザでサポートされていません。");
        return;
    }

    try {
        const adapter = await navigator.gpu.requestAdapter();
        if (!adapter) {
            console.log("WebGPUアダプターの取得に失敗しました。");
            return;
        }
        
        const device = await adapter.requestDevice();
        const outputElement = document.getElementById('output');

navigator.gpu を使ってWebGPUがサポートされているか確認します。
サポートされていない場合は、メッセージをコンソールに表示します。
WebGPUアダプターを取得し、デバイスをリクエストします。取得できない場合はエラーメッセージを表示します。

行列の初期化

        const matrixSize = 4;
        const matrixA = new Float32Array(matrixSize * matrixSize);
        const matrixB = new Float32Array(matrixSize * matrixSize);
        const matrixC = new Float32Array(matrixSize * matrixSize);

        for (let i = 0; i < matrixSize * matrixSize; i++) {
            matrixA[i] = Math.random();
            matrixB[i] = Math.random();
        }

4x4の行列を表すために、Float32Arrayを使用して行列A、B、Cを作成します。
行列AとBにランダムな値を設定します。

WebGPU バッファの作成とデータ転送

        const gpuBufferA = device.createBuffer({
            size: matrixA.byteLength,
            usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
            mappedAtCreation: true
        });
        new Float32Array(gpuBufferA.getMappedRange()).set(matrixA);
        gpuBufferA.unmap();

        const gpuBufferB = device.createBuffer({
            size: matrixB.byteLength,
            usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
            mappedAtCreation: true
        });
        new Float32Array(gpuBufferB.getMappedRange()).set(matrixB);
        gpuBufferB.unmap();

        const gpuBufferC = device.createBuffer({
            size: matrixC.byteLength,
            usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST
        });

        const gpuBufferRead = device.createBuffer({
            size: matrixC.byteLength,
            usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
        });

行列AとBのデータをWebGPUバッファに転送します。GPUBufferUsage.STORAGE と GPUBufferUsage.COPY_DST は、バッファの使用用途を指定します。
行列Cを保存するためのバッファと、結果を読み取るためのバッファを作成します。

シェーダコードの作成

        const shaderCode = `
            @group(0) @binding(0) var<storage, read> matrixA : array<f32>;
            @group(0) @binding(1) var<storage, read> matrixB : array<f32>;
            @group(0) @binding(2) var<storage, read_write> matrixC : array<f32>;

            @compute @workgroup_size(1)
            fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
                let row = global_id.x;
                let col = global_id.y;
                if (row < 4u && col < 4u) {
                    let size = 4u;
                    var sum = 0.0;
                    for (var i = 0u; i < size; i = i + 1u) {
                        sum = sum + matrixA[row * size + i] * matrixB[i * size + col];
                    }
                    matrixC[row * size + col] = sum;
                }
            }
        `;

WGSL（WebGPU Shading Language）で行列の積を計算するシェーダコードを作成します。
@group(0) @binding(x) でバッファのバインディングを指定し、@compute でコンピュートシェーダを定義します。

コンピュートパイプラインとバインドグループの作成

        const shaderModule = device.createShaderModule({
            code: shaderCode
        });

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

        const bindGroup = device.createBindGroup({
            layout: computePipeline.getBindGroupLayout(0),
            entries: [
                { binding: 0, resource: { buffer: gpuBufferA }},
                { binding: 1, resource: { buffer: gpuBufferB }},
                { binding: 2, resource: { buffer: gpuBufferC }}
            ]
        });

シェーダモジュールを作成し、コンピュートパイプラインを作成します。
バインドグループを作成し、行列A、B、Cのバッファをシェーダにバインドします。

コマンドエンコーディングと実行

        const commandEncoder = device.createCommandEncoder();
        const passEncoder = commandEncoder.beginComputePass();
        passEncoder.setPipeline(computePipeline);
        passEncoder.setBindGroup(0, bindGroup);
        passEncoder.dispatchWorkgroups(matrixSize, matrixSize);
        passEncoder.end();

        commandEncoder.copyBufferToBuffer(gpuBufferC, 0, gpuBufferRead, 0, matrixC.byteLength);

        const start = performance.now();
        const commands = commandEncoder.finish();
        device.queue.submit([commands]);

        await gpuBufferRead.mapAsync(GPUMapMode.READ);
        const arrayBuffer = gpuBufferRead.getMappedRange();
        const result = new Float32Array(arrayBuffer);
        const end = performance.now();

        outputElement.textContent = `行列C（結果）:\n${result}\n処理時間: ${(end - start).toFixed(2)} ms`;

        gpuBufferRead.unmap();

コマンドエンコーダを使って、コンピュートパスを開始し、パイプラインとバインドグループを設定します。
行列Cの計算結果をバッファからCPUにコピーし、処理が完了するまで待機します。
結果を読み取り、画面に表示します。

まとめ
このコードは、WebGPUを用いて行列の掛け算を並列に計算し、その結果と処理時間をブラウザに表示します。

コード。

初期化部分: navigator.gpu を使って WebGPU がサポートされているか確認し、サポートされていない場合はメッセージを表示。
デバイスの取得: WebGPU アダプターとデバイスをリクエストし、取得できない場合はエラーメッセージを表示。
行列の初期化: 行列AとBをランダムな値で初期化し、それを WebGPU バッファにアップロード。
シェーダコード: WGSL（WebGPU Shading Language）で行列の積を計算するシェーダコードを作成。
パイプラインとバインドグループ: コンピュートパイプラインとバインドグループを作成して、計算のための設定を行う。
コマンドエンコーダ: 計算のコマンドを作成し、GPUに送信して処理を実行。
結果の読み取り: 結果を CPU にコピーし、画面に表示。

<!DOCTYPE html>
<html lang="en">
<head>
    <meta charset="UTF-8">
    <meta name="viewport" content="width=device-width, initial-scale=1.0">
    <title>WebGPU Matrix Multiplication</title>
</head>
<body>
    <h1>WebGPU Matrix Multiplication</h1>
    <pre id="output"></pre>
    <script type="module">
        async function initWebGPU() {
            // WebGPUがサポートされているか確認する
            if (!navigator.gpu) {
                console.log("WebGPUはこのブラウザでサポートされていません。");
                return;
            }

            try {
                // WebGPUアダプターを取得
                const adapter = await navigator.gpu.requestAdapter();
                if (!adapter) {
                    console.log("WebGPUアダプターの取得に失敗しました。");
                    return;
                }
                
                // デバイスをリクエスト
                const device = await adapter.requestDevice();
                const outputElement = document.getElementById('output');

                // 行列サイズを設定（4x4行列）
                const matrixSize = 4;
                const matrixA = new Float32Array(matrixSize * matrixSize);
                const matrixB = new Float32Array(matrixSize * matrixSize);
                const matrixC = new Float32Array(matrixSize * matrixSize);

                // 行列AとBをランダムな値で初期化
                for (let i = 0; i < matrixSize * matrixSize; i++) {
                    matrixA[i] = Math.random();
                    matrixB[i] = Math.random();
                }

                // WebGPUバッファを作成し、行列Aをアップロード
                const gpuBufferA = device.createBuffer({
                    size: matrixA.byteLength,
                    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
                    mappedAtCreation: true
                });
                new Float32Array(gpuBufferA.getMappedRange()).set(matrixA);
                gpuBufferA.unmap();

                // WebGPUバッファを作成し、行列Bをアップロード
                const gpuBufferB = device.createBuffer({
                    size: matrixB.byteLength,
                    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
                    mappedAtCreation: true
                });
                new Float32Array(gpuBufferB.getMappedRange()).set(matrixB);
                gpuBufferB.unmap();

                // 行列Cを保存するためのWebGPUバッファを作成
                const gpuBufferC = device.createBuffer({
                    size: matrixC.byteLength,
                    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST
                });

                // 行列Cの結果を読み取るためのバッファを作成
                const gpuBufferRead = device.createBuffer({
                    size: matrixC.byteLength,
                    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
                });

                // WGSLシェーダコードを作成
                const shaderCode = `
                    @group(0) @binding(0) var<storage, read> matrixA : array<f32>;
                    @group(0) @binding(1) var<storage, read> matrixB : array<f32>;
                    @group(0) @binding(2) var<storage, read_write> matrixC : array<f32>;

                    @compute @workgroup_size(1)
                    fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
                        let row = global_id.x;
                        let col = global_id.y;
                        if (row < 4u && col < 4u) {
                            let size = 4u;
                            var sum = 0.0;
                            for (var i = 0u; i < size; i = i + 1u) {
                                sum = sum + matrixA[row * size + i] * matrixB[i * size + col];
                            }
                            matrixC[row * size + col] = sum;
                        }
                    }
                `;

                // シェーダモジュールを作成
                const shaderModule = device.createShaderModule({
                    code: shaderCode
                });

                // コンピュートパイプラインを作成
                const computePipeline = device.createComputePipeline({
                    layout: "auto",
                    compute: {
                        module: shaderModule,
                        entryPoint: "main"
                    }
                });

                // バインドグループを作成
                const bindGroup = device.createBindGroup({
                    layout: computePipeline.getBindGroupLayout(0),
                    entries: [
                        { binding: 0, resource: { buffer: gpuBufferA }},
                        { binding: 1, resource: { buffer: gpuBufferB }},
                        { binding: 2, resource: { buffer: gpuBufferC }}
                    ]
                });

                // コマンドエンコーダを作成
                const commandEncoder = device.createCommandEncoder();
                const passEncoder = commandEncoder.beginComputePass();
                passEncoder.setPipeline(computePipeline);
                passEncoder.setBindGroup(0, bindGroup);
                passEncoder.dispatchWorkgroups(matrixSize, matrixSize);
                passEncoder.end(); // 修正：end()を使用

                // 行列Cの結果をバッファにコピー
                commandEncoder.copyBufferToBuffer(gpuBufferC, 0, gpuBufferRead, 0, matrixC.byteLength);

                // コマンドを送信して実行
                const start = performance.now();
                const commands = commandEncoder.finish();
                device.queue.submit([commands]);

                // 結果を読み取る
                await gpuBufferRead.mapAsync(GPUMapMode.READ);
                const arrayBuffer = gpuBufferRead.getMappedRange();
                const result = new Float32Array(arrayBuffer);
                const end = performance.now();

                // 結果を表示
                outputElement.textContent = `行列C（結果）:\n${result}\n処理時間: ${(end - start).toFixed(2)} ms`;

                gpuBufferRead.unmap();
            } catch (error) {
                console.error("エラーが発生しました:", error);
            }
        }

        initWebGPU();
    </script>
</body>
</html>

参考。

変更が必要な箇所のまとめ
データの型定義（バッファや変数の型、array から array> などの変更）
計算ロジック（ベクトルの加算から行列の加算への変更など）
ワークグループサイズとスレッドの割り当て方法（処理の並列化方法に応じて）
これらのポイントを適切に変更すれば、異なる計算を行うシェーダを容易に作成できます。

WGSLでは、スカラー、ベクトル、行列（マトリックス）、および配列として3次元以上のデータも扱えますが、直接「テンソル」という型の概念は存在しません。ただし、テンソルのようなデータ構造（多次元配列）を扱うことは可能です。

以下、WGSLがサポートする具体的なデータ型について説明します。

  1. スカラー（Scalar）
    スカラーは、単一の値を持つ基本的なデータ型です。WGSLでは以下のスカラー型がサポートされています。

i32 : 32ビット符号付き整数
u32 : 32ビット符号なし整数
f32 : 32ビット浮動小数点数
bool : 真偽値
例: スカラー変数の宣言

var x: f32;
var y: i32;
2. ベクトル（Vector）
ベクトルは、スカラー値の並びです。WGSLでは以下の形式でベクトルを定義します。

vec2 : 2要素のベクトル
vec3 : 3要素のベクトル
vec4 : 4要素のベクトル
ここで T は i32, u32, f32 のいずれかです。
例: ベクトル変数の宣言

var velocity: vec3;
var position: vec4;
3. 行列（Matrix）
行列は、ベクトルの並びで構成され、数値演算に使われます。WGSLでは、matNxM の形式で行列を定義します。N は行数、M は列数を表します。

mat2x2 : 2×2行列
mat3x3 : 3×3行列
mat4x4 : 4×4行列
例: 行列変数の宣言

var transformation: mat4x4;
4. 配列（Array）
配列を使うことで、3次元やそれ以上のデータを扱うことができます。WGSLの配列は、テンソルのように多次元データを表現できますが、配列の要素自体はベクトルや行列であるため、柔軟に多次元データを扱うことができます。

配列の定義は array 形式で、T は要素の型、N は要素数です。
例: 配列変数の宣言

var data: array, 100>; // 100個の3次元ベクトル
3次元データの配列（テンソルのように扱う例）

var tensor: array, 4>, 4>; // 4×4×4の3次元配列
ここで、tensorは4×4×4の3次元配列として機能し、テンソルのように扱えます。このように、WGSLは多次元の配列をネストして使用することで、実質的にテンソルと同様の構造を表現することができます。

  1. テンソルの扱い
    WGSL自体に「テンソル」という専用の型はありませんが、多次元配列（ネストされた配列）や、ベクトル・行列を組み合わせることでテンソルに相当するデータ構造を作ることができます。

例: 3次元テンソルのようなデータの宣言

var tensor3d: array, 4>, 4>; // 4x4x4のテンソル相当
この例では、各要素が4次元ベクトルである3次元配列を使用し、テンソルのように扱います。

まとめ
WGSLはスカラー、ベクトル、行列の基本型をサポートしています。
配列を使用して、3次元以上のデータ構造を扱うことができます。配列はテンソルのように扱うことができ、多次元データに対応可能です。
テンソルそのものの型はないものの、多次元配列をネストしてテンソルと同様の計算が可能です。
このように、WGSLを使ってスカラー、ベクトル、行列、さらにはテンソル的な多次元配列を使った計算を行うことができます。

2
3
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
2
3