コンピュートシェーダー

このチュートリアルでは、最小限のコンピュートシェーダーを作成するプロセスについて説明します。その前にコンピュートシェーダーと、それがGodotとどのように連携するかについて少し説明します。

注釈

このチュートリアルは読者がシェーダー全般について精通していることを前提としています。シェーダーを初めて使用する場合は、このチュートリアルに進む前に、シェーダーの紹介 と 初めてのシェーダー をお読みください。

コンピュートシェーダーは、汎用プログラミングを対象とした特別なタイプのシェーダープログラムです。言い換えればこれらは固定された目的 (頂点の変換や画像への色の書き込みなど) を持たないため、頂点シェーダーやフラグメントシェーダーよりも柔軟です。フラグメントシェーダーや頂点シェーダーとは異なり、コンピュートシェーダーは舞台裏でほとんど何も行われません。作成したコードはGPUが実行するものであり、それ以外ではほとんど実行されません。これは重い計算をGPUにオフロードするための非常に便利なツールになります。

それでは、短いコンピュートシェーダーを作成することから始めましょう。

まず、任意の 外部 テキスト エディターで、プロジェクトフォルダに compute_example.glsl という新しいファイルを作成します。 Godotでコンピュートシェーダを記述する場合はGLSLで直接記述します。 Godotシェーダ言語はGLSLに基づいています。 Godotの通常のシェーダに慣れている場合は、以下の構文に見覚えがあるでしょう。

注釈

コンピュートシェーダーは、RenderingDeviceベースのレンダラー (Forward+ または Mobileレンダラー) からのみ使用できます。このチュートリアルを進めるには、Forward+ または Mobile レンダラーを使用していることを確認してください。この設定はエディターの右上隅にあります。

モバイルデバイスのコンピュートシェーダーのサポートは、技術的にサポートされている場合でも(ドライバーのバグなどのため)不十分であることに注意してください。

このコンピュートシェーダーのコードを見てみましょう。

#[compute]
#version 450

// Invocations in the (x, y, z) dimension
layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;

// A binding to the buffer we create in our script
layout(set = 0, binding = 0, std430) restrict buffer MyDataBuffer {
    float data[];
}
my_data_buffer;

// The code we want to execute in each invocation
void main() {
    // gl_GlobalInvocationID.x uniquely identifies this invocation across all work groups
    my_data_buffer.data[gl_GlobalInvocationID.x] *= 2.0;
}

このコードはfloatの配列を取得して、各要素を2で乗算し、結果をバッファ配列に格納します。では一行ずつ見てみましょう。

#[compute]
#version 450

この2行は、次の2つのことを意味します。

1. このコードはコンピュートシェーダーです。これはエディタがシェーダーファイルを適切にインポートするために必要なGodot固有のヒントです。

2. このコードはGLSLバージョン450を使用しています。

コンピュートシェーダーをカスタムする場合も、この2 行を変更する必要はありません。

// Invocations in the (x, y, z) dimension
layout(local_size_x = 2, local_size_y = 1, local_size_z = 1) in;

次に各ワークグループで使用される実行インスタンスの数を伝えます。実行インスタンスは同じワークグループ内で実行されるシェーダーのインスタンスです。CPUからコンピュートシェーダーを起動するとき、実行するワークグループの数を指示します。複数のワークグループは相互に並列実行されます。1つのワークグループの実行中は、別のワークグループの情報にアクセスできません。しかし同じワークグループ内では、他の実行インスタンスへのアクセスが制限される場合があります。

ワークグループと実行インスタンスを、このように入れ子になった巨大な for ループとして考えてください。

for (int x = 0; x < workgroup_size_x; x++) {
  for (int y = 0; y < workgroup_size_y; y++) {
     for (int z = 0; z < workgroup_size_z; z++) {
        // Each workgroup runs independently and in parallel.
        for (int local_x = 0; local_x < invocation_size_x; local_x++) {
           for (int local_y = 0; local_y < invocation_size_y; local_y++) {
              for (int local_z = 0; local_z < invocation_size_z; local_z++) {
                 // Compute shader runs here.
              }
           }
        }
     }
  }
}

