/*

main.js
  webgpuを使ってみる

*/

async function init() {
  let ren = new webgpuControl("test")
  await ren.setup()
  ren.setCode(`
  struct Input {
    data: array<f32>,
  };

  struct Output {
    data: array<f32>,
  };

  @group(0) @binding(0)
  var<storage, read> input : Input;

  @group(0) @binding(1)
  var<storage, read_write> output : Output;

  @compute @workgroup_size(4)
  fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
    output.data[global_id.x] = input.data[global_id.x] * 2.0;
  }
  `)
  ren.setBuf(0, 16, "DST", {data : new Float32Array([1, 2, 3, 4])})
  ren.setBuf(1, 16, "SRC")
  ren.run(4, 1, 1, "main")
  let obj = await ren.getBuf(1, 16, 0)

  let ary = new Float32Array(obj[0])
  console.log(ary.byteLength)
  console.log(ary)
  for(let i = 0; i < 4; i++) {
    console.log(i + " : " + ary[i])
  }

  ren.unmap(obj)

  console.log(ren.debug())
}

init()

WebGPU(webgpuControl.js)

/*

webgpuControl.js
  WebGPUのComputerShaderを使いやすくしたラッパークラス。
  exanple : setup -> setCode -> setBuf -> run -> getBuf -> unmap

*/

class webgpuControl {
  constructor(label = "") {
    this.label = label
    this.bufList = []
    this.error = undefined
  }

  async setup() {
    if (!navigator.gpu) {
      this.error = "WebGPU not supported."
      return 0
    }
    else {
      this.adapter = await navigator.gpu.requestAdapter()
      if (!this.adapter) {
        this.error = "Couldn't request WebGPU adapter."
        return 0
      }
      this.device = await this.adapter.requestDevice({label : this.label})
    }
    return 0
  }

  setCode(code = " ") {
    if(this.error === undefined) {
      this.computeShaderModule = this.device.createShaderModule({code: code, label : (this.label + " : setCode")})
    }
  }

  setBuf(bindingIndex = 0, byteSize = 0, bufTypebuf = undefined, buf = undefined) {
    if(this.error === undefined) {
      let flag;
      if(bufTypebuf === "DST") {
        flag = GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
      }
      else if(bufTypebuf === "SRC") {
        flag = GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
      }
      else if(bufTypebuf === "UNIFORM") {
        flag = GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST
      }
      else {//"NON"
        flag = GPUBufferUsage.STORAGE
      }
      this.bufList[bindingIndex] = {
        binding : bindingIndex,
        resource : {
          buffer : this.device.createBuffer({
            size : byteSize, usage : flag,
            label : (this.label + " : setBuf(" + bindingIndex + ")")
          })
        }
      }

      if(buf !== undefined) {
        this.device.queue.writeBuffer(this.bufList[bindingIndex].resource.buffer, 0, buf.data, buf.offset, buf.size)
      }
    }
  }

  run(xLen = 1, yLen = 1, zLen = 1, entryPoint = "main") {
    if(this.error === undefined) {
      let computePipeline = this.device.createComputePipeline({
        layout : "auto",
        compute : {
          module : this.computeShaderModule,
          entryPoint : entryPoint
        },
        label : (this.label + " : run_createComputePipeline")
      })

      let bindGroup = this.device.createBindGroup({
        layout : computePipeline.getBindGroupLayout(0),
        entries: this.bufList,
        label : (this.label + " : run_createBindGroup")
      })
      
      let commandEncoder = this.device.createCommandEncoder({label : (this.label + " : run_createCommandEncoder")})
      let passEncoder = commandEncoder.beginComputePass({label : (this.label + " : run_beginComputePass")})
      passEncoder.setPipeline(computePipeline)
      passEncoder.setBindGroup(0, bindGroup)
      passEncoder.dispatchWorkgroups(xLen, yLen, zLen)
      passEncoder.end()
      
      this.device.queue.submit([commandEncoder.finish()])
    }
  }

  async getBuf(bindingIndex = 0, size = 0, offset = 0) {
    if(this.error === undefined) {
      let stagingBuffer = this.device.createBuffer({
        mappedAtCreation : false,
        size : size,
        usage : GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
        label : (this.label + " : getBuf(" + bindingIndex + ")")
      })
      
      let copyEncoder = this.device.createCommandEncoder({label : (this.label + " : getBuf_createCommandEncoder")})
      copyEncoder.copyBufferToBuffer(
        this.bufList[bindingIndex].resource.buffer, offset,
        stagingBuffer, 0, size
      )

      this.device.queue.submit([copyEncoder.finish()])
      await stagingBuffer.mapAsync(GPUMapMode.READ)
      let copyArrayBuffer = stagingBuffer.getMappedRange()
      
      return [copyArrayBuffer, stagingBuffer]
    }
  }

  unmap(obj = undefined) {
    if(this.error === undefined) {
      obj[1].unmap()
    }
  }

  debug() {
    return this
  }
}

c++AMP

//amp.hをincludeする。defineで警告を止める
#define _SILENCE_AMP_DEPRECATION_WARNINGS
#include <amp.h>
#include <amp_math.h>
#include <windows.h>
#include <new>
#include <iostream>
#include <cmath>
#include <vector>

double time() {
    LARGE_INTEGER freq;//速度計測用
    QueryPerformanceFrequency(&freq);
    LARGE_INTEGER start;
    QueryPerformanceCounter(&start);

    return static_cast<double>(start.QuadPart) * 1000.0 / freq.QuadPart;
}

int main() {
    std::cout << "gpu test" << std::endl;
    float t = 12345;
    float s = 1.2345;
    std::vector<float> work;

    for (int size = 1; size < (9000000); size *= 3) {
        work.resize(size);

        for (int i = 0; i < size; i++) work[i] = i;

        double t1 = time();
        for (int i = 0; i < size; i++) work[i] = sqrt(pow(work[i], s)) / t;
        std::cout << "size = " << size << std::endl << time() - t1 << std::endl;

        for (int i = 0; i < size; i++) work[i] = i;

        t1 = time();
        concurrency::array_view<float, 1> ta_(size, work);
        //gpuでの処理をラムダ式で渡す。
        concurrency::parallel_for_each(
            ta_.extent,
            [=](concurrency::index<1> ix) restrict(amp) {
                ta_[ix] = concurrency::fast_math::sqrt(concurrency::fast_math::pow(ta_[ix], s)) / t;
            });
        std::cout << time() - t1 << std::endl;
    }

    system("pause");
	return 0;
}

hlsl(ComputeShader_M.hpp)

/*

ComputeShader関連

*/

#pragma once
#include <string>
#include <vector>
#include <windows.h>
#include <d3d11.h>
#include <d3dcompiler.h>
#pragma comment(lib,"d3d11.lib")
#pragma comment(lib,"d3dcompiler.lib")

class ComputeShader_M {

	//direct11固有のオブジェクトとか
	ID3D11Device* dev = nullptr;
	ID3D11DeviceContext* devContext = nullptr;
	ID3D11ComputeShader* pComputeShader = nullptr;

	//cpuとgpuを行き来するbuff用
	std::vector<ID3D11Buffer*> pBuf_SRV;
	std::vector<ID3D11ShaderResourceView*> pBufSRV;
	std::vector<ID3D11Buffer*> pBuf_UAV;
	std::vector<ID3D11UnorderedAccessView*> pBufUAV;

	std::string error_str = "code_error::";
	int error_num = 0;//エラー情報
	/*
	0 エラーなし
	-1 CreateDevice失敗
	-2 デバイスがcompute shaderに対応していない
	-3 hlslファイルのコンパイル失敗
	-4 シェーダーオブジェクトの作成失敗
	-5 SRVのCreateBuffer失敗
	-6 CreateShaderResourceView失敗
	-7 UAVのCreateBuffer失敗
	-8 CreateUnorderedAccessView失敗
	-9 gpu→cpuのCreateBuffer失敗
	-10 devContext->Map失敗
	*/

public:

	//コンストラクタ。デバイス情報を取得する
	ComputeShader_M() {
		UINT createDeviceFlags = 0;
#ifdef _DEBUG//when debug
		createDeviceFlags |= D3D11_CREATE_DEVICE_DEBUG;
#endif

		D3D_FEATURE_LEVEL lvl[] = { D3D_FEATURE_LEVEL_11_1, D3D_FEATURE_LEVEL_11_0,
									D3D_FEATURE_LEVEL_10_1, D3D_FEATURE_LEVEL_10_0 };

		HRESULT hr = D3D11CreateDevice(
			nullptr, D3D_DRIVER_TYPE_HARDWARE, nullptr,
			createDeviceFlags, &lvl[0], _countof(lvl),
			D3D11_SDK_VERSION, &dev, nullptr, &devContext);

		//lvl[0]がだめなときにlvl[1]でも試してみる
		if (hr == E_INVALIDARG) {
			hr = D3D11CreateDevice(
				nullptr, D3D_DRIVER_TYPE_HARDWARE, nullptr,
				createDeviceFlags, &lvl[1], _countof(lvl) - 1,
				D3D11_SDK_VERSION, &dev, nullptr, &devContext);
		}

		if (FAILED(hr)) {
			error_num = -1;
		}
		else {
			//デバイスがcompute shaderに対応しているか確認する
			if (dev->GetFeatureLevel() < D3D_FEATURE_LEVEL_11_0) {
				D3D11_FEATURE_DATA_D3D10_X_HARDWARE_OPTIONS hwopts = { 0 };
				dev->CheckFeatureSupport(D3D11_FEATURE_D3D10_X_HARDWARE_OPTIONS, &hwopts, sizeof(hwopts));
				if (!hwopts.ComputeShaders_Plus_RawAndStructuredBuffers_Via_Shader_4_x) {//ここらへん、まったくわからん。https://learn.microsoft.com/ja-jp/windows/win32/direct3d11/direct3d-11-advanced-stages-compute-create
					error_num = -2;
				}
			}
		}
	}

