WebGPUでComputeShader

# WebGPUでComputeShader
イングリッシュな記事が多くて困ったりしたので、少し纏めることにしました。間違いがありましたらご指摘お願します。

:::note warn
**WebGPUの機能**は今後さらに改善されていく可能性があるので、**最新情報のチェック**を忘れずに。
:::

## 参考文献
先に**参考文献**を置いておきます。
<details><summary>WebGPU(ComputeShader)</summary>

・[WebGPU - コアの全てを canvas 抜きで](https://inzkyk.xyz/misc/webgpu/)
・[コンピュートシェーダーを使う](https://zenn.dev/emadurandal/books/cb6818fd3a1b2e/viewer/compute_only)
</details>

<details><summary>WebGPU(こちらは英語です。)</summary>

・[WebGPU API](https://developer.mozilla.org/en-US/docs/Web/API/WebGPU_API#browser_compatibility)
・[WebGPU Shading Language](https://gpuweb.github.io/gpuweb/wgsl/)
・[WebGPU WGSL](https://webgpufundamentals.org/webgpu/lessons/webgpu-wgsl.html)
・[WebGPU Error Handling best practices](https://toji.dev/webgpu-best-practices/error-handling)
</details>

<details><summary>WebGPUのサンプルコードはこちらが参考になります。</summary>

・[WebGPU Samples](https://webgpu.github.io/webgpu-samples/)
</details>

<details><summary>GPGPU初学者の方にはc++AMPが分かり易くておすすめです。が、c++AMPはこれから廃れていく一方なので注意が必要です。</summary>

・[C++ AMPでGPGPU！](https://qiita.com/wx257osn2/items/6ad749f194db12c98467)
・[C++ AMP の概要](https://learn.microsoft.com/ja-jp/cpp/parallel/amp/cpp-amp-overview?view=msvc-170)
・[C++ AMPを用いたGPUプログラミングwiki](https://aeroastro.sd.tmu.ac.jp/hydrodynamics/main/colums/AMPindex.html)
・[C++ AMPで画像処理](https://any-programming.hatenablog.com/entry/2018/01/01/174840)
・[C++ AMPの死について](https://zenn.dev/wx257osn2/articles/rip_cxx_amp-ksaudhawigubweinfwklaeiuhfawelifh)

</details>

<details><summary>windowsでしたらhlslでもいいかもしれないです。情報が豊富なので学習しやすいです。</summary>

・[上位レベル シェーダー言語 (HLSL)](https://learn.microsoft.com/ja-jp/windows/win32/direct3dhlsl/dx-graphics-hlsl)
・[Direct3D 11.0　コンピュートシェーダー](http://maverickproj.web.fc2.com/d3d11_17.html)
・[DirectXの話 第108回](https://sites.google.com/site/monshonosuana/directx%E3%81%AE%E8%A9%B1/directx%E3%81%AE%E8%A9%B1-%E7%AC%AC108%E5%9B%9E)
・[シェーダを使った画面クリア](https://tositeru.github.io/ImasaraDX11/part/clear-screen)
・[方法: コンピューティング シェーダーを作成する](https://learn.microsoft.com/ja-jp/windows/win32/direct3d11/direct3d-11-advanced-stages-compute-create)
・[Direct3D 11でのバッファリソースタイプについて](https://qiita.com/kodai100/items/89b0e08ab99b5f6a7cac)
・[Direct3D11/DirectX11 (6) D3D11 の ComputeShader を使ってみる](https://wlog.flatlib.jp/2008/11/11/n1220/)
・[ShaderでのDirectX Shaderリソースビュー](http://ja.voidcc.com/question/p-rzupqvay-br.html)
・[DirectX 11 Compute Shaders](https://www.codeproject.com/articles/42612/directx-compute-shaders)
・[Direct3D11/DirectX11 ComputeShader 4.0 を使う](https://wlog.flatlib.jp/2009/06/12/n1342/)
・[[Unity][ComputeShader]GroupMemoryBarrierWithGroupSyncとgroupsharedについて](https://qiita.com/keito_takaishi/items/3fb6c527fa46974027cf)
・[Variable Syntax](https://learn.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-variable-syntax)
・[HLSLリファレンスまとめ](https://hlslref.wiki.fc2.com/)
・[DirectX11の自作シェーダーの定数](https://dxlib.xsrv.jp/cgi/patiobbs/patio.cgi?mode=view&no=4032)
・[リソース制限 (Direct3D 11)](https://learn.microsoft.com/ja-jp/windows/win32/direct3d11/overviews-direct3d-11-resources-limits)
・[定数バッファ](https://tositeru.github.io/ImasaraDX11/part/constant-buffer)
</details>

<details><summary>GPUを使用したアルゴリズムなどです。GPUの理解などの助けになります。</summary>

・[GPGPUでSeamCarvingをしてみる](https://qiita.com/thammin/items/cc8319dacd4a4fb6dd48)
・[Compute Shaderでのソートアルゴリズムの処理時間比較](https://qiita.com/Kuyuri-Iroha/items/057ddd61556190d021fa)
・[GPU最速ソート! Radix Sort その①](https://qiita.com/tommyecguitar/items/3c1897bceda4a06beef2)
・[DirectX11 GPUパーティクル(ComputeShader)の実装](https://qiita.com/Kanade00/items/939a3d47bf16dd1a6d63)
</details>

　
以下からは**サンプルコード**です。正確性は保証しませんが、学習のお役に立つと思ます。
<details><summary>WebGPU(main.js)</summary>

```js
/*

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()
```
</details>

<details><summary>WebGPU(webgpuControl.js)</summary>

```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
  }
}
```
</details>

<details><summary>c++AMP</summary>

```c++
//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;
}
```
</details>

<details><summary>hlsl(ComputeShader_M.hpp)</summary>

```c++
/*

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();
	}

};
```
</details>

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

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

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

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

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

具体的な得意分野を挙げると…
```
CPU
・if文やfor文の多い処理。
・分岐がいっぱいの複雑な処理。

GPU
・if文やfor文の少ない簡単な処理。
・並列に実行できる処理。
```
:::note info
**GPU**には**並列に実行できる簡単な処理**を任せると良きです。
:::

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

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

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

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

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

以下はイメージです。

```plantuml
@startuml
title CPU -> GPU
participant CPU as cpu
participant WebGPU as wgsl
participant GPU as gpu
cpu -> wgsl: WebGPUに配列を渡す
wgsl -> gpu: CPUにある値をGPUのメモリにコピー
@enduml
```

```plantuml
@startuml
title GPU -> CPU
participant CPU as cpu
participant WebGPU as wgsl
participant GPU as gpu
cpu -> wgsl: gpuにある値を知りたい
gpu -> cpu: WebGPUがgpuにあるメモリをcpuにコピーする
@enduml
```
:::note info
**CPU**と**GPU**のデータのやりとりには**WebGPU**を使います。
:::

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

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

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

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

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

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

```plantuml
@startuml
participant CPU as cpu
participant WebGPU as wgsl
participant GPU as gpu
cpu -> cpu: **1**. 初期化。\n**2**. **WGSLコード**をコンパイルする。\n**3**. **GPU**の計算に必要なバッファを用意する。
cpu -> wgsl: **4**. 2と3を引数に**シェーダーを実行**する。
wgsl -> gpu: 4の情報をもとにWebGPUがメモリのコピーなどを行う。
gpu <- gpu: シェーダーが並列に実行される。
cpu -> wgsl: GPUにある値を知りたい。
gpu -> cpu: **5**. 計算結果が**GPU**のメモリにあるので、それを**CPUのメモリにコピー**する。
@enduml
```
:::note warn
この項では一般的な処理の流れを紹介しています。実際には3や5の処理を飛ばしても問題はありません。(実用性もありませんが笑)
:::

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

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

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

  //論理デバイスの取得
  device = await adapter.requestDevice()
```
　
まずは、**ブラウザがWebGPUに対応しているかどうか**を調べます。グローバル変数である`navigator.gpu`を確認することで調べることができます。
```js
  if (!navigator.gpu) {
    //ブラウザがGPUのサポートをしていないときの処理
    //...
  }
```
　
次に全体のベースとなるオブジェクトを取得します。
**adapterは物理デバイス**を表し、そこから`requestDevice関数`で**deviceという論理デバイス**を取得します。**返り値がpromise**であることに注意してください。

また、**物理デバイスがnull**ならデバイスに**WebGPUに対応したGPUがない**ということになります。
```js
  //物理デバイスの取得
  const adapter = await navigator.gpu.requestAdapter()
  if (!adapter) {
    //WebGPUに対応したGPUがない場合の処理
    //...
  }

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

### WGSLコードのコンパイル
**WGSL**(GPUで実行するためのプログラム)を**コンパイル**します。**WGSL**の書き方や文法については後述します。
```js
  //WGSLコード
  const shader_code = ``

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

### バッファの準備
**GPU**で使いたいメモリを`createBuffer関数`で作成します。**CPU上のメモリ**を**GPU上のメモリ**に**コピー**したい場合は`writeBuffer関数`でコピーします。**GPUBuffer**(WebGPUで扱うためのバッファ)は使いたいバッファの分だけ用意する必要があります。
```js
  //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関数`の**フラグ**には以下のようなものがあります。
```js
//コピー操作の宛先となる場合(writeBuffer使用時)
GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST

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

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

//GPU内で使用できるバッファ(バッファの値は未定義)
GPUBufferUsage.STORAGE
```
　
`writeBuffer関数`では引数でコピー位置などを指定することができます。
```js
device.queue.writeBuffer(
  inputBuffer,//コピー先のGPUBuffer
  0,//GPUBufferのどこからコピーを始めるか(バイト単位)
  inputArray,//コピー元のバッファ(ArrayBuffer,TypedArray, DataViewのいずれか)
  0,//コピーの開始位置(※)
  inputArray.length//コピーするサイズ(※)
)
```
:::note warn
※**第4引数**と**第5引数**の値は**コピー元のバッファの種類によって指定方法が変わります**。コピー元のバッファが**TypedArrayの場合は要素の数**、**TypedArray以外の場合にはバイト単位**で指定します。また、第4引数と第5引数は省略可能です。(省略時の既定値は、第4引数は0、第5引数は第4引数で指定された値からバッファの最後までのサイズ)
:::

### シェーダーの実行
シェーダーを実行します。以下は例です、
```js
  //パイプラインの作成
  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()])
```
　
まずは**コンピュートシェーダー用のパイプライン**を作成します。
```js
  //パイプラインの作成
  const computePipeline = device.createComputePipeline({
    layout: "auto",
    compute: {
      module: computeShaderModule,
        //createShaderModuleで作成したものです
      entryPoint: "main"
        //WGSLコード内で実行する関数名です
    }
  })
```
`layout`にはパイプラインで使用される**GPUリソース(GPUBufferなど)の構造情報**を設定します。
構造情報を設定する方法は、
**1. createPipelineLayout関数で構造情報を作成する。
2. "auto"文字列を設定することでWGSLコードから自動的に構造情報を作成する。**
…の二種類があります。

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

詳細は以下を確認してください。
[createPipelineLayout](https://developer.mozilla.org/en-US/docs/Web/API/GPUDevice/createPipelineLayout)
[createBindGroupLayout](https://developer.mozilla.org/en-US/docs/Web/API/GPUDevice/createBindGroupLayout)
　
　
次は**BindGroup**(GPUリソースのリスト)を作成していきます。
```js
  //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を渡します。**

:::note warn
**bindingの数字はWGSLのタグと対応**するようにします。
:::
　
最後に**コマンド**を書き込んでいきます。
```js
  //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のシェーダー**は**非同期**に実行されます。
```plantuml
@startuml
participant CPU as cpu
participant WebGPU as wgsl
participant GPU as gpu
cpu -> wgsl: WebGPUにコマンドを発行
wgsl -> gpu: コマンドをもとにGPUで処理
note over gpu: シェーダーの実行
/ note over cpu: コマンドを発行した後、\nシェーダーの終了を待たずに次の処理を始める
@enduml
```
そこで、**submit関数を使用してシェーダーの終了まで待つ**ことができます。
```plantuml
@startuml
participant CPU as cpu
participant WebGPU as wgsl
participant GPU as gpu
cpu -> wgsl: WebGPUにコマンドを発行
wgsl -> gpu: コマンドをもとにGPUで処理
note over gpu: シェーダーの実行
/ note over cpu: コマンドを発行した後、\nシェーダーの終了を待つ
gpu -> wgsl: シェーダー終了
wgsl -> cpu: シェーダーの終了を伝える
note over cpu: シェーダーが終了したので\n次の処理を始める
@enduml
```

### GPUのメモリにあるバッファを読み取る
**GPU**にあるメモリを読み取るには、**GPUとCPUどちらもアクセスすることができるステージングバッファ**を介します。
```js
  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の参照を取得**することができます。用事が済んだら**ステージングバッファのマップを解除**しましょう。
```plantuml
@startuml
participant CPU as cpu
participant WebGPU as wgsl
participant stagingBuffer as sta
participant GPU as gpu
cpu -> wgsl: コマンドを発行
gpu -> sta: WebGPUがコピー処理を行う
sta -> cpu: マップすることでcpuから参照できる
sta -> cpu: マップ解除後は参照できない
@enduml
```

**ステージングバッファ**のデータは通常の配列のように**書き込みできない**ため、そのような操作をしたい場合には**CPUのバッファにコピー**することで**WebGPU**から完全に独立したバッファを作成できます。
```js
  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)側のコード**について解説していきます。[公式リファレンス](https://gpuweb.github.io/gpuweb/wgsl/)もありますが少し長いので簡単にまとめています。

### 変数の型
変数の型にはいくつか種類があります。
| keyword | type |
|:-------:|:----:|
| i32  | 32ビット符号付き整数 |
| u32  | 32ビット符号なし整数 |
| f32  | 32ビット浮動小数点数 |
| bool | ブール値            |

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

### 構文
#### 0. コメント
多くの方にとっては馴染み深い方法だと思います。
```js
//この行はコメント

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

#### 1. 変数宣言

##### データ型
`let`は定数値を宣言するときに使います。`var`は普通の変数、`const`はコンパイル時定数です。
```js
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つの同じ型の変数をまとめたベクトル型というものを使用できます。
```js
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`しか使用できません。
```js
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));
```

デフォルトで行列用の関数が用意されていますが、数が少ないので外部ライブラリの使用をおすすめします。また、簡単な関数であれば自分で実装するのも一つの選択肢です。
[コンストラクタ、距離](https://www.w3.org/TR/WGSL/#determinant-builtin)
[転値](https://www.w3.org/TR/WGSL/#transpose-builtin)
[行列ライブラリ](https://github.com/greggman/wgpu-matrix)

##### 配列
型と要素数の指定で宣言できます。動的配列は今のところありません。(ストレージバッファは別です)
```js
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);//型推論
```

##### 構造体
構造体はメンバ変数と型を対応させて書いていきます。最後のコンマはあってもなくても変わりません。
```js
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バイトになります。
```js
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<f32> | 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になります。
```js
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以上でアライメントされる必要があります。
```js
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;
```

各データ型のアライメントサイズについては[こちら](https://gpuweb.github.io/gpuweb/wgsl/#memory-layouts)にあるアライメント表を参考にしてください。

#### 2. 四則演算
変数だけでなく、ベクトル型も四則演算することができます。累乗は`＾`ではなく`pow関数`であることに注意してください。
```js
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)
```
　
異なる型同士(行列とベクトル型など)では四則記号の意味合いが変わる場合があります。詳細は[こちら](https://gpuweb.github.io/gpuweb/wgsl/#arithmetic-expr)を確認してください。

#### 3. 代入
WGSLは型定義が厳密です。違う型同士で代入を行うとエラーになります。違う型の変数を代入したいときは`f32関数`などのコンストラクタを使用して型変換します。
```js
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
```

コンストラクタ一覧は[こちら](https://gpuweb.github.io/gpuweb/wgsl/#value-constructor-builtin-function)を参照してください。

#### 4. 条件分岐
##### if文
```js
  let a = 10;
  if(a == 1) {
    //...
  }
  else if(a == 2) {
    //...
  }
  else {
    //...
  }

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

##### switch
`case`は`u32`または`i32`の定数である必要があります。
```js
let x : i32 = 10;
switch x {
  case 0: {//caseは定数の整数である必要がある
    //...
  }
  default {//defaultの順番はどこでもよい
    //...
  }
  case 1, 2, {//複数指定も可
    //...
  }
  case 3, {//カンマはあってもなくてもよい
    //...
  }
  case 4 {
    //...
  }
}
```
```js
const c = 3;
let x : i32 = 10;
switch x {
  case c: {//constで宣言した変数は使用可
    //...
  }
  case 0, default, {//変数とdefaultでも組める
    //...
  }
}
```

##### all
引数のベクトル型がすべて`true`であるときに`true`を返します。
```js
let a = all(vec2<bool>(true, true));//true
```

##### any
引数のベクトル型のいずれかの要素が`true`であるときに`true`を返します。
```js
let a = any(vec2<bool>(true, false));//true
```

##### select
三項演算子のように処理されます。`t`と`f`は通常のデータ型かベクトルです。`flag`は`true`か`false`である必要があります。
```js
let a = flag ? t : f;//JavaScript
let a = select(f, t, flag);//WGSL
```

#### 5. 反復処理

##### break
反復処理と`switch文`でのみ使用できます。
```js
  var i = 0;
  loop {
    i = i + 1;
    if i >= 10 { break; }
  }
```

##### continue
反復処理でのみ使用できます。`continue`した後は`continuing`内の処理を行ってから次のループに入ります。
```js
  var i = 0;
  loop {
    i = i + 1;
    if i >= 5 { continue; }

    continuing {//continuingはなくてもよい
      //...
      break if i >= 10;//※
    }
  }
```
:::note warn
※`continuing`はあってもなくても大丈夫ですが、`continuing`内で`break`する場合は最後に`break if`を使用します。
:::

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

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

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

#### 6. 論理演算子
`if文`などで使用すると思います。
```js
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. 関数
自分でオリジナルの関数を宣言することができます。**引数と戻り値の型を決める**必要があります。
```js
fn a(b: f32) -> i32 {//戻り値の型を`->`の後に指定する
  return i32(b * 2);
}

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

//複数の引数も可
fn d(b: f32, e: vec2<f32>) { return; }
```
また、三角関数などのいくつかの関数はデフォルトで用意されています。
関数の一覧は[リファレンス](https://gpuweb.github.io/gpuweb/wgsl/#numeric-builtin-functions)で確認できます。

#### 8. ConputeShaderで実行される関数
JavaScript(cpu)側のコードでエントリポイントとして指定される関数の記述は、普通の関数と異なります。
```js
  @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個のスレッドが走ることになります。
```js
  //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`で必要な値を引数として参照できます。
```js
  //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個のスレッドが動きます。
```js
  //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`に関しては図で確認すると分かり易いです。
![du.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/3005402/1159a59d-47fa-ba23-4d37-e135d7663e15.png)

#### 9. bindされた変数
JavaScript(CPU)側のコードでbindされた変数は、bindの番号とgrupeの番号から参照することができます。この際、アドレス空間と型を指定します。アクセス修飾子は任意で指定します。また、配列の場合は要素数を明示しないことで、JavaScript(CPU)側のコードから配列の長さを決めることができます。
```js
//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する際のフラグを参考にすることが多いです。
```js
//コピー操作の宛先となる場合(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. ビット操作
```js
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`です。参照型とポインタ型のどちらもポインタを扱う型ですが、参照型はポインタ型と比べて制約の多いものとなります。
```js
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変数は参照型にできない。
・関数の引数にすることができない。
・宣言時に初期化が必要。
・再代入不可。

構造体などのメンバ変数を参照するには以下のようにします。
```js
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
}
```
　
ポインタ型は関数の引数にすることができます。
```js
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
}
```
　
特定のポインタにアクセスするようなコードはコンパイルエラーになります。
```js
let b: ref<private, i32> = 0;//error : 0番地にアクセスしようとしている
```

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

例えば、スレッド1とスレッド2が変数aにアクセスする際、処理の順番によって結界が変わります。これはスレッド1と2がそれぞれ非同期に実行されるためです。
```plantuml
@startuml
title 順番に実行された場合
participant スレッド1 as one
participant メモリ as me
participant スレッド2 as two
note over me: a==1
one <- me: aの値を読み取る\na==1
one -> one: a=a+1\na==2
one -> me: メモリに書き込む\na==2
note over me: a==2
two <- me: aの値を読み取る\na==2
two <- two: a=a+1\na==3
two -> me: メモリに書き込む\na==3
note over me: a==3
@enduml
```
```plantuml
@startuml
title ほぼ同じタイミングで実行された場合
participant スレッド1 as one
participant メモリ as me
participant スレッド2 as two
note over me: a==1
one <- me: aの値を読み取る\na==1
two <- me: aの値を読み取る\na==1
one -> one: a=a+1\na==2
two <- two: a=a+1\na==2
one -> me: メモリに書き込む\na==2
two -> me: メモリに書き込む\na==2
note over me: a==2
@enduml
```
これでは結果が2になるときと3になるときが発生してしまいます。
そこでatomic操作をすることで、片方のスレッドが変数aにアクセスしている間は他のスレッドが変数aにアクセスできないようにすることができます。
```plantuml
@startuml
title ほぼ同じタイミングで実行された場合
participant スレッド1 as one
participant メモリ as me
participant スレッド2 as two
note over me: a==1
one <- me: aの値を読み取る\n(atomic操作開始)\na==1
two <- me: スレッド1が変数aに\nアクセス中はアクセスできない
one -> one: a=a+1\na==2
two <- me: スレッド1が変数aに\nアクセス中はアクセスできない
one -> me: メモリに書き込む\n(atomic操作終了)\na==2
note over me: a==2
two <- me: aの値を読み取る\na==2
two <- two: a=a+1\na==3
two -> me: メモリに書き込む\na==3
note over me: a==3
@enduml
```
　
atomic操作はatomic型でのみ使用可能です。変数の型は`u32`か`i32`である必要があり、アドレス空間は`workgroup`か`storage`でアクセス修飾子は`read_write`である必要があります。
```js
@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型の値と第二引数の値とでそれぞれの演算を行い、結果をポインタ先に入れます。戻り値はもともとポインタ先にあった値です。
　
一部の実装では[等価比較が上手くいかない](https://www.w3.org/TR/WGSL/#atomic-rmw)場合がありますが、以下のような関数もあります。
```js
  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は基本的に非同期に実行されますが、アルゴリズムによっては処理の足並みを揃えたいときがあります。
```plantuml
@startuml
title 非同期処理(実行順はばらばら)
participant スレッド1 as one
participant スレッド2 as two
participant スレッド3 as th
participant スレッド4 as f
one -> one: スレッド1の実行\n(順不動)
th -> th: スレッド3の実行\n(順不動)
f -> f: スレッド4の実行\n(順不動)
two -> two: スレッド2の実行\n(順不動)
@enduml
```
```plantuml
@startuml
title 同期処理(同期関数のとき同期する)
participant スレッド1 as one
participant スレッド2 as two
participant スレッド3 as th
participant スレッド4 as f
one -> one: 同期関数まで\nスレッド1の実行\n(順不動)
th -> th: 同期関数まで\nスレッド3の実行\n(順不動)
f -> f: 同期関数まで\nスレッド4の実行\n(順不動)
two -> two: 同期関数まで\nスレッド2の実行\n(順不動)
note over one: 同期関数の実行
/ note over th: 同期関数の実行
/ note over f: 同期関数の実行
/ note over two: 同期関数の実行
one -> one: 同期関数の\n直後の処理から\nスレッド1の実行\n(順不動)
th -> th: 同期関数の\n直後の処理から\nスレッド3の実行\n(順不動)
f -> f: 同期関数の\n直後の処理から\nスレッド4の実行\n(順不動)
two -> two: 同期関数の\n直後の処理から\nスレッド2の実行\n(順不動)
@enduml
```
　
同期関数には三種類あります。どの関数も同じworkgroup内のスレッドが同期します。(異なるworkgroup同士での同期は今のところできません。)
```js
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だと引数と戻り値の値は同じになる)
}
```
　
以下のような処理はコンパイルエラーになります。同期関数が実行されないスレッドがでてくるからです。
```js
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内で使用できるメモリはそこまで大きくありません。そのため、できるだけ少ないメモリで多くの情報量を持てると嬉しいです。そこで複数の変数を一つ分のメモリを持った変数にパッケージ化することができます。(もとの型に比べて精度は落ちますが、そこまで精度が大事でない場合はパッケージ化しても影響が少ない場合が多いです。)

```js
  //-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はエラーを出します。プログラムを書くときはこれらの制限以内の値になるように気を付けなければならないです。制限の一覧は[こちら](https://gpuweb.github.io/gpuweb/#limits)を確認してください。

### 制限を引き上げる
ブラウザ以外にもGPUを使うソフトが存在する場合があります。
```plantuml
@startuml
title GPUを使うソフト
participant ブラウザ as bro
participant GPU as gpu
participant ソフト1 as game1
participant ソフト2 as game2
participant ソフト3 as game3
bro -> gpu: WebGPUからGPUを使う
game1 -> gpu: OpenGLからGPUを使う
game2 -> gpu: OSのAPIからGPUを使う
note over bro: GPUを使用します
/ note over game1: GPUを使用します
/ note over game2: GPUを使用します
/ note over game3: GPUを使用しません
@enduml
```

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

現在のlimitの制限をGPUの制限ギリギリまで引き上げたい際などは`requestDevice関数`の引数にlimitオブジェクトを渡すことで、制限された値を上げることができます。
```js
  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を使います。
```js
    //制限を引き上げた論理デバイスの取得
    device = await adapter.requestDevice({requiredLimits: adapter.limits})
```

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

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

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

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

[こちら](https://webgpu.github.io/webgpu-samples/)のサンプル集では3つのシェーダーを組み合わせたコードも見ることができます。


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

#### 1.label
WebGPUの多くの関数の引数には`label`プロパティがあり、この`label`に文字列を設定するとエラー時のコンソールに`label`名が表示されます。
```js
  device = await adapter.requestDevice({label : "test1"})
    //関数が失敗したとき、エラー文と"test1"という文字列が表示される
```

複数のシェーダーを実行した際に、どの`label`でエラーが起きたのかを確認することができるので便利です。
```js
  device1 = await adapter1.requestDevice()//1
  device2 = await adapter2.requestDevice()//2
    //エラーが出ても、どっちでエラーが出ているか分からない
```
```js
  device1 = await adapter1.requestDevice({label : "test1"})//1
  device2 = await adapter2.requestDevice({label : "test2"})//2
    //エラーが出るとラベル名が表示されるので、どこでエラーが発生したか明確
```

また、同じ関数を繰り返し使用する場合にも有効です。
```js
  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`を使用するといい場合があります。
```js
  //コマンドエンコーダーの作成
  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`しましょう。
```js
  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が識別できなかった場合に出るエラー。[制限](https://gpuweb.github.io/gpuweb/#gpusupportedlimits)に違反したことを示す。(現在はパイプラインの作成によってのみ発生)

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

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

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

すべてのエラーを補足したい場合は`push`と`pop`で囲ってしまいます。
```js
  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)
```

実際の実装では以下のようにすることで処理が失敗した時のプログラムを記述します。
```js
  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を使う際は非同期に実行される他の処理に気を付ける必要があります。
```js
//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)にいってしまう
```
```js
//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のスタックに文字列がない場合に機能します。一括でエラーを補足したい場合に便利です。
```js
  //deviceにイベントを登録
  device.addEventListener("uncapturederror", (event) => {
    console.log(event.error.message)
  })
```
また、物理デバイス毎によるバグの頻度の検知などする際は[こちら](https://developer.mozilla.org/en-US/docs/Web/API/GPUAdapter/requestAdapterInfo)と組み合わせることも有効なようです。

#### 5.Shaders
`getCompilationInfo関数`はWGSL内のエラーを補足します。(3と4の方法ではWGSLでのエラーを補足できません) 
```js
  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"`の三種類があります。

プロパティについての詳細は[こちら](https://developer.mozilla.org/en-US/docs/Web/API/GPUCompilationMessage)を確認してください。

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

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

<details><summary>main.html</summary>

```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>
```
</details>

<details><summary>main.js</summary>

```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()
```
</details>

#### 2.GPUTexture
[GPUTexture](https://developer.mozilla.org/en-US/docs/Web/API/GPUTexture)を使う方法もあります。`createImageBitmap関数`で作成したビットマップデータを`craeteTexture関数`で作成した`GPUTexture`に`copyExternalImageToTexture関数`でコピーします。

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

WGSLで`GPUTexture`を宣言するには[こちら](https://www.w3.org/TR/WGSL/#sampled-texture-type)にある型を使用します。

参照するには[Texture Built-in Functions](https://www.w3.org/TR/WGSL/#texture-builtin-functions)を使用します。[sampler](https://www.w3.org/TR/WGSL/#sampled-texture-type)を使うことで行列などを用いた座標変換などを行うこともできるようです。

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

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

### 定数の上書き
WGSLコード内での定数宣言の際、`override`を付けることでJavaScript側から上書きできるようになります。対応している型は`bool`、`i32`、`u32`、`f32`、`f16`です。(`f16`はデフォルトでは使用不可)
```js
//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;
```
```js
//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](https://twitter.com/metaphysicalba1)