ワークグループと実行インスタンスは高度なトピックです。現時点ではワークグループごとに2つの実行インスタンスを実行することに注意してください。

// A binding to the buffer we create in our script
layout(set = 0, binding = 0, std430) restrict buffer MyDataBuffer {
    float data[];
}
my_data_buffer;

ここではコンピュート シェーダーがアクセスできるメモリに関する情報を提供します。 layout プロパティを使用すると、シェーダーにバッファの参照する場所を指示できます。後でCPU側から set と binding の位置を一致させる必要があります。

restrict キーワードは、このバッファがこのシェーダ内の1つの場所からのみアクセスされることをシェーダに伝えます。言い換えれば、このバッファを別の set または binding インデックスにバインドしません。これはシェーダーコンパイラがシェーダーコードを最適化できるようにするために重要です。可能な場合は常に restrict を使用してください。

また、これは サイズなし バッファです。つまり任意のサイズにすることができます。したがってバッファのサイズより大きなインデックスから読み取らないように注意する必要があります。

// The code we want to execute in each invocation
void main() {
    // gl_GlobalInvocationID.x uniquely identifies this invocation across all work groups
    my_data_buffer.data[gl_GlobalInvocationID.x] *= 2.0;
}

Finally, we write the main function which is where all the logic happens. We access a position in the storage buffer using the gl_GlobalInvocationID built-in variables. gl_GlobalInvocationID gives you the global unique ID for the current invocation.

続けるには、新しく作成した compute_example.glsl ファイルに上記のコードを書き込みます。

ローカルの RenderingDevice を作成する

コンピュートシェーダーと対話して実行するにはスクリプトが必要です。好きな言語で新しいスクリプトを作成しシーン内の任意のノードにアタッチします。

ここでコンピュートシェーダーを実行するには、RenderingServer を使用して作成できるローカルの RenderingDevice が必要です。

GDScriptC#

# Create a local rendering device.
var rd := RenderingServer.create_local_rendering_device()

その後、新しく作成したシェーダーファイル compute_example.glsl をロードし、これを使用してプリコンパイル済みのシェーダー(SPIR-V)を作成できます。

GDScriptC#

# Load GLSL shader
var shader_file := load("res://compute_example.glsl")
var shader_spirv: RDShaderSPIRV = shader_file.get_spirv()
var shader := rd.shader_create_from_spirv(shader_spirv)

警告

ローカルのRenderingDeviceは、RenderDoc などのツールを使用してデバッグすることができません。

入力データを指定する

念のためやることを確認すると、入力する配列をシェーダーに渡し各要素を2で乗算して結果を取得したいと考えています。

コンピュートシェーダーに値を渡すためのバッファーを作成する必要があります。 この例ではfloatの配列を扱っているためストレージバッファを使用します。ストレージバッファはバイト配列を受け取り、CPUとGPUの間でデータを転送できるようにします。

それではfloatの配列を初期化し、ストレージ バッファを作成しましょう:

GDScriptC#

# Prepare our data. We use floats in the shader, so we need 32 bit.
var input := PackedFloat32Array([1, 2, 3, 4, 5, 6, 7, 8, 9, 10])
var input_bytes := input.to_byte_array()

# Create a storage buffer that can hold our float values.
# Each float has 4 bytes (32 bit) so 10 x 4 = 40 bytes
var buffer := rd.storage_buffer_create(input_bytes.size(), input_bytes)

バッファを配置したら、このバッファを使用するようにRenderingDeviceに指示する必要があります。これを行うには(通常のシェーダーと同様に) Uniformを作成し、それをUniformSetに割り当てる必要があります。これを後でシェーダーに渡します。

GDScriptC#

# Create a uniform to assign the buffer to the rendering device
var uniform := RDUniform.new()
uniform.uniform_type = RenderingDevice.UNIFORM_TYPE_STORAGE_BUFFER
uniform.binding = 0 # this needs to match the "binding" in our shader file
uniform.add_id(buffer)
var uniform_set := rd.uniform_set_create([uniform], shader, 0) # the last parameter (the 0) needs to match the "set" in our shader file