	//基本情報のセット、及びhlslのコンパイル
	bool set(char* code, long code_size, char* Entrypoint, long SRV_num, long UAV_num) {
		//error情報の有無
		if (error_num < 0) return false;

		//SRVとUAVの数に応じてvectorを初期化
		{
			//余分な分は解放しておく
			if (pBuf_SRV.size() > SRV_num)
				for (auto i = SRV_num - 1; i < pBuf_SRV.size(); i++)
					delete_SRV(i);

			if (pBuf_UAV.size() > UAV_num)
				for (auto i = UAV_num - 1; i < pBuf_UAV.size(); i++)
					delete_UAV(i);

			pBuf_SRV.resize(SRV_num, nullptr);
			pBufSRV.resize(SRV_num, nullptr);
			pBuf_UAV.resize(UAV_num, nullptr);
			pBufUAV.resize(UAV_num, nullptr);
		}

		//前のやつがあれば解放
		clean();

		//hlslファイルをコンパイル
		ID3DBlob* ppCode = nullptr;
		{
			LPCSTR pTarget = (dev->GetFeatureLevel() >= D3D_FEATURE_LEVEL_11_0) ? "cs_5_0" : "cs_4_0";
			UINT flag1 = D3DCOMPILE_ENABLE_STRICTNESS;
#if defined( DEBUG ) || defined( _DEBUG )
			flag1 |= D3DCOMPILE_DEBUG;
#endif
			ID3DBlob* ppErrorMsgs = nullptr;

			HRESULT hr = D3DCompile(
				code, code_size, nullptr, nullptr,
				D3D_COMPILE_STANDARD_FILE_INCLUDE,
				Entrypoint, pTarget, flag1,
				0, &ppCode, &ppErrorMsgs);

			/*HRESULT hr = D3DCompileFromFile(
				L"ComputeShader.hlsl", NULL, D3D_COMPILE_STANDARD_FILE_INCLUDE,
				"main", pTarget, flag1,
				0, &ppCode, &ppErrorMsgs);*/

			if (FAILED(hr)) {
				if (ppErrorMsgs) {
					//エラーメッセージ
					//printf("error:%s\n", (char*)ppErrorMsgs->GetBufferPointer());
					std::string t = (char*)ppErrorMsgs->GetBufferPointer();
					error_str = error_str + t;
					ppErrorMsgs->Release();
				}

				if (ppCode) ppCode->Release();

				error_num = -3; return false;
			}
		}

		//シェーダーオブジェクトの作成
		{
			HRESULT hr = dev->CreateComputeShader(ppCode->GetBufferPointer(), ppCode->GetBufferSize(), nullptr, &pComputeShader);
			ppCode->Release();
			if (FAILED(hr)) {
				error_num = -4; return false;
			}
		}

		return true;
	}

	//シェーダーリソースビューの設定
	bool set_SRV_buff(void* buff, long struct_size, long num, long index) {
		//error情報の有無
		if (error_num < 0) return false;

		//前のものがあれば解放
		delete_SRV(index);

		ID3D11Buffer* t_buff = nullptr;
		ID3D11ShaderResourceView* tbuff = nullptr;

		D3D11_BUFFER_DESC INdesc;
		ZeroMemory(&INdesc, sizeof(D3D11_BUFFER_DESC));
		INdesc.ByteWidth = struct_size * num;//バッファのサイズ
		INdesc.BindFlags = D3D11_BIND_UNORDERED_ACCESS | D3D11_BIND_SHADER_RESOURCE;
		INdesc.MiscFlags = D3D11_RESOURCE_MISC_BUFFER_STRUCTURED;
		INdesc.StructureByteStride = struct_size;

		HRESULT hr;
		if (buff == nullptr) {
			hr = dev->CreateBuffer(&INdesc, nullptr, &t_buff);
		}
		else {
			D3D11_SUBRESOURCE_DATA pInitialData;
			pInitialData.pSysMem = buff;
			hr = dev->CreateBuffer(&INdesc, &pInitialData, &t_buff);
		}

		if (FAILED(hr)) {
			error_num = -5; return false;
		}

		D3D11_SHADER_RESOURCE_VIEW_DESC INdesc2;
		ZeroMemory(&INdesc2, sizeof(D3D11_SHADER_RESOURCE_VIEW_DESC));
		INdesc2.Format = DXGI_FORMAT_UNKNOWN;
		INdesc2.ViewDimension = D3D11_SRV_DIMENSION_BUFFEREX;
		INdesc2.BufferEx.FirstElement = 0;
		INdesc2.BufferEx.NumElements = num;

		hr = dev->CreateShaderResourceView(t_buff, &INdesc2, &tbuff);

		if (FAILED(hr)) {
			t_buff->Release(); error_num = -6; return false;
		}

		//vectorに書きこむ
		pBuf_SRV[index] = t_buff;
		pBufSRV[index] = tbuff;

		return true;
	}

	//アンオーダードアクセスビューの設定
	bool set_UAV_buff(void* buff, long struct_size, long num, long index) {
		//error情報の有無
		if (error_num < 0) return false;

		//前のものがあれば解放
		delete_UAV(index);

		ID3D11Buffer* t_buff = nullptr;
		ID3D11UnorderedAccessView* tbuff = nullptr;

		D3D11_BUFFER_DESC OUTdesc;
		ZeroMemory(&OUTdesc, sizeof(D3D11_BUFFER_DESC));
		OUTdesc.ByteWidth = struct_size * num;//バッファのサイズ
		OUTdesc.BindFlags = D3D11_BIND_UNORDERED_ACCESS | D3D11_BIND_SHADER_RESOURCE;
		OUTdesc.MiscFlags = D3D11_RESOURCE_MISC_BUFFER_STRUCTURED;
		OUTdesc.StructureByteStride = struct_size;

		HRESULT hr;
		if (buff == nullptr) {
			hr = dev->CreateBuffer(&OUTdesc, nullptr, &t_buff);
		}
		else {
			D3D11_SUBRESOURCE_DATA pInitialData;
			pInitialData.pSysMem = buff;
			hr = dev->CreateBuffer(&OUTdesc, &pInitialData, &t_buff);
		}

		if (FAILED(hr)) {
			error_num = -7; return false;
		}

		D3D11_UNORDERED_ACCESS_VIEW_DESC OUTdesc2;
		ZeroMemory(&OUTdesc2, sizeof(D3D11_UNORDERED_ACCESS_VIEW_DESC));
		OUTdesc2.Format = DXGI_FORMAT_UNKNOWN;
		OUTdesc2.ViewDimension = D3D11_UAV_DIMENSION_BUFFER;
		OUTdesc2.Buffer.FirstElement = 0;
		OUTdesc2.Buffer.NumElements = num;

		hr = dev->CreateUnorderedAccessView(t_buff, &OUTdesc2, &tbuff);

		if (FAILED(hr)) {
			t_buff->Release(); error_num = -8; return false;
		}

		//vectorに書きこむ
		pBuf_UAV[index] = t_buff;
		pBufUAV[index] = tbuff;

		return true;
	}

	//hlslの実行
	bool run(long x, long y, long z) {
		//error情報の有無
		if (error_num < 0) return false;

		devContext->CSSetShader(pComputeShader, nullptr, 0);
		if (pBufSRV.size() > 0) devContext->CSSetShaderResources(0, pBufSRV.size(), &pBufSRV[0]);//シェーダーリソースビューを設定
		if (pBufUAV.size() > 0) devContext->CSSetUnorderedAccessViews(0, pBufUAV.size(), &pBufUAV[0], nullptr);//アンオーダードアクセスビューを設定
		devContext->Dispatch(x, y, z);//実行

		//後処理(バインドされたものを解放？するみたい
		devContext->CSSetShader(nullptr, nullptr, 0);
		if (pBufSRV.size() > 0) {
			std::vector<ID3D11UnorderedAccessView*> pIN(pBufSRV.size(), nullptr);
			devContext->CSSetUnorderedAccessViews(0, pBufSRV.size(), &pIN[0], nullptr);
		}
		if (pBufUAV.size() > 0) {
			std::vector<ID3D11ShaderResourceView*>pOUT(pBufUAV.size(), nullptr);
			devContext->CSSetShaderResources(0, pBufUAV.size(), &pOUT[0]);
		}
		ID3D11Buffer* pBUff[1] = { nullptr };
		devContext->CSSetConstantBuffers(0, 1, pBUff);//定数バッファはこれでいいのかな？？？

		return true;
	}

