コンピュートシェーダー
コンピュートシェーダーは、レンダリングパイプラインとは独立して、GPU上で汎用計算を実行するプログラムです。頂点シェーダーやフラグメントシェーダーとは異なり、コンピュートシェーダーはジオメトリやピクセルに縛られず、任意のデータを操作するため、パーティクルシミュレーション、画像処理、物理計算、プロシージャルコンテンツ生成などのタスクに理想的です。
コンピュートシェーダーはWebGPUプラットフォームでのみサポートされています。WebGLを使用している場合は利用できません。
サポートの確認
コンピュートシェーダーを使用する前に、デバイスがサポートしているか確認してください:
if (device.supportsCompute) {
// コンピュートシェーダーが利用可能
}
WGSL 言語拡張
ブラウザが任意の WGSL 機能(例: 線形ワーカー / 呼び出しインデックス、サブグループ、半精度 float)を公開している場合、エンジンは対応する device.supports* フラグと CAPS_* プリプロセッサ定義を設定します。一覧と注意点は WGSL 言語拡張 を参照してください。
コンピュートシェーダーの作成
コンピュートシェーダーは、WGSL の cshader ソースを使用して Shader クラスで作成します。使用するリソース(ユニフォーム、ストレージバッファ、テクスチャ、ストレージテクスチャ)を簡略化された WGSL 構文(@group/@binding なし)で宣言すると、エンジンがソースからそれらを反映し、バインドグループを自動的に構築します。リソース宣言構文の詳細は WGSL リフレクション を参照してください。
基本的なシェーダー定義
const shader = new pc.Shader(device, {
name: 'MyComputeShader',
shaderLanguage: pc.SHADERLANGUAGE_WGSL,
cshader: `
// 簡略化された構文で宣言されたリソースは自動的に反映されます -
// computeBindGroupFormat は不要です
uniform count: u32;
var<storage, read_write> data: array<f32>;
@compute @workgroup_size(64, 1, 1)
fn main(@builtin(global_invocation_id) global_id: vec3u) {
let i = global_id.x;
if (i >= uniform.count) { return; }
data[i] = data[i] * 2.0;
}
`
});
バインドグループを明示的に制御する必要がある場合(手動で記述する computeBindGroupFormat、明示的な @group/@binding インデックス、または手動バインドのリソースと反映されるリソースの混在)は、コンピュートシェーダー(応用) を参照してください。
エントリポイント
デフォルトでは、エンジンはエントリポイント関数の名前がmainであることを期待します。computeEntryPointを使用して別の関数名を指定することもでき、これにより単一のシェーダーソースに複数のエントリポイントを含めることもできます:
const shaderSource = `
@compute @workgroup_size(64, 1, 1)
fn initParticles(@builtin(global_invocation_id) global_id: vec3u) {
// パーティクルの初期化
}
@compute @workgroup_size(64, 1, 1)
fn updateParticles(@builtin(global_invocation_id) global_id: vec3u) {
// パーティクルの更新
}
`;
// 同じソースから、それぞれ異なるエントリポイントを使用して別々のシェーダーを作成
const initShader = new pc.Shader(device, {
name: 'InitParticles',
shaderLanguage: pc.SHADERLANGUAGE_WGSL,
cshader: shaderSource,
computeEntryPoint: 'initParticles'
});
const updateShader = new pc.Shader(device, {
name: 'UpdateParticles',
shaderLanguage: pc.SHADERLANGUAGE_WGSL,
cshader: shaderSource,
computeEntryPoint: 'updateParticles'
});
コンピュートインスタンスの作成
Computeクラスは、関連するパラメータを持つコンピュートシェーダーの実行可能なインスタンスを表します:
const compute = new pc.Compute(device, shader, 'MyComputeInstance');
パラメータの設定
setParameterを使用してリソースをバインドし、ユニフォーム値を設定します。リソースは名前によってシェーダーの宣言に対応付けられます:
// ストレージバッファをバインド
compute.setParameter('data', storageBuffer);
// テクスチャをバインド
compute.setParameter('inputTexture', texture);
// ユニフォーム値を設定
compute.setParameter('count', 1024);
compute.setParameter('tint', [1.0, 0.5, 0.0, 1.0]);
コンピュートシェーダーが宣言するすべてのユニフォームとリソースには、コンピュートインスタンスで値を設定する必要があります。コンピュートシェーダーはグラフィックスデバイスのスコープにグローバルに設定された値を使用しません。また、デバッグビルドでは、値が設定されていない宣言済みのユニフォームやリソースが報告されます。
ストレージバッファの作成
ストレージバッファは、コンピュートシェーダーが読み書きできるデータを保持します:
const storageBuffer = new pc.StorageBuffer(
device,
bufferSizeInBytes,
pc.BUFFERUSAGE_COPY_SRC | // CPUへの読み戻しを有効化
pc.BUFFERUSAGE_COPY_DST // CPUからの書き込みを有効化
);
// 初期データを書き込み
const data = new Float32Array([...]);
storageBuffer.write(0, data);
// バッファをクリア
storageBuffer.clear();
ストレージテクスチャの作成
ストレージテクスチャはstorage: trueオプションで作成されます:
const storageTexture = new pc.Texture(device, {
name: 'StorageTexture',
width: 512,
height: 512,
format: pc.PIXELFORMAT_RGBA8,
mipmaps: false,
minFilter: pc.FILTER_LINEAR,
magFilter: pc.FILTER_LINEAR,
storage: true // ストレージテクスチャとして有効化
});
コンピュートシェーダーのディスパッチ
コンピュートシェーダーを実行するには、まずディスパッチ次元を設定してからディスパッチします:
// ディスパッチ次元を設定(X、Y、Zのワークグループ数)
compute.setupDispatch(width, height, 1);
// コンピュートシェーダーをディスパッチ
device.computeDispatch([compute], 'MyDispatch');
複数のコンピュートシェーダーを単一のコンピュートパスでまとめてディスパッチできます:
compute1.setupDispatch(64, 64);
compute2.setupDispatch(128, 128);
device.computeDispatch([compute1, compute2], 'BatchedDispatch');
device.computeDispatch は現在のフレームのコマンドエンコーダに記録されるため、レンダリングフレーム内(通常は app.on('update', ...) ハンドラ内)で呼び出す必要があります。フレーム外(setTimeout や切り離された Promise など)から呼び出すと信頼できず、ディスパッチが暗黙的にスキップされることがあります。
コンピュートインスタンスは1フレームに最大1回ディスパッチされます。1フレーム内でコンピュートシェーダーを複数回ディスパッチするには、ディスパッチごとに別のコンピュートインスタンスを使用してください。インスタンス間でシェーダーを共有でき、シェーダーのコンパイルは一度だけです:
const computes = [0.25, 0.5].map((strength) => {
const compute = new pc.Compute(device, shader, `Blur-${strength}`);
compute.setParameter('strength', strength);
compute.setupDispatch(width, height);
return compute;
});
device.computeDispatch(computes, 'Blur');
ワークグループサイズ
総呼び出し回数はdispatchSize × workgroupSizeです。例えば、(width, height)でディスパッチし、シェーダーが@workgroup_size(1, 1, 1)の場合、width × height回の呼び出しが行われます。
大規模なデータセットでのパフォーマンス向上のために、より大きなワークグループサイズを使用してください。例えば、1次元のデータには64、画像には8 × 8が適しています:
@compute @workgroup_size(64, 1, 1)
fn main(@builtin(global_invocation_id) global_id: vec3u) {
// global_id.xの要素を処理
}
スレッド数を指定したディスパッチ
setupDispatchThreadsは、ワークグループ数の代わりに、各次元でディスパッチするスレッド数を受け取ります。スレッドはシェーダーの1回の呼び出しで、通常は1つの要素またはピクセルを処理します。ワークグループ数は、シェーダーが宣言するワークグループサイズから求められます。そのため、ワークグループサイズは@workgroup_size(8, 8)のように整数リテラルで記述する必要があります。cdefinesで注入した値(定義とインクルードを参照)は、ワークグループサイズが読み取られる前にリテラルに置き換えられます。
// 要素ごとに1スレッド
compute.setupDispatchThreads(numElements);
// 画像のピクセルごとに1スレッド
compute.setupDispatchThreads(width, height);
ディスパッチはワークグループ単位に切り上げられるため、シェーダーは要求された数を超えるスレッドをスキップします。dispatchThreadsCSインクルードの関数は、このスキップとスレッドのインデックス計算を行います:
#include "dispatchThreadsCS"
var<storage, read_write> values: array<f32>;
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) id: vec3u) {
if (!dispatchThreadInRange(id)) {
return;
}
let index = dispatchThreadIndex(id);
values[index] = values[index] * 2.0;
}
dispatchThreadsCSインクルードは次の関数を提供します:
| 関数 | 説明 |
|---|---|
dispatchThreadCount() -> vec3u | ディスパッチに設定されたスレッド数。 |
dispatchThreadInRange(globalId: vec3u) -> bool | スレッドが、ディスパッチに設定されたスレッドの1つである場合にtrue。 |
dispatchThreadIndex(globalId: vec3u) -> u32 | スレッドの線形インデックス。xが最も速く変化します。 |
シェーダーは、書き込むテクスチャのサイズなど、独自の範囲でスレッドをスキップすることもできます。ただし、デバイスが各次元でサポートするワークグループ数には上限があり、少なくとも65535です。それを超えるワークグループが必要な1次元のディスパッチは複数行のワークグループに分割されます。その場合、シェーダーはdispatchThreadsCSをインクルードし、分割を考慮するdispatchThreadIndexでスレッドのインデックスを求める必要があります。
間接ディスパッチ
間接ディスパッチは、あるコンピュートシェーダーが別のコンピュートシェーダーのディスパッチパラメータを生成することを可能にし、CPUへの読み戻しなしで完全にGPU駆動のワークロードを実現します。これは以下の場合に便利です:
- GPUで決定される可変ワークロードサイズ
- タイル数が動的に計算されるタイルベースの処理
- GPUカリング後の可視要素のみの処理
ディスパッチスロットの予約
デバイスは間接ディスパッチパラメータ用の組み込みバッファを提供します。毎フレームスロットを予約してください:
const slot = device.getIndirectDispatchSlot();
各スロットは、x、y、zワークグループ数を表す3つの32ビット符号なし整数を保持します。最大スロット数はdevice.maxIndirectDispatchCountで制御されます(デフォルト:256)。
ディスパッチパラメータの書き込み
コンピュートシェーダー内で間接バッファを宣言し(自動的に反映されます)、予約したスロットにディスパッチパラメータを書き込みます。間接ディスパッチレイアウトに一致する構造体を定義します:
struct DispatchIndirectArgs {
x: u32,
y: u32,
z: u32
};
var<storage, read_write> indirectBuffer: array<DispatchIndirectArgs>;
uniform slot: u32; // 書き込み先のスロットインデックス
@compute @workgroup_size(1)
fn main() {
// ワークロードサイズを動的に計算
let workloadSize = calculateWorkload();
// ディスパッチパラメータをスロットに書き込み
indirectBuffer[uniform.slot].x = workloadSize;
indirectBuffer[uniform.slot].y = 1u;
indirectBuffer[uniform.slot].z = 1u;
}
間接ディスパッチの使用
setupIndirectDispatchを使用して、2番目のコンピュートシェーダーがバッファからディスパッチパラメータを読み取るように設定します:
// このフレームのスロットを予約
const slot = device.getIndirectDispatchSlot();
// 最初のパス:コンピュートシェーダーがディスパッチパラメータを書き込む
prepareCompute.setParameter('indirectBuffer', device.indirectDispatchBuffer);
prepareCompute.setParameter('slot', slot);
prepareCompute.setupDispatch(1, 1, 1);
device.computeDispatch([prepareCompute]);
// 2番目のパス:バッファからのパラメータを使用してディスパッチ
processCompute.setupIndirectDispatch(slot);
device.computeDispatch([processCompute]);
デバイスの組み込み間接バッファを使用する場合、スロットは現在のフレームでのみ有効なため、setupIndirectDispatchを毎フレーム呼び出す必要があります。
カスタム間接バッファ
レンダリングフレーム外での複雑なスケジューリングなどの高度なユースケースでは、独自のストレージバッファを提供できます:
// 間接ディスパッチ用のカスタムバッファを作成
const customBuffer = new pc.StorageBuffer(device, 3 * 4, pc.BUFFERUSAGE_INDIRECT);
// 間接ディスパッチにカスタムバッファを使用
compute.setupIndirectDispatch(0, customBuffer);
カスタムバッファを使用する場合、そのライフタイムと内容を自分で管理します。フレーム検証は行われません。
Engine 2.23以降では、デバイスの復旧時にコンピュートリソースが再作成されますが、ストレージバッファの内容は再度書き込むか再生成する必要があります。中断された読み取りへの対応や明示的なリソースの破棄については、デバイスロストと復旧を参照してください。
CPUへのデータの読み戻し
ストレージバッファからCPUに結果を読み戻すには:
const resultData = new Float32Array(numElements);
storageBuffer.read(0, undefined, resultData).then((data) => {
// データを処理
console.log('First value:', data[0]);
});
read()はGPU操作が非同期であるためPromiseを返します。データはGPUがコンピュートシェーダーの実行を完了した後に利用可能になり、これは数フレーム後になる場合があります。
時間に厳しい読み取りの場合、4番目のパラメータとしてimmediate: trueを渡すことができます:
storageBuffer.read(0, undefined, resultData, true).then((data) => {
// データがより早く利用可能になるが、パフォーマンスコストがある
});
デフォルト(immediate: false)では、読み取りはGPUコマンドバッファが自然に送信される次のイベント処理サイクルまで延期されます。immediate: trueでは、コマンドバッファが即座に送信され、読み取りがすぐに実行されます。
immediate: trueを使用すると、コマンドバッファの早期送信を強制するため、パフォーマンスに影響があります。低レイテンシの読み取りが不可欠な場合にのみ使用してください。
シーンの深度マップとカラーマップ
コンピュートシェーダーはカメラがレンダリングしたシーンの深度マップとカラーマップを読み取ることができます。たとえば、スクリーンスペースエフェクトの実装や、レンダリングされた画像の解析に使用できます。
カメラは、カメラコンポーネントでリクエストされるとこれらのマップをレンダリングします:
cameraEntity.camera.requestSceneDepthMap(true);
cameraEntity.camera.requestSceneColorMap(true);
CameraFrameでレンダリングするカメラは、そのように設定されている場合にマップをレンダリングします:
cameraFrame.rendering.sceneDepthMap = true;
cameraFrame.rendering.sceneColorMap = true;
cameraFrame.update();
カメラのマップをコンピュートインスタンスにアタッチします:
compute.setSceneDepthMap(cameraEntity.camera.sceneDepthMapHandle);
compute.setSceneColorMap(cameraEntity.camera.sceneColorMapHandle);
マップのアタッチは一度だけで済みます。各ディスパッチでは、カメラが最後にレンダリングしたマップが使用されます。そのため、updateイベントなどからカメラのレンダリング前にディスパッチされたコンピュートは、前のフレームのマップを使用します。コンピュートがマップを使用しなくなったとき、およびカメラが破棄される前には、nullを渡してマップをデタッチしてください。
コンピュートシェーダーは、sceneDepthCSとsceneColorCSインクルードの関数を使用してマップにアクセスします:
#include "sceneDepthCS"
#include "sceneColorCS"
var output: texture_storage_2d<rgba16float, write>;
@compute @workgroup_size(8, 8, 1)
fn main(@builtin(global_invocation_id) id: vec3u) {
// カメラがレンダリングしたビューポートのテクセルを処理
let viewport = sceneDepthViewport();
if (id.x >= viewport.z || id.y >= viewport.w) {
return;
}
let texel = vec2i(viewport.xy + id.xy);
// カメラ前方にあるサーフェスの深度と、線形に変換したシーンカラー
let depth = sceneDepthLinear(texel);
let color = sceneColorToLinear(sceneColorLoad(texel, 0).rgb);
// 距離に応じてシーンをグレーにフェード
let fogged = mix(color, vec3f(0.5), saturate(depth / 100.0));
textureStore(output, texel, vec4f(fogged, 1.0));
}
sceneDepthCSインクルードは次の関数を提供します:
| 関数 | 説明 |
|---|---|
sceneDepthSize() -> vec2u | 深度マップのサイズ。 |
sceneDepthViewport() -> vec4u | カメラが深度マップをレンダリングしたビューポート(テクセル単位)。最初のテクセルのxとy、および幅と高さです。 |
sceneDepthNearClip() -> f32 | 深度マップのレンダリングに使用されたニアクリップ面。 |
sceneDepthFarClip() -> f32 | 深度マップのレンダリングに使用されたファークリップ面。 |
sceneDepthLinear(texel: vec2i) -> f32 | テクセル位置にあるサーフェスの線形深度。カメラの視線方向に沿った、ワールド単位の値です。 |
sceneDepthWorldPosition(texel: vec2i) -> vec3f | テクセル位置にあるサーフェスのワールド座標。テクセルはビューポート内にある必要があります。 |
sceneColorCSインクルードは次の関数を提供します:
| 関数 | 説明 |
|---|---|
sceneColorSize(lod: i32) -> vec2u | カラーマップのミップレベルのサイズ。 |
sceneColorLoad(texel: vec2i, lod: i32) -> vec4f | テクセルのカラー。 |
sceneColorSample(uv: vec2f, lod: f32) -> vec4f | フィルタリングしてサンプリングしたカラー。カラーマップのフォーマットがデバイスでフィルタリングできない場合は、最も近いテクセルのカラーを返します。 |
sceneColorToLinear(color: vec3f) -> vec3f | カラーマップから読み取ったカラーを線形に変換します。 |
sceneColorToDisplay(color: vec3f) -> vec3f | カラーマップから読み取ったカラーをガンマエンコードに変換します。 |
深度関数は、カメラが深度をどのように保存したかに関係なく同じ値を返します。カラーは、カメラが保存したとおりに返されます。CameraFrameでレンダリングするカメラでは線形、それ以外ではガンマエンコードです。ガンマエンコードの場合、コンピュートシェーダーではSCENE_COLORMAP_GAMMAが定義されます。2つの変換関数はカラーマップから読み取ったカラーを受け取り、すでに要求された色空間にある場合はそのまま返します。
ターゲットの一部にのみレンダリングするカメラ(rectを参照)は、マップのうち自身のビューポートだけをカバーします。マップの残りの部分には、同じターゲットにレンダリングする他のカメラのビューなど、そのカメラがレンダリングしなかった内容が含まれます。カメラのビューだけを処理するには、上の例のように呼び出しをビューポートに対応付け、ビューポートのサイズでディスパッチします。このサイズは、カメラのrectにターゲットのサイズを掛けたものと一致します。
プリプロセッサ
コンピュートシェーダーは、頂点シェーダーやフラグメントシェーダーと同じシェーダープリプロセッサをサポートしており、#define、#ifdef、#if、#includeなどが含まれます。
組み込みインクルード
エンジンは、コンピュートシェーダーで自動的に利用可能な組み込みシェーダーチャンクを提供します:
| インクルード | 説明 |
|---|---|
halfTypesCS | 半精度型エイリアス(half、half2など)。サポートされている場合はf16に、そうでない場合はf32に解決されます。半精度型を参照。 |
sceneDepthCS | カメラのシーン深度マップへのアクセス。シーンの深度マップとカラーマップを参照。 |
sceneColorCS | カメラのシーンカラーマップへのアクセス。シーンの深度マップとカラーマップを参照。 |
dispatchThreadsCS | setupDispatchThreadsで設定したディスパッチのスレッドのスキップとインデックス計算。スレッド数を指定したディスパッチを参照。 |
例:
#include "halfTypesCS"
@compute @workgroup_size(64, 1, 1)
fn main(@builtin(global_invocation_id) global_id: vec3u) {
// 計算にhalf型を使用
var color: half3 = half3(1.0, 0.5, 0.0);
// ...
}
定義とインクルード
cdefinesを使用して定義を渡し、cincludesを使用してインクルードコンテンツを提供します:
const shader = new pc.Shader(device, {
name: 'ComputeShader',
shaderLanguage: pc.SHADERLANGUAGE_WGSL,
cshader: `
#include "myHelper"
@compute @workgroup_size({WORKGROUP_SIZE}, 1, 1)
fn main(@builtin(global_invocation_id) global_id: vec3u) {
var<workgroup> sharedData: array<f32, {WORKGROUP_SIZE}>;
// ...
}
`,
cdefines: new Map([
['{WORKGROUP_SIZE}', '64']
]),
cincludes: pc.ShaderChunks.get(device, pc.SHADERLANGUAGE_WGSL)
});
{WORKGROUP_SIZE}プレースホルダーはコンパイル前に64に置き換えられます。通常の定義とインジェクション定義の詳細については、プリプロセッサドキュメントを参照してください。
サンプル
様々なコンピュートシェーダーのユースケースを示すライブサンプルを探索してください:
- Edge Detect - エッジ検出による画像処理
- Particles - 衝突検出付きGPUベースのパーティクルシミュレーション
- Histogram - アトミック操作を使用した画像ヒストグラムの計算
- Texture Generation - コンピュートシェーダーでテクスチャを生成・変更
- Vertex Update - メッシュ頂点バッファのリアルタイム変更
- Indirect Draw - 間接描画呼び出しによるGPU駆動レンダリング
- Indirect Dispatch - 深度ベースのタイル分類によるGPU駆動コンピュートディスパッチ
- Scene Maps - シーンの深度マップとカラーマップを使用し、シーンを走査するスキャナーパルス