コンピュートパイプラインの定義

次のステップは、GPUが実行できる一連の命令を作成することです。そのためにはパイプラインとコンピュートリストが必要です。

コンピュートシェーダーで結果を得るために必要な手順は次のとおりです:

1. 新しいパイプラインの作成。

2. GPUが実行する命令のリストを開始。

3. コンピュートリストをパイプラインにバインド

4. バッファのUniformをパイプラインにバインド

5. 使用するワークグループの数を指定

6. 命令のリストを終了

GDScriptC#

# Create a compute pipeline
var pipeline := rd.compute_pipeline_create(shader)
var compute_list := rd.compute_list_begin()
rd.compute_list_bind_compute_pipeline(compute_list, pipeline)
rd.compute_list_bind_uniform_set(compute_list, uniform_set, 0)
rd.compute_list_dispatch(compute_list, 5, 1, 1)
rd.compute_list_end()

X軸に5つのワーク グループ、その他に1つのワーク グループを含むコンピュートシェーダーをディスパッチしていることに注意してください。 X 軸 (シェーダーで指定) に2つのローカル呼び出しがあるため、合計10個のコンピュートシェーダーの実行コンテキストが起動されます。バッファの範囲外のインデックスの読み取りまたは書き込みを行うと、シェーダー制御下の外側のメモリや他の変数の一部にアクセスする可能性があり、一部のハードウェアで問題が発生する可能性があります。

コンピュートシェーダーを実行する

これでほぼ完了ですが、これからパイプラインを実行する必要があります。これまでGPUに実行させたいことだけを記録しました。まだ実際にシェーダープログラムを実行していません。

コンピュートシェーダーを実行するには、パイプラインをGPUに送信して実行が完了するのを待つ必要があります。

GDScriptC#

# Submit to GPU and wait for sync
rd.submit()
rd.sync()

すぐに sync() を呼び出してRenderingDeviceを同期させないほうが理想的にはいいです。これはCPUがGPUの処理を完了するまで待機することになるからです。この例ではデータをすぐに読み取るため直後に同期しています。一般的にはGPUがCPUと並行して実行できるように、同期する前に 少なくとも 2～3フレーム待機する必要があります。

警告

コンピュート処理に時間がかかると、WindowsによってトリガーされるTDR (タイムアウト検出と回復) が原因でWindowsグラフィックスドライバーが "クラッシュ" する可能性があります。これはグラフィックスドライバーからのアクティビティがない状態で一定の時間が経過した後 (通常は5～10秒)、グラフィックスドライバーを再初期化するメカニズムです。

コンピューシェーダーの実行にかかる時間に応じて、各ディスパッチにかかる時間を短縮し、TDR がトリガーされる可能性を減らすため、コンピュートシェーダーを複数のディスパッチに分割する必要がある場合があります。 TDRが時間に依存することを考えると、低速なGPUは高速なGPUに比べて、特定のコンピュートシェーダーを実行するときにTDRが発生しやすい可能性があります。

結果の取得

この例のシェーダーでは、ストレージバッファの内容を変更していることに気づいたかもしれません。言い換えればシェーダーは配列からデータを読み取り、同じ配列にデータを保存しているので結果はそこにあります。データを取得して結果をコンソールに出力しましょう。

GDScriptC#

# Read back the data from the buffer
var output_bytes := rd.buffer_get_data(buffer)
var output := output_bytes.to_float32_array()
print("Input: ", input)
print("Output: ", output)

Freeing memory

The buffer, pipeline, and uniform_set variables we've been using are each an RID. Because RenderingDevice is meant to be a lower-level API, RIDs aren't freed automatically. This means that once you're done using buffer or any other RID, you are responsible for freeing its memory manually using the RenderingDevice's free_rid() method.

これでコンピュートシェーダーの使用を開始するために必要なものがすべて揃いました。

参考

The demo projects repository contains a Compute Shader Heightmap demo This project performs heightmap image generation on the CPU and GPU separately, which lets you compare how a similar algorithm can be implemented in two different ways (with the GPU implementation being faster in most cases).