	//アンオーダードアクセスビューの内容をcpuに書き出す
	bool write_buff(void* buff, long buff_size, long index) {
		//error情報の有無
		if (error_num < 0) return false;

		ID3D11Buffer* buff0 = nullptr;
		D3D11_BUFFER_DESC desc;
		ZeroMemory(&desc, sizeof(D3D11_BUFFER_DESC));
		pBuf_UAV[index]->GetDesc(&desc);
		desc.Usage = D3D11_USAGE_STAGING;
		desc.BindFlags = 0;
		desc.CPUAccessFlags = D3D11_CPU_ACCESS_READ;
		desc.MiscFlags = 0;

		HRESULT hr = dev->CreateBuffer(&desc, nullptr, &buff0);//buff0を作成

		if (FAILED(hr)) {
			error_num = -9; return false;
		}

		devContext->CopyResource(buff0, pBuf_UAV[index]);//pBuf_OUTの内容をbuff0にコピーする

		D3D11_MAPPED_SUBRESOURCE pt;
		if (FAILED(devContext->Map(buff0, 0, D3D11_MAP_READ, 0, &pt))) {//buff0の生ポインタを取得する
			buff0->Release();
			error_num = -10; return false;
		}

		unsigned char* t = reinterpret_cast<unsigned char*>(buff);
		unsigned char* p = reinterpret_cast<unsigned char*>(pt.pData);
		for (auto i = 0; i < buff_size; i++) *(t + i) = *(p + i);

		devContext->Unmap(buff0, 0);
		buff0->Release();

		return true;
	}

	//エラー情報を渡す
	std::string debug_info(int* num) {
		*num = error_num;
		return error_str;
	}

	//前のシェーダーをReleaseする
	void clean() {
		if (pComputeShader != nullptr) { pComputeShader->Release(); pComputeShader = nullptr; }
	}

	//SRVをReleaseする
	void delete_SRV(int index) {
		if (pBuf_SRV[index] != nullptr) { pBuf_SRV[index]->Release();  pBuf_SRV[index] = nullptr; }
		if (pBufSRV[index] != nullptr) { pBufSRV[index]->Release();  pBufSRV[index] = nullptr; }
	}

	//UAVをReleaseする
	void delete_UAV(int index) {
		if (pBuf_UAV[index] != nullptr) { pBuf_UAV[index]->Release();  pBuf_UAV[index] = nullptr; }
		if (pBufUAV[index] != nullptr) { pBufUAV[index]->Release();  pBufUAV[index] = nullptr; }
	}

	//後処理
	void end() {
		for (auto i = 0; i < pBuf_SRV.size(); i++) delete_SRV(i);
		for (auto i = 0; i < pBuf_UAV.size(); i++) delete_UAV(i);
		clean();
		if (devContext != nullptr) { devContext->Release(); devContext = nullptr; }
		if (dev != nullptr) { dev->Release(); dev = nullptr; }
	}

	//デストラクタ
	~ComputeShader_M() {
		end();
	}

};

導入

導入ではGPGPUについて簡単に説明します。

GPGPUってなに？

普通のプログラムはCPUというパソコンの部品で実行されます。ですが、最近のパソコンにはGPUという部品もあって、GPUもCPUも四則演算を始めとした演算を行うことができます。

CPUとGPUにはそれぞれ得意分野があって、CPUが得意な処理はCPUで、GPUが得意な処理はGPUで実行しようというのがGPGPUの考え方です。

GPUが得意な処理はGPUに任せるというのがGPGPUの考え方です。

GPUでやると良い処理

CPUとGPUでは得意な処理が異なります。ざっくり言うとCPUは複雑な処理が得意で、GPUは単純な処理が得意です。

具体的な得意分野を挙げると…

CPU
・if文やfor文の多い処理。
・分岐がいっぱいの複雑な処理。

GPU
・if文やfor文の少ない簡単な処理。
・並列に実行できる処理。

GPUには並列に実行できる簡単な処理を任せると良きです。

WebGPUのComputeShader

GPGPUを行う方法はいくつかありますが、この記事ではWebGPUのComputeShaderを用いた方法を紹介します。WebGPUはGPUをブラウザから使用できるJavaScriptAPIです。

WebGPUは現在実験的に導入されいる機能なので、ブラウザごとに互換性を確認する必要があります。

全体像

全体像ではWebGPUの仕組みなどについて触れていきます。大まかな流れは他のシェーダー言語と変わらないです。

CPUとGPU

CPUとGPUのメモリは別の場所にあります。なので、CPUで宣言した値をGPUで参照することができません。同様に、GPUで計算した値をCPUから参照することができません。

そこでWebGPUの機能を使用することで、CPUのメモリ上の値をGPUのメモリにコピーしたり、GPUのメモリ上の値をCPUのメモリにコピーすることでデータのやりとりを行います。

以下はイメージです。

CPUとGPUのデータのやりとりにはWebGPUを使います。

シェーダー言語

CPUで実行するプログラムはJavaScriptなどで記述されますが、GPUで実行するプログラムはシェーダー言語というもので記述されます。WebGPUで採用されているシェーダー言語はWGSLというものです。

通常、シェーダー言語(WGSL)は処理の実行前にコンパイルすることでデバイスごとの最適化を図ります。

WebGPUでは、GPUで実行されるプログラムを記述するためにWGSLというシェーダー言語を使用します。

処理の流れ

WebGPUでComputeShaderをする際の大まかな流れを確認します。

1. 初期化。
2. WGSLコードをコンパイルする。
3. GPUの計算に必要なバッファを用意する。
4. 2と3を引数にシェーダーを実行する。
5. 計算結果がGPUのメモリにあるので、それをCPUのメモリにコピーする。

やることは以上の5つです。

この項では一般的な処理の流れを紹介しています。実際には3や5の処理を飛ばしても問題はありません。(実用性もありませんが笑)

JavaScript(CPU)側のコード

JavaScript(CPU)側のコードについて解説していきます。

初期化

  if (!navigator.gpu) {
    //ブラウザがGPUのサポートをしていないときの処理
    //...
  }

  //物理デバイスの取得
  const adapter = await navigator.gpu.requestAdapter()
  if (!adapter) {
    //WebGPUに対応したGPUがない場合の処理
    //...
  }

  //論理デバイスの取得
  device = await adapter.requestDevice()

　
まずは、ブラウザがWebGPUに対応しているかどうかを調べます。グローバル変数であるnavigator.gpuを確認することで調べることができます。

  if (!navigator.gpu) {
    //ブラウザがGPUのサポートをしていないときの処理
    //...
  }

　
次に全体のベースとなるオブジェクトを取得します。
adapterは物理デバイスを表し、そこからrequestDevice関数でdeviceという論理デバイスを取得します。返り値がpromiseであることに注意してください。

また、物理デバイスがnullならデバイスにWebGPUに対応したGPUがないということになります。

  //物理デバイスの取得
  const adapter = await navigator.gpu.requestAdapter()
  if (!adapter) {
    //WebGPUに対応したGPUがない場合の処理
    //...
  }

  //論理デバイスの取得
  device = await adapter.requestDevice()

WGSLコードのコンパイル

WGSL(GPUで実行するためのプログラム)をコンパイルします。WGSLの書き方や文法については後述します。

  //WGSLコード
  const shader_code = ``

  //コンパイル
  const computeShaderModule = device.createShaderModule({code: shader_code})

バッファの準備

GPUで使いたいメモリをcreateBuffer関数で作成します。CPU上のメモリをGPU上のメモリにコピーしたい場合はwriteBuffer関数でコピーします。GPUBuffer(WebGPUで扱うためのバッファ)は使いたいバッファの分だけ用意する必要があります。

  //GPUに送りたいバッファ
  inputArray = new Float32Array([1, 2, 3, 4])

  //GPUBuffer(webgpuで扱うためのバッファ)作成
  inputBuffer = device.createBuffer({
    size: inputArray.byteLength,//バイト数
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST//フラグ
  })

  //コピー
  device.queue.writeBuffer(inputBuffer, 0, inputArray)

　
createBuffer関数のフラグには以下のようなものがあります。

//コピー操作の宛先となる場合(writeBuffer使用時)
GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST

//コピー操作のソースとなる場合(GPU -> CPU のようにバッファをコピーする場合)
GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC

//定数バッファとして使う場合
GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST

//GPU内で使用できるバッファ(バッファの値は未定義)
GPUBufferUsage.STORAGE

　
writeBuffer関数では引数でコピー位置などを指定することができます。

device.queue.writeBuffer(
  inputBuffer,//コピー先のGPUBuffer
  0,//GPUBufferのどこからコピーを始めるか(バイト単位)
  inputArray,//コピー元のバッファ(ArrayBuffer,TypedArray, DataViewのいずれか)
  0,//コピーの開始位置(※)
  inputArray.length//コピーするサイズ(※)
)

※第4引数と第5引数の値はコピー元のバッファの種類によって指定方法が変わります。コピー元のバッファがTypedArrayの場合は要素の数、TypedArray以外の場合にはバイト単位で指定します。また、第4引数と第5引数は省略可能です。(省略時の既定値は、第4引数は0、第5引数は第4引数で指定された値からバッファの最後までのサイズ)

シェーダーの実行

