WebGPUでComputeShader
イングリッシュな記事が多くて困ったりしたので、少し纏めることにしました。間違いがありましたらご指摘お願します。
WebGPUの機能は今後さらに改善されていく可能性があるので、最新情報のチェックを忘れずに。
参考文献
先に参考文献を置いておきます。
WebGPU(ComputeShader)
WebGPU(こちらは英語です。)
・WebGPU API
・WebGPU Shading Language
・WebGPU WGSL
・WebGPU Error Handling best practices
WebGPUのサンプルコードはこちらが参考になります。
GPGPU初学者の方にはc++AMPが分かり易くておすすめです。が、c++AMPはこれから廃れていく一方なので注意が必要です。
・C++ AMPでGPGPU!
・C++ AMP の概要
・C++ AMPを用いたGPUプログラミングwiki
・C++ AMPで画像処理
・C++ AMPの死について
windowsでしたらhlslでもいいかもしれないです。情報が豊富なので学習しやすいです。
・上位レベル シェーダー言語 (HLSL)
・Direct3D 11.0 コンピュートシェーダー
・DirectXの話 第108回
・シェーダを使った画面クリア
・方法: コンピューティング シェーダーを作成する
・Direct3D 11でのバッファリソースタイプについて
・Direct3D11/DirectX11 (6) D3D11 の ComputeShader を使ってみる
・ShaderでのDirectX Shaderリソースビュー
・DirectX 11 Compute Shaders
・Direct3D11/DirectX11 ComputeShader 4.0 を使う
・[Unity][ComputeShader]GroupMemoryBarrierWithGroupSyncとgroupsharedについて
・Variable Syntax
・HLSLリファレンスまとめ
・DirectX11の自作シェーダーの定数
・リソース制限 (Direct3D 11)
・定数バッファ
GPUを使用したアルゴリズムなどです。GPUの理解などの助けになります。
・GPGPUでSeamCarvingをしてみる
・Compute Shaderでのソートアルゴリズムの処理時間比較
・GPU最速ソート! Radix Sort その①
・DirectX11 GPUパーティクル(ComputeShader)の実装
以下からはサンプルコードです。正確性は保証しませんが、学習のお役に立つと思ます。
WebGPU(main.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()
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から上書きされることもなければエラーになる)
むすび
英語で書かれたよくわからんチョコがおいしいです。疲れたら甘いものを是非に。((謎宣伝