シェーダーを実行します。以下は例です、

  //パイプラインの作成
  const computePipeline = device.createComputePipeline({
    layout: "auto",
    compute: {
      module: computeShaderModule,
      entryPoint: "main"
    }
  })

  //GPUBufferをBindGroupにする(WGSLのタグと対応)
  const bindGroup = device.createBindGroup({
    layout: computePipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: { buffer: inputBuffer } },
      { binding: 1, resource: { buffer: outputBuffer } }
      //...ここは必要な分だけ
    ]
  })

  //コマンドを書き込んでいく
  const commandEncoder = device.createCommandEncoder()
  const passEncoder = commandEncoder.beginComputePass()
  passEncoder.setPipeline(computePipeline)
  passEncoder.setBindGroup(0, bindGroup)
  passEncoder.dispatchWorkgroups(4, 1, 1)
  passEncoder.end()//コマンド書き込み終了

  //処理が終わるのを待つ
  device.queue.submit([commandEncoder.finish()])

　
まずはコンピュートシェーダー用のパイプラインを作成します。

  //パイプラインの作成
  const computePipeline = device.createComputePipeline({
    layout: "auto",
    compute: {
      module: computeShaderModule,
        //createShaderModuleで作成したものです
      entryPoint: "main"
        //WGSLコード内で実行する関数名です
    }
  })

layoutにはパイプラインで使用されるGPUリソース(GPUBufferなど)の構造情報を設定します。
構造情報を設定する方法は、
1. createPipelineLayout関数で構造情報を作成する。
2. "auto"文字列を設定することでWGSLコードから自動的に構造情報を作成する。
…の二種類があります。

今回は"auto"文字列を使用していますが、createPipelineLayout関数で構造情報を作成する場合は構造情報を使いまわすことができるので、シェーダーを複数回実行する場合などに適しています。

詳細は以下を確認してください。
createPipelineLayout
createBindGroupLayout
　
　
次はBindGroup(GPUリソースのリスト)を作成していきます。

  //GPUBufferをBindGroupにする(WGSLのタグと対応)
  const bindGroup = device.createBindGroup({
    layout: computePipeline.getBindGroupLayout(0),
    entries: [
      { binding: 0, resource: { buffer: inputBuffer } },
      { binding: 1, resource: { buffer: outputBuffer } }
      //...ここは必要な分だけ
    ]
  })

createBindGroup関数のlayoutに渡すものは、createComputePipeline関数のlayoutに何を設定したかで変わります。
1. createPipelineLayout関数だった場合
　-> createBindGroupLayout関数で作成したGPUBindGroupLayoutを渡します。
2. "auto"文字列を指定した場合
　-> computePipeline.getBindGroupLayout関数で作成したGPUBindGroupLayoutを渡します。

bindingの数字はWGSLのタグと対応するようにします。

　
最後にコマンドを書き込んでいきます。

  //GPUCommandEncoderの作成
  const commandEncoder = device.createCommandEncoder()

  //コンピューティングパスのエンコードを開始
  const passEncoder = commandEncoder.beginComputePass()

  //パイプラインを設定
  passEncoder.setPipeline(computePipeline)

  //bindGroupを設定
  passEncoder.setBindGroup(0, bindGroup)

  //シェーダーの実行
  passEncoder.dispatchWorkgroups(4, 1, 1)

  //パスのエンコード終了
  passEncoder.end()

  //処理が終わるのを待つ
  device.queue.submit([commandEncoder.finish()])

WebGPUのシェーダーはGPUCommandEncoderオブジェクトにコマンドを発行することで実行されます。

setBindGroup関数の第一引数はWGSLのタグと対応している必要があります。
dispatchWorkgroups関数の引数はx、y、zの順に実行するスレッドの数です。createComputePipeline関数のentryPointで指定された関数が、xyzの乗算の分だけ実行されます。つまり、(100,1,1)と(25,2,2)では同じ回数実行されることになります。

WebGPUのシェーダーは非同期に実行されます。

そこで、submit関数を使用してシェーダーの終了まで待つことができます。

GPUのメモリにあるバッファを読み取る

GPUにあるメモリを読み取るには、GPUとCPUどちらもアクセスすることができるステージングバッファを介します。

  const stagingBuffer = device.createBuffer({
    mappedAtCreation: false,
      //trueで即時にマップします。今回は後からマップします
    size: 16,
    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
      //読み取り用にマップ
  })

  const copyEncoder = device.createCommandEncoder()
  copyEncoder.copyBufferToBuffer(//コピーする
    outputBuffer, //source buffer,
    0, //source offset,
    stagingBuffer, //destination buffer,
    0, //destination offset,
    16 //size
  )
  device.queue.submit([copyEncoder.finish()])//処理終了まで待つ

  await stagingBuffer.mapAsync(GPUMapMode.READ)//マップ

  //ArrayBufferの参照を取得できます
  const copyArrayBuffer = stagingBuffer.getMappedRange()

  //ここでなんらかの処理

  stagingBuffer.unmap()//マップ解除(copyArrayBufferの参照も切れる)

読みたいGPU上のメモリをcopyBufferToBuffer関数でステージングバッファにコピーした後、マップすることでArrayBufferの参照を取得することができます。用事が済んだらステージングバッファのマップを解除しましょう。

ステージングバッファのデータは通常の配列のように書き込みできないため、そのような操作をしたい場合にはCPUのバッファにコピーすることでWebGPUから完全に独立したバッファを作成できます。

  const copyArrayBuffer = stagingBuffer.getMappedRange()
  let data = new Float32Array(copyArrayBuffer.slice(0))
  stagingBuffer.unmap()//マップ解除(copyArrayBufferの参照も切れる)

  //マップ解除したので参照できない
  let data2 = new Float32Array(copyArrayBuffer)
  console.log(data2[0])

  //dataはバッファがCPUにあるので参照できる
  console.log(data[0])

WGSL(GPU)側のコード

WGSL(GPU)側のコードについて解説していきます。公式リファレンスもありますが少し長いので簡単にまとめています。

変数の型

変数の型にはいくつか種類があります。

keyword	type
i32	32ビット符号付き整数
u32	32ビット符号なし整数
f32	32ビット浮動小数点数
bool	ブール値

WGSLは型定義が厳密なので、常に変数の型を意識する必要があります。

構文

0. コメント

多くの方にとっては馴染み深い方法だと思います。

//この行はコメント

/*複数行
あ
い
う
ここまで*/

1. 変数宣言

データ型

letは定数値を宣言するときに使います。varは普通の変数、constはコンパイル時定数です。

let a = 10;//型は自動で決まる
let b: f32 = 10;//型指定
var c: f32 = 10;//普通の変数
const d: i32 = 10;//コンパイル時定数

b = b + 1;//letは定数値なのでエラー
c = c + 1;//ok

fn add(a: f32) -> f32 {
  const result = a + 1;//コンパイル時に決まらないのでエラー
  return result;
}

ベクトル型

2~4つの同じ型の変数をまとめたベクトル型というものを使用できます。

let a = vec2<i32>(0, -2);
let b = vec4<f32>(20, 3.4, 3.1, -32);

//xyzw、rgba、indexのいずれかで各要素にアクセスできます
let c = a.x;//0
let d = b.a;//-32
let e = b[0];//20

//複数の要素にアクセスできます
let f = vec2<f32>(b[2], b[0]);
let g = b.zx;
let h = b.br;

行列

行列を扱うことができます。ベクトル型の配列のようなイメージです。行も列も2~4の範囲で作成することができます。デフォルトではf32しか使用できません。

let a = mat4x4f(1, 0, 0, 0,
                0, 1, 0, 0,
                0, 0, 1, 0,
                0, 0, 0, 1);

let b = a[1][1];//1
let c = a[2];//vec4<f32>(0, 0, 1, 0)

//dとeは同じです
let d = mat2x3f(2, -3, 3,
                1, 0, 2.1);
let e = mat2x3f(vec3<f32>(2, -3, 3),
                vec3<f32>(1, 0, 2.1));

デフォルトで行列用の関数が用意されていますが、数が少ないので外部ライブラリの使用をおすすめします。また、簡単な関数であれば自分で実装するのも一つの選択肢です。
コンストラクタ、距離
 転値
 行列ライブラリ

配列

型と要素数の指定で宣言できます。動的配列は今のところありません。(ストレージバッファは別です)

let a = array<f32, 2>(2, 5);
let b = array<vec4f, 3>;

let c = a[1];//5

//eとfは同じです
let d = vec3f(0, 1, 2);
let e = array<vec3f, 3>(d, d, d);
let f = array(d, d, d);//型推論

構造体

構造体はメンバ変数と型を対応させて書いていきます。最後のコンマはあってもなくても変わりません。

struct DataA {
  a: i32,
  b: f32,
  c: u32,
}

struct DataB {//DataAと同じ
  a: i32,
  b: f32,
  c: u32
}

let a: DataA;//DataA構造体の宣言
let b = a.b;

struct DataC {
  a: i32,
  b: vec2<f32>,
  c: mat2x3f,
  d: DataA
}

アライメント

アライメント(構造体メモリアクセスの効率を上げるための仕様)に注意する必要があります。アライメントを考慮すると、DataDのサイズは24バイトになります。

struct DataD {
  a: f32,        //size 4
                 //padding size 4
  b: vec2<f32>,  //size 8
  c: f32,        //size 4
                 //padding size 4
}

member	size	align	offset
DataD	24	8	/
a: f32	4	4	0
padding	4	/	4
b: vec2	8	8	8
c: f32	4	4	12
padding	4	/	16

アライメントの仕様は以下の二点です。
1. データ型やベクトル型に応じて固有のアライメントがある。メンバ変数のoffset位置は常にアライメントの倍数である必要がある。倍数でない場合は倍数になるようにパディング(空のメモリ)で調整される。
2. 構造体のアライメントは、その要素の中で一番大きいアライメントが採用される。
　
配列もアライメントに注意しなければならないときがあります。vec3<f32>のsizeは12ですが、alignは16です。vec3<f32>には4のパディング(空のメモリ)があるので、bのサイズは48になります。

let a: array<f32, 3>;//size 4*3
/*
f32  size 4  align 4

a[0] size 4  offset 0
a[1] size 4  offset 4
a[2] size 4  offset 8
*/

let b: array<vec3<f32>, 3>;//size 16*3
/*
vec3<f32>  size 12  align 16

b[0]       size 12  offset 0
   padding size 4   offset 12
b[1]       size 12  offset 16
   padding size 4   offset 28
b[2]       size 12  offset 32
   padding size 4   offset 44
*/

また、uniform属性を持った定数バッファはすべてのメンバ変数が16以上でアライメントされる必要があります。

struct DataA {
  a: f32,
  @align(16) b: f32//align 4 -> 16
}
@group(0) @binding(0) var<uniform> data: DataA;


//配列も16でアライメントされる必要があるので、DataBのようにします。
struct DataB {
  @size(16) elem: f32
}
struct DataC {
  a: array<DataB, 8>
}
@group(0) @binding(1) var<uniform> data2: DataC;

各データ型のアライメントサイズについてはこちらにあるアライメント表を参考にしてください。

2. 四則演算

変数だけでなく、ベクトル型も四則演算することができます。累乗は＾ではなくpow関数であることに注意してください。

var a: f32 = 5;
var b: f32 = 2;
var c: u32 = 4;

var d: f32 = 0;
d = a + b;//7
d = a - b;//3
d = a * b;//10
d = a / b;//2.5
d = a % b;//1

var e = vec2<i32>(0, -2);
var f = vec2<i32>(20, 30);
var g = e + f;//(20, 28)

　
異なる型同士(行列とベクトル型など)では四則記号の意味合いが変わる場合があります。詳細はこちらを確認してください。

3. 代入

WGSLは型定義が厳密です。違う型同士で代入を行うとエラーになります。違う型の変数を代入したいときはf32関数などのコンストラクタを使用して型変換します。

var a: u32 = 10;
var b: f32 = 20;

var c: f32 = a;//error
var c: f32 = f32(a)//ok

var d: f32 = a + b;//error
var d: f32 = f32(a) + b;//ok

コンストラクタ一覧はこちらを参照してください。

4. 条件分岐

if文

  let a = 10;
  if(a == 1) {
    //...
  }
  else if(a == 2) {
    //...
  }
  else {
    //...
  }

  //条件式の()はなくてもよい
  if a == 0 {
    //...
  }

switch

caseはu32またはi32の定数である必要があります。

let x : i32 = 10;
switch x {
  case 0: {//caseは定数の整数である必要がある
    //...
  }
  default {//defaultの順番はどこでもよい
    //...
  }
  case 1, 2, {//複数指定も可
    //...
  }
  case 3, {//カンマはあってもなくてもよい
    //...
  }
  case 4 {
    //...
  }
}

const c = 3;
let x : i32 = 10;
switch x {
  case c: {//constで宣言した変数は使用可
    //...
  }
  case 0, default, {//変数とdefaultでも組める
    //...
  }
}

all

引数のベクトル型がすべてtrueであるときにtrueを返します。

let a = all(vec2<bool>(true, true));//true

any

引数のベクトル型のいずれかの要素がtrueであるときにtrueを返します。

let a = any(vec2<bool>(true, false));//true

select

三項演算子のように処理されます。tとfは通常のデータ型かベクトルです。flagはtrueかfalseである必要があります。

let a = flag ? t : f;//JavaScript
let a = select(f, t, flag);//WGSL

5. 反復処理

break

反復処理とswitch文でのみ使用できます。

  var i = 0;
  loop {
    i = i + 1;
    if i >= 10 { break; }
  }

continue

反復処理でのみ使用できます。continueした後はcontinuing内の処理を行ってから次のループに入ります。

  var i = 0;
  loop {
    i = i + 1;
    if i >= 5 { continue; }

    continuing {//continuingはなくてもよい
      //...
      break if i >= 10;//※
    }
  }

※continuingはあってもなくても大丈夫ですが、continuing内でbreakする場合は最後にbreak ifを使用します。

for

  for(var i = 0; i < 10; i++) {
    //...
  }

loop

{//←スコープがあるとiはスコープ外で参照できない
  var i = 0;
  loop {
    i = i + 1;
    if(i >= 10) {
      break;
    }
  }
}

while

  var i = 0;
  while(i < 10) {
    i++;
  }

6. 論理演算子

if文などで使用すると思います。

a == b //イコール
a || b //論理和
a && b //論理積
a != b //論理否定
a > b  //aはbより大きい
a < b  //aはbより小さい
a >= b //aはb以上
a <= b //aはb以下

7. 関数

自分でオリジナルの関数を宣言することができます。引数と戻り値の型を決める必要があります。

fn a(b: f32) -> i32 {//戻り値の型を`->`の後に指定する
  return i32(b * 2);
}

//戻り値はなくてもよい
fn c(b: f32) {
  return;//returnはあってもなくてもよい
}

//複数の引数も可
fn d(b: f32, e: vec2<f32>) { return; }

また、三角関数などのいくつかの関数はデフォルトで用意されています。
関数の一覧はリファレンスで確認できます。

8. ConputeShaderで実行される関数

JavaScript(cpu)側のコードでエントリポイントとして指定される関数の記述は、普通の関数と異なります。

  @compute @workgroup_size(4, 1, 1)
  fn main() {
    //...
  }

  @compute @workgroup_size(4)
  fn nin(@builtin(global_invocation_id) global_id : vec3<u32>) {
    //...
  }

関数宣言の前に@computeと書くことででコンピュートシェーダーであることを示します。@workgroup_sizeではx、y、zのそれぞれの方向にいくつのスレッドを走らせるかを決めます。yとzは省略可能です。(既定値は1)
　
実際に実行されるスレッド数はdispatchWorkgroups関数で指定された値と@workgroup_sizeで指定した値との積になります。例えば、以下のような場合には100個のスレッドが走ることになります。

  //WGSL(一つのworkgroupで2*5*1=10のスレッドが走る)
  ...
  @compute @workgroup_size(2, 5, 1)
  fn....

  //Javascript(workgroupが10*1*1=10走る)
  ...
  passEncoder.dispatchWorkgroups(10, 1, 1)
  ...

  //実際に走るのは、workgroup_sizeのxyzとdispatchWorkgroupsのxyzとの乗算の結果
  //一つのworkgroup内で走るスレッド 2*5*1=10
  //workgroupの数 10*1*1=10
  //全体で走るスレッドの数 2*5*1*10*1*1=100

スレッドの数指定は6次元の自由度があることになります。
　
コンピュートシェーダーの引数は@builtinを使用したものを使います。@builtinで必要な値を引数として参照できます。

  //JavaScript(JX, JY, JZ)
  @compute @workgroup_size(JX, JY, JZ)

  //WGSL(WX, WY, WZ)
  @compute @workgroup_size(WX, WY, WZ)
  fn nin(
    @builtin(local_invocation_id) local_id : vec3<u32>,
      //workgroup内での位置を三次元で表したもの。(0,0,0)~(WX-1,WY-1,WZ-1)
    @builtin(local_invocation_index) local_index : u32,
      //workgroup内での位置を一次元で表したもの。0~WX*WY*WZ-1
    @builtin(global_invocation_id) global_id : vec3<u32>,
      //workgroup_id * workgroup_size + local_invocation_id
    @builtin(workgroup_id) workgroupId : vec3<u32>,
      //workgroupの位置。(0,0,0)~(JX-1,JY-1,JZ-1)
    @builtin(num_workgroups) numWork : vec3<u32>
      //dispatchWorkgroups関数で指定された値。(JX,JY,JZ)
    ) {
    //...
    let a = local_id.x;//関数内で引数を参照できる
  }

　
例えば、以下のようなコードでは16個のスレッドが動きます。

  //WGSL(一つのworkgroupで2*2*1=4のスレッドが走る)
  @compute @workgroup_size(2, 2, 1)

  //Javascript(workgroupが2*2*1=4走る)
  passEncoder.dispatchWorkgroups(2, 2, 1)

workgroup_id	local_invocation_id	local_invocation_index	global_invocation_id
(0,0,0)	(0,0,0)	0	(0,0,0)
(0,0,0)	(1,0,0)	1	(1,0,0)
(0,0,0)	(0,1,0)	2	(0,1,0)
(0,0,0)	(1,1,0)	3	(1,1,0)

workgroup_id	local_invocation_id	local_invocation_index	global_invocation_id
(1,0,0)	(0,0,0)	0	(2,0,0)
(1,0,0)	(1,0,0)	1	(3,0,0)
(1,0,0)	(0,1,0)	2	(2,1,0)
(1,0,0)	(1,1,0)	3	(3,1,0)

workgroup_id	local_invocation_id	local_invocation_index	global_invocation_id
(0,1,0)	(0,0,0)	0	(0,2,0)
(0,1,0)	(1,0,0)	1	(1,2,0)
(0,1,0)	(0,1,0)	2	(0,3,0)
(0,1,0)	(1,1,0)	3	(1,3,0)

workgroup_id	local_invocation_id	local_invocation_index	global_invocation_id
(1,1,0)	(0,0,0)	0	(2,2,0)
(1,1,0)	(1,0,0)	1	(3,2,0)
(1,1,0)	(0,1,0)	2	(2,3,0)
(1,1,0)	(1,1,0)	3	(3,3,0)

global_invocation_idに関しては図で確認すると分かり易いです。

9. bindされた変数

JavaScript(CPU)側のコードでbindされた変数は、bindの番号とgrupeの番号から参照することができます。この際、アドレス空間と型を指定します。アクセス修飾子は任意で指定します。また、配列の場合は要素数を明示しないことで、JavaScript(CPU)側のコードから配列の長さを決めることができます。

//readを指定した読み込み専用バッファ
@group(0) @binding(0)
var<storage, read> data : array<f32>;

//read_writeを指定した読み書き両方できるバッファ
@group(0) @binding(1)
var<storage, read_write> data2 : array<u32>;

@compute @workgroup_size(1, 1, 1)
fn main(@builtin(workgroup_id) workgroupId : vec3<u32>) {
  data2[workgroupId.x] = data[workgroupId.x];//data2は書き込み可
  data[workgroupId.x] = f32(10);//error : dataには書き込めない
}

アドレス空間はcreateBufferする際のフラグを参考にすることが多いです。

//コピー操作の宛先となる場合(writeBuffer使用時)
GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
  -> storage, read

//コピー操作のソースとなる場合(GPU -> CPU のようにバッファをコピーする場合)
GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
  -> storage, read_write

//定数バッファとして使う場合
GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST
  -> uniform

//GPU内で使用できるバッファ(バッファの値は未定義)
GPUBufferUsage.STORAGE
  -> storage, read_write

アドレス空間には以下のようなものがあります。

type	Default access mode	Notes
function	read_write	ローカル変数(関数内で宣言)
private	read_write	グローバル変数(関数外で宣言)
workgroup	read_write	worker毎に用意される専用バッファ
uniform	read	定数バッファ
storage	read	ストレージバッファ
handle	read	テクスチャ用

アクセス修飾子は三種類あります。

type	Notes
read	read onry
write	write onry
read_write	read and write

10. ビット操作

var a: u32 = 0x00000001;//十六進数宣言
var b: u32 = 0x00000100;

var c = a << 1;//左シフト 0x00000010
var d = b >> 1;//右シフト 0x00000010
var e = a | b;//論理和 0x00000101
var f = a & b;//論理積 0x00000000
var g = ~a;//ビット反転 0xfffffffe
var h = a ^ b;//排他的論理和 0x00000101

11. ポインタ

参照型とポインタ型があります。アドレス空間がstorageの場合、デフォルトのアクセス修飾子はreadです。参照型とポインタ型のどちらもポインタを扱う型ですが、参照型はポインタ型と比べて制約の多いものとなります。

var<private> a: i32 = 10;
let b: ref<private, i32> = &a;//変数aの参照型
var c: ptr<private, i32> = &a;//変数aのポインタ型

var d = *b;//10
var e = *c;//10

通常の変数のポインタを得るには&を使用して&aのようにします。逆にポインタから変数の値を得るには*を使用して*bのようにします。
　
参照型の制約には以下のようなものがあります。
・let変数は参照型にできない。
・関数の引数にすることができない。
・宣言時に初期化が必要。
・再代入不可。

構造体などのメンバ変数を参照するには以下のようにします。

struct Data {
  a: i32,
  b: vec2<f32>,
  c: mat2x3f,
}

fn a(b: f32) -> i32 {
  var<function> a: Data;
  let b: ref<function, i32> = &a;//変数aの参照型

  a.a = i32(10);
  return (*b).a;//10
}

　
ポインタ型は関数の引数にすることができます。

fn f(a: ptr<function, i32>) {
  *a = *a + 1;
}

@compute @workgroup_size(1, 1, 1)
fn main() {
  var i: i32 = 0;
  f(&i);
  let b: i32 = i;//1
}

　
特定のポインタにアクセスするようなコードはコンパイルエラーになります。

let b: ref<private, i32> = 0;//error : 0番地にアクセスしようとしている

atomic

ComputeShaderで実行される関数は並列に実行されます。そのため、複数のスレッドが同じ変数に書き込みを行う際は注意する必要があります。

例えば、スレッド1とスレッド2が変数aにアクセスする際、処理の順番によって結界が変わります。これはスレッド1と2がそれぞれ非同期に実行されるためです。

これでは結果が2になるときと3になるときが発生してしまいます。
そこでatomic操作をすることで、片方のスレッドが変数aにアクセスしている間は他のスレッドが変数aにアクセスできないようにすることができます。

　
atomic操作はatomic型でのみ使用可能です。変数の型はu32かi32である必要があり、アドレス空間はworkgroupかstorageでアクセス修飾子はread_writeである必要があります。

@group(0) @binding(0)
var<storage, read_write> data : array<atomic<u32>>;
  //atomic型の配列

@compute @workgroup_size(1, 1, 1)
fn main(@builtin(workgroup_id) workgroupId : vec3<u32>) {
  var a: ptr<storage, u32> = &data[0];//ポインタ型
  var b: u32 = atomicLoad(a);//atomic変数の読み取り
  atomicStore(a, u32(10));//atomic変数に書き込む
  var c: u32 = atomicExchange(a, u32(10));//atomic変数の置き換え
    //戻り値はもともとatomic変数にあった値

  //演算用のatomic関数
  var c: u32 = atomicAdd(a, u32(1));//加算
  var d: u32 = atomicSub(a, u32(1));//減算
  var e: u32 = atomicMax(a, u32(1));//MAX関数の挙動
  var f: u32 = atomicMin(a, u32(1));//MIN関数の挙動
  var g: u32 = atomicAnd(a, 0x01);//論理積
  var h: u32 = atomicOr(a, 0x01);//論理和
  var i: u32 = atomicXor(a, 0x01);//排他的論理和
}

atomic型はatomic関数でのみアクセスできます。atomic関数の中ではatomic操作を行うので、他のスレッドとぶつかることはないです。演算用のatomic関数はポインタの指すatomic型の値と第二引数の値とでそれぞれの演算を行い、結果をポインタ先に入れます。戻り値はもともとポインタ先にあった値です。
　
一部の実装では等価比較が上手くいかない場合がありますが、以下のような関数もあります。

  var a: ptr<storage, u32> = &data[0];//ポインタ型
  var b: u32 = atomicCompareExchangeWeak(a, u32(7), u32(10));
    /*
      1.ポインタ先の値と第二引数を等価比較
      2. -> true : ポインタ先に第三引数の値を入れる
         -> false : なにもしない
      3.戻り値は構造体
        b.old_value : もともとポインタ先にあった値
        b.exchanged : bool型で等価比較の結果
    */

同期処理

ConputeShaderは基本的に非同期に実行されますが、アルゴリズムによっては処理の足並みを揃えたいときがあります。

　
同期関数には三種類あります。どの関数も同じworkgroup内のスレッドが同期します。(異なるworkgroup同士での同期は今のところできません。)

var<workgroup> data: array<i32, 10>;

@compute @workgroup_size(4, 4, 4)
fn main() {
  storageBarrier();
    //この関数が呼ばれる前の処理で、
    //ストレージバッファの操作があった場合に
    //同じworkgroup内のスレッドがその操作を
    //終わらせるまで待機する

  workgroupBarrier();
    //この関数が呼ばれる前の処理で、
    //ワークグループバッファの操作があった場合に
    //同じworkgroup内のスレッドがその操作を
    //終わらせるまで待機する

  var p: ptr<workgroup, i32> = &data[0];
  let a: i32 = workgroupUniformLoad(p);
    //workgroupBarrierのように同期します
    //戻り値はポインタ先の値です
    //引数のポインタの位置は同じworkgroup内で
    //統一する必要があります
    //(同じworkgroupだと引数と戻り値の値は同じになる)
}

　
以下のような処理はコンパイルエラーになります。同期関数が実行されないスレッドがでてくるからです。

var<workgroup> data: array<i32, 10>;

@compute @workgroup_size(4, 4, 4)
fn main(@builtin(local_invocation_index) local_index : u32) {
  if(local_index % 2 == 0) {
    workgroupBarrier();
  }

  /*
  local_invocation_indexが奇数のとき、
  同期関数は呼ばれないので、関数はすぐ終了する。
  
  local_invocation_indexが偶数のとき、
  同期関数が呼ばれるので、
  関数は同じworkgroupのスレッドの処理がworkgroupBarrierにくるまで待機。
  -> local_invocation_indexが奇数のスレッドはworkgroupBarrierを呼ばない。
  -> 偶数スレッドは永遠に待機することになってしまう。
  */
}

pack

GPU内で使用できるメモリはそこまで大きくありません。そのため、できるだけ少ないメモリで多くの情報量を持てると嬉しいです。そこで複数の変数を一つ分のメモリを持った変数にパッケージ化することができます。(もとの型に比べて精度は落ちますが、そこまで精度が大事でない場合はパッケージ化しても影響が少ない場合が多いです。)

  //-1.0~1.0のvec4をu32にする
  let a = vec4<f32>(0, 1, 0.5, -0.3);
  let b: u32 = pack4x8snorm(a);//pack
  let c: vec4<f32> = unpack4x8snorm(b);//unpack

  //0.0~1.0のvec4をu32にする
  let d = vec4<f32>(0, 1, 0.5, -0.3);
    //-0.3は範囲外なので0.0になる
  let e: u32 = pack4x8unorm(d);//pack
  let f: vec4<f32> = unpack4x8unorm(e);//unpack

  //-1.0~1.0のvec2をu32にする
  let g = vec2<f32>(0, 1);
  let h: u32 = pack2x16snorm(g);//pack 
  let i: vec2<f32> = unpack2x16snorm(h);//unpack

  //0.0~1.0のvec2をu32にする
  let j = vec2<f32>(0, 1);
  let k: u32 = pack2x16unorm(j);//pack
  let l: vec2<f32> = unpack2x16unorm(k);//unpack

  //f32のvec2をu32にする(精度は落ちる)
  let n = vec2<f32>(100, -0.03);
  let m: u32 = pack2x16float(n);//pack
  let o: vec2<f32> = unpack2x16float(m);//unpack

どの関数も範囲外の場合は丸められます。

pack2x16floatは他の関数と違って、値の上限と下限が実行環境によって変わります。そのためもとの式がconst式とパイプラインの作成時に評価できる式の場合にエラーになります。

pack2x16floatの大雑把な処理方法は以下です。
16bitで表わせる場合 -> そのまま。
16bitで表わせない場合 ->下の方の精度は捨てられる。また、大きい(上限以上)又は小さい(下限以上)の値の場合無限として扱われる

制限について

GPUの性能によっては実行できないコードというものが存在します。例えば、最大スレッドが10万のGPUでは20万のスレッドを実行する処理ができません。

limit

WebGPUでは論理デバイスのlimitsプロパティを確認することでGPUの制限を確認することができます。

name	notes
maxComputeInvocationsPerWorkgroup	workgrupe_sizeの最大数(xyzの乗算)
maxComputeWorkgroupSizeX	workgrupe_sizeの最大数(x)
maxComputeWorkgroupSizeY	workgrupe_sizeの最大数(y)
maxComputeWorkgroupSizeZ	workgrupe_sizeの最大数(z)
maxComputeWorkgroupsPerDimension	dispatchWorkgroupsの最大数(xyzそれぞれがこれ以下)
maxBufferSize	GPUBufferの最大サイズ
maxUniformBufferBindingSize	uniformバッファの最大サイズ
maxBindGroups	1つのパイプラインで作成できるBindGroupの最大数

上記は一例ですが、これらの制限に引っかかるとWebGPUはエラーを出します。プログラムを書くときはこれらの制限以内の値になるように気を付けなければならないです。制限の一覧はこちらを確認してください。

制限を引き上げる

ブラウザ以外にもGPUを使うソフトが存在する場合があります。

上図のような場合にブラウザがGPUのリソースを占有してしまうと、ブラウザ以外のソフトがGPUを使うことができません。そこで、WebGPUでは実際のGPUの物理的制限よりも低い値をlimitとして設けています。(limitのデフォルト値と一覧についてはこちらを確認してください。)

現在のlimitの制限をGPUの制限ギリギリまで引き上げたい際などはrequestDevice関数の引数にlimitオブジェクトを渡すことで、制限された値を上げることができます。

  if (!navigator.gpu) {
    //ブラウザがGPUのサポートをしていないときの処理
    //...
  }

  //物理デバイスの取得
  const adapter = await navigator.gpu.requestAdapter()
  if (!adapter) {
    //WebGPUに対応したGPUがない場合の処理
    //...
  }

  //GPUがBindGroups10個を扱えるかどうかを物理デバイスのlimitから確認
  if(adapter.limits.maxBindGroups >= 10) {
    let nlimit = {}
    nlimit.maxBindGroups = 10//メンバはdevice.limitsと同じ
      //空のメンバはデフォルト値になる

    //制限を引き上げた論理デバイスの取得
    device = await adapter.requestDevice({requiredLimits: nlimit})
  }
  else {
    //GPUがBindGroups10個を扱えない場合の処理
  }

現在使用しているGPUの最大の制限を適用したい場合は物理デバイスを表すadapterのlimitを使います。

    //制限を引き上げた論理デバイスの取得
    device = await adapter.requestDevice({requiredLimits: adapter.limits})

物理デバイスと仮想デバイスの違いに注意してください。シェーダーを実行するのは仮想デバイスなので、シェーダーの制限は仮想デバイスのlimitです。

tips

細かいテクニック？の紹介です。

他のshader

WebGPUにはComputeShaderの他にも2つのシェーダーがあります。VertexShader(頂点シェーダー)とFragmentShader(フラグメントシェーダー)です。どちらもグラフィックス用のシェーダーです。

汎用計算をするのならばComputeShaderを使用しますが、なにかをcanvas要素に描画したいときにはVertexShaderとFragmentShaderを使うことが多いです。

こちらのサンプル集では3つのシェーダーを組み合わせたコードも見ることができます。

エラー処理

WebGPUでエラーが発生した場合、多くの場合はエラー情報がコンソールに表示されます。ですが、それだけでは物足りない場合があります。WebGPUの便利なエラー処理方法について4つ紹介しましす。

1.label

WebGPUの多くの関数の引数にはlabelプロパティがあり、このlabelに文字列を設定するとエラー時のコンソールにlabel名が表示されます。

  device = await adapter.requestDevice({label : "test1"})
    //関数が失敗したとき、エラー文と"test1"という文字列が表示される

複数のシェーダーを実行した際に、どのlabelでエラーが起きたのかを確認することができるので便利です。

  device1 = await adapter1.requestDevice()//1
  device2 = await adapter2.requestDevice()//2
    //エラーが出ても、どっちでエラーが出ているか分からない

  device1 = await adapter1.requestDevice({label : "test1"})//1
  device2 = await adapter2.requestDevice({label : "test2"})//2
    //エラーが出るとラベル名が表示されるので、どこでエラーが発生したか明確

また、同じ関数を繰り返し使用する場合にも有効です。

  let buf = []
  for(let index = 0; i < 3; i++) {
    device.createBuffer({
      size : byteSize,
      usage : flag,
      label : "test1 : " + index
    })
  }

labelの使用例は参考文献にあるサンプルコード(webgpuControl.js)を見ると分かり易いです。labelはとても軽量であるため、リリースバージョンのアプリでも使用することが推奨されています。

2.Debug Groups

labelプロパティだけではエラーがどこで生じたのかを補足できない場合があります。その際は、DebugGroupsを使用するといい場合があります。

  //コマンドエンコーダーの作成
  let commandEncoder = device.createCommandEncoder()

  //スタックに"in the commandEncoder"を追加
  commandEncoder.pushDebugGroup("in the commandEncoder")

  //(1)
  //...

  //スタックから"in the commandEncoder"を削除
  commandEncoder.popDebugGroup()

  //処理の終了を待つ
  device.queue.submit([commandEncoder.finish()])

DebugGroupsはコマンドエンコーダーとコンピューティングパスエンコーダー、レンダーパスエンコーダー(ComputeShaderでは使用しない)の3つのオブジェクトに設定することができます。コマンドの書き込みなどでエラーが発生した際、スタックに文字列があればエラー情報と共にスタック内の文字列がコンソールに表示されます。例えば(1)でコマンドを書き込み、エラーが発生した場合は"in the commandEncoder"という文字列とエラー情報が表示されます。
　
DebugGroupsのpushDebugGroup関数とpopDebugGroup関数の数は対になっている必要があります。pushしたら必ずpopしましょう。

  let commandEncoder = device.createCommandEncoder()
  commandEncoder.pushDebugGroup("in the commandEncoder")//a_s

    let passEncoder = commandEncoder.beginComputePass()
    passEncoder.pushDebugGroup("in the passEncoder")//b_s

     passEncoder.setPipeline(computePipeline)
     passEncoder.setBindGroup(0, bindGroup)//(1)

     passEncoder.pushDebugGroup("dispatchWorkgroups")//c_s
       passEncoder.dispatchWorkgroups(xLen, yLen, zLen)//(2)
     passEncoder.popDebugGroup()//c_f

    passEncoder.popDebugGroup()//b_f
    passEncoder.end()

  commandEncoder.popDebugGroup()//a_f
  device.queue.submit([commandEncoder.finish()])

(1)でエラーが発生すると"in the commandEncoder"と"in the passEncoder"がエラー情報と一緒に表示されます。

(2)でエラーが発生すると"in the commandEncoder"、"in the passEncoder"と"dispatchWorkgroups"がエラー情報と一緒に表示されます。

DebugGroupsはlabelと同様にとても軽量であるため、リリースバージョンのアプリでも使用することが推奨されています。

3.Error Scopes

ErrorScopesでエラーを補足します。エラーが発生した時の処理を記述したい場合に便利です。

エラーの種類は三種類です。

type	notes
validation	引数が不正
out-of-memory	メモリ不足
internal	上記以外(※)

※引数が不正かメモリ不足かをWebGPUが識別できなかった場合に出るエラー。制限に違反したことを示す。(現在はパイプラインの作成によってのみ発生)

どの種類のエラーを補足するかはスタックのようなもので管理されています。pushErrorScope関数でスタックに追加することでエラーを補足でき、popErrorScope関数でスタックから削除します。戻り値のpromiseにエラーが発生した際の処理を書きます。

  //キューにvalidationを追加
  device.pushErrorScope("validation")

  //パイプラインを作成したり、
  //GPUBufferを作成してdispatchWorkgroupsしたり
  //...
  //...

  //promiseでエラー時の処理を
  device.popErrorScope().then(
    (error) => {console.log(error.message)}
  );

すべてのエラーを補足したい場合はpushとpopで囲ってしまいます。

  function e(error) {
    console.log(error.message)
  }

  device.pushErrorScope("validation")
  device.pushErrorScope("out-of-memory")
  device.pushErrorScope("internal")

  //...
  //...

  device.popErrorScope().then(e)
  device.popErrorScope().then(e)
  device.popErrorScope().then(e)

実際の実装では以下のようにすることで処理が失敗した時のプログラムを記述します。

  function e(...) {
    //internalになったときの処理
    /*(internalはシェーダーがデバイスに対して複雑な場合に起こることがあるので、単純な処理に置き換える)*/
    let computePipeline = device.createComputePipeline(...)
    //...
    //...
  }

  device.pushErrorScope("out-of-memory")

  device.pushErrorScope("internal")
  let computePipeline = device.createComputePipeline(...)
  device.popErrorScope().then(
    (error) => {e(...)}
  )

  //...
  device.pushErrorScope("validation")
  passEncoder.dispatchWorkgroups(xLen, yLen, zLen)
  device.popErrorScope().then((error) => {
    let maxNum = device.limits.maxComputeWorkgroupsPerDimension
    passEncoder.dispatchWorkgroups(
      xLen <= maxNum ? xLen : maxNum;,
      yLen <= maxNum ? yLen : maxNum;,
      zLen <= maxNum ? zLen : maxNum;)
  })
  //...

  device.popErrorScope().then(
    (error) => {console.log("out-of-memory : " + error.message)}
  )

ErrorScopesは非同期関数を考慮した処理を行いません。非同期関数内でErrorScopesを使う際は非同期に実行される他の処理に気を付ける必要があります。

//bad example
  async function create_a() {
    device.pushErrorScope("validation")
    let buf = await getBuffer(...)
      //非同期処理

    device.createBuffer({size : byteSize, usage : flag})

    device.popErrorScope().then(...)//(1)
  }

create_a()
computePipeline = device.createComputePipeline(...)
//処理順によっては、createComputePipelineのエラーが(1)にいってしまう

//better example
  async function create_a() {
    let buf = await getBuffer(...)
      //非同期処理

    //同期処理の部分だけ囲う
    device.pushErrorScope("validation")
    device.createBuffer({size : byteSize, usage : flag})
    device.popErrorScope().then(...)
  }

create_a()
computePipeline = device.createComputePipeline(...)

4.Uncaptured Errors

イベントリスナーにエラー時の処理を追加する方法です。ErrorScopesのスタックに文字列がない場合に機能します。一括でエラーを補足したい場合に便利です。

  //deviceにイベントを登録
  device.addEventListener("uncapturederror", (event) => {
    console.log(event.error.message)
  })

また、物理デバイス毎によるバグの頻度の検知などする際はこちらと組み合わせることも有効なようです。

5.Shaders

getCompilationInfo関数はWGSL内のエラーを補足します。(3と4の方法ではWGSLでのエラーを補足できません)

  let code ='...'//WGSLで記述されたプログラム
  let computeShaderModule = device.createShaderModule({code: code})
    //シェーダーコードをコンパイルする

  let error_info = await computeShaderModule.getCompilationInfo()
    //エラー情報を取得

  //エラー情報の表示
  for (let mes of error_info.messages) {
    console.log(mes.lineNum + " : " + mes.type + " : " + mes.message)
  }

getCompilationInfo関数はpromiseでGPUCompilationInfoオブジェクトを返します。GPUCompilationInfoオブジェクトのmessagesプロパティはエラーの詳細が含まれるGPUCompilationMessageオブジェクトの配列です。

GPUCompilationMessageオブジェクトのプロパティには以下のものがあります。

name	notes
length	エラー箇所の文字列の長さ
lineNum	行番号
linePos	UTF-16で行内でのエラー位置
message	エラー情報
offset	UTF-16でのオフセット位置
type	メッセージの種類(※)

※"error"、"info"、"warning"の三種類があります。

プロパティについての詳細はこちらを確認してください。

テクスチャ

画像や動画などの要素をGPU内で参照したい場合があります。その際にテクスチャをGPUで扱う方法がいくかあります。

1.GPUBuffer

canvas要素は画像データを表すArrayBufferを取得できます。それをGPUBufferに書き込むことでGPU内で参照することができます。また、GPUBufferをcanvasに書き込むことで画像処理のようなこともできます。

main.html

<!DOCTYPE html>
<html lang = "ja">
    <head>
        <title>test</title>
        <meta charset="UTF-8">

    <script src="webgpuControl.js"></script>

    </head>
    <body>
	<body bgcolor="#ffffff" text="000000">

        <h3>
            debug
        </h3>

	<!-- dom操作用の箱　-->
	<div id="domBox">
	</div>

	<!-- スクリプトの実行 -->
	<script src="main.js"></script>

    </body>
</html>

main.js

async function init() {
  //canvas作成
  let dataBox = document.getElementById("domBox")
  let bc = document.createElement("canvas")

  //idと位置を作成
  bc.id = "BaseCanvas"
  bc.width = 100
  bc.height = 100
  bc.style.position = "fixed"
  bc.style.left = 0 + "px"
  bc.style.top = 0 + "px"

  //色を塗る
  let ct = bc.getContext("2d")
  ct.fillStyle = "#000000"
  ct.fillRect(0, 0, bc.width, bc.height)
  ct.fillStyle = "#dddddd"
  ct.fillRect(10, 10, bc.width - 20, bc.height - 20)

  //domに追加
  dataBox.appendChild(bc)

  let ren = new webgpuControl("test")
  await ren.setup()

  ren.setCode(`
  struct Input {
    data: array<f32>,
  };

  struct Output {
    data: array<u32>,
  };

  @group(0) @binding(0)
  var<storage, read> input : Input;

  @group(0) @binding(1)
  var<storage, read_write> output : Output;

  @compute @workgroup_size(1)
  fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
    let x = input.data[i];
    let color: u32 = 0xff0000;//bgr
    output.data[global_id.x] = 0xff000000 | color;//a
  }
  `)

  let num = bc.width * bc.height
  let numo = bc.width * bc.height
  let buf = new Float32Array(num)
  for(let i = 0; i < num; i++) {
    buf[i] = i
  }

  ren.setBuf(0, num * 4, "DST", {data : buf})
  ren.setBuf(1, numo * 4, "SRC")
  ren.run(num, 1, 1, "main")
  let obj = await ren.getBuf(1, numo * 4, 0)

  //canvasに書き込む
  let new_array = obj[0].slice(0)
  let data = new ImageData(new Uint8ClampedArray(new_array), bc.width, bc.height)
  ct.putImageData(data, 0, 0)

  ren.unmap(obj)
}

init()

2.GPUTexture

GPUTextureを使う方法もあります。createImageBitmap関数で作成したビットマップデータをcraeteTexture関数で作成したGPUTextureにcopyExternalImageToTexture関数でコピーします。

createBindGroup関数のresourceプロパティにはcreateViewの戻り値を渡します。

WGSLでGPUTextureを宣言するにはこちらにある型を使用します。

参照するにはTexture Built-in Functionsを使用します。samplerを使うことで行列などを用いた座標変換などを行うこともできるようです。

3.GPUExternalTexture

HTMLVideoElementを参照するには、createBindGroup関数のresourceプロパティにimportExternalTexture関数の戻り値を渡します。

WGSLでの宣言及び参照方法はGPUTextureと同じです。

定数の上書き

WGSLコード内での定数宣言の際、overrideを付けることでJavaScript側から上書きできるようになります。対応している型はbool、i32、u32、f32、f16です。(f16はデフォルトでは使用不可)

//WGSL
override a: bool = true;
override baseColor: u32 = 0xff00ff00;
@id(0) dot: f32 = 0.01;
@id(100) po: f32 = -100;
override cf = po * 0.1;//型推論
override suf: bool;

//JavaScript
  let computePipeline = this.device.createComputePipeline({
    layout : "auto",
      compute : {
        module : this.computeShaderModule,
        entryPoint : entryPoint,
        constants: {
          a: true,
          0: 0,
          100: 9.9,
          cf: 0.01,
          suf: false
        }
      }
      })

JavaScript側から上書きするにはパイプラインを作成する際にconstantsプロパティに値を入れておきます。@idで指定されいる場合はその数字、@idがない場合は変数名で入れます。constantsプロパティで値を指定しない場合はデフォルトの定数値が使用されます。(デフォルトの数値がなく、JavaScriptから上書きされることもなければエラーになる)

むすび

英語で書かれたよくわからんチョコがおいしいです。疲れたら甘いものを是非に。((謎宣伝

twitter