WebGPU シェーディング言語

1. 序論

WebGPU Shading Language (WGSL) は、[WebGPU] のシェーダー言語です。 つまり、WebGPU API を使用するアプリケーションは、GPU 上で実行されるシェーダーと呼ばれるプログラムを WGSL を使用して表現します。

// ポイントライトでテクスチャ付きジオメトリを照明するフラグメントシェーダー。

// ストレージバッファーバインディングからのライト。
struct PointLight {
  position : vec3f,
  color : vec3f,
}

struct LightStorage {
  pointCount : u32,
  point : array<PointLight>,
}
@group(0) @binding(0) var<storage> lights : LightStorage;

// テクスチャーとサンプラー。
@group(1) @binding(0) var baseColorSampler : sampler;
@group(1) @binding(1) var baseColorTexture : texture_2d<f32>;

// 関数の引数は頂点シェーダーからの値。
@fragment
fn fragmentMain(@location(0) worldPos : vec3f,
                @location(1) normal : vec3f,
                @location(2) uv : vec2f) -> @location(0) vec4f {
  // テクスチャーからサーフェスのベースカラーをサンプリングする。
  let baseColor = textureSample(baseColorTexture, baseColorSampler, uv);

  let N = normalize(normal);
  var surfaceColor = vec3f(0);

  // シーンのポイントライトを反復処理する。
  for (var i = 0u; i < lights.pointCount; i++) {
    let worldToLight = lights.point[i].position - worldPos;
    let dist = length(worldToLight);
    let dir = normalize(worldToLight);

    // このライトのサーフェスカラーへの寄与を決定する。
    let radiance = lights.point[i].color * (1 / pow(dist, 2));
    let nDotL = max(dot(N, dir), 0);

    // ライトの寄与をサーフェスカラーに累積する。
    surfaceColor += baseColor.rgb * radiance * nDotL;
  }

  // 累積したサーフェスカラーを返す。
  return vec4(surfaceColor, baseColor.a);
}

1.1. 概要

WebGPU は、GPU コマンドの形式で作業単位を GPU に発行します。 WGSL は 2 種類の GPU コマンドを対象とします:

• 描画コマンドは、 レンダー パイプラインを、 入力、出力、および接続されたリソースのコンテキストで実行します。

• ディスパッチ コマンドは、コンピュートパイプラインを、 入力および接続されたリソースのコンテキストで実行します。

どちらの種類のパイプラインも、WGSL で記述されたシェーダーを使用します。

シェーダーは、 パイプライン内でシェーダー ステージを実行する WGSL プログラムの一部分です。 シェーダーは次のものから構成されます:

• エントリーポイントの関数。

• エントリーポイントから始まる、呼び出されるすべての関数の推移閉包。 この集合には、ユーザー定義関数と組み込み関数の両方が含まれます。 （より厳密な定義については、「シェーダーステージ内の関数」を参照してください。）

• それらすべての関数によって静的にアクセスされる変数および定数の集合。

• それらすべての関数、変数、および定数を定義または解析するために使用される型の集合。

注: WGSL プログラムにエントリーポイントは必須ではありません。ただし、そのような プログラムは、GPUProgrammableStage の作成にエントリーポイントが必要なため、API によって実行できません。

シェーダーステージを実行するとき、実装は次のことを行います:

• モジュールスコープで宣言された定数の値を計算します。

• リソースを シェーダーのリソースインターフェイス内の変数にバインドし、 実行中にそれらのリソースの内容をシェーダーで利用できるようにします。

• その他のモジュールスコープ変数にメモリーを割り当て、 指定された初期値でそのメモリーを設定します。

• エントリーポイントの仮引数が存在する場合、それらにシェーダーステージの入力を設定します。

• エントリーポイントの戻り値が存在する場合、それをシェーダーステージの出力に接続します。

• その後、エントリーポイントを呼び出します。

WGSL プログラムは次の要素で構成されます:

• ディレクティブ。モジュールレベルの動作制御を指定します。

• 関数。実行動作を指定します。

• 文。宣言または実行可能な動作の単位です。

• リテラル。純粋な数学的値のテキスト表現です。

• 定数。それぞれ、特定の時点で計算された値に名前を付けます。

• 変数。それぞれ、値を保持するメモリーに名前を付けます。

• 式。それぞれ、一連の値を組み合わせて結果値を生成します。

• 型。それぞれ、次のものを記述します:

  • 値の集合。

  • サポートされる式に対する制約。

  • それらの式のセマンティクス。

• 属性。オブジェクトを変更して、次のような追加情報を指定します:

  • エントリーポイントへのインターフェイスを指定する。

  • 診断フィルターを指定する。

注: WGSL プログラムは現在、単一の WGSL モジュールで構成されます。

WGSL は命令型言語です。動作は、実行する文のシーケンスとして指定されます。 文では次のことができます:

• 定数または変数を宣言する。

• 変数の内容を変更する。

• 構造化プログラミング構文を使用して実行順序を変更する:

  • 選択的実行: if（任意の else if および else 節を伴う）、switch。

  • 反復: loop、while、for。

  • ネストされた実行構文からの脱出: continue、break、break if。

  • リファクタリング: 関数呼び出しおよびreturn。

• 上記の動作の一部として式を評価して値を計算する。

• シェーダー作成時に、定数式に対する仮定を検査する。

WGSL は静的型付けです。特定の式によって計算される各値は、 プログラムソースを調べることだけで決定される特定の型に属します。

WGSL には、ブール値および数値 （整数と浮動小数点数）を記述する型があります。 これらの型は、複合型 （ベクトル、行列、 配列、および構造体）に集約できます。 WGSL には、固有の操作を提供する特殊な型（例: アトミック型）があります。 WGSL は、メモリーに格納できる型をメモリービューとして記述します。 WGSL は、一般的に使用されるレンダリング型を テクスチャーとサンプラーの形式で提供します。 これらの型には、グラフィックスレンダリング向けに一般的に提供される GPU ハードウェア機能を公開する 組み込み関数が関連付けられています。

WGSL の暗黙の変換と昇格は非常に限定されています。 暗黙の変換と昇格は、抽象 型、および関数呼び出しの引数として渡されるバッファーポインターに対してのみ提供されます。 値をある具象数値型またはブール型から 別の型へ変換するには、明示的な変換、 値コンストラクター、または ビットの再解釈が必要です。ただし、WGSL にはスカラー 型をベクトル型へ昇格するための限定的な機能があります。 これは複合型にも適用されます。

シェーダーステージの処理は、1 つ以上の呼び出しに分割され、 それぞれがエントリーポイントを実行しますが、条件はわずかに異なります。 シェーダーステージ内の呼び出しは、特定の変数へのアクセスを共有します:

• ステージ内のすべての呼び出しは、シェーダーインターフェイス内のリソースを共有します。

• コンピュート シェーダーでは、同じ ワークグループ内の呼び出しが、 workgroup アドレス空間内の変数を共有します。 異なるワークグループ内の呼び出しは、それらの変数を共有しません。

ただし、呼び出しは異なるシェーダーステージ入力の集合に作用します。これには、 呼び出しを他の呼び出しと区別する識別値を提供する組み込み入力も含まれます。 各呼び出しには、private およびfunction アドレス空間内の変数という形式で、 独立した独自のメモリー空間があります。

シェーダーステージ内の呼び出しは同時に実行され、多くの場合並列に実行されることがあります。 シェーダー作者は、シェーダーステージ内の呼び出しの動的な動作が次を満たすようにする責任があります:

• テクスチャーサンプリングや制御バリアなど、特定のプリミティブ操作の 一様性要件を満たす。

• 共有変数への競合する可能性のあるアクセスを調整し、データ競合を回避する。

WGSL は、特定の機能について複数の可能な動作を許可することがあります。 異なる実装が異なる動作を示す可能性があるため、これは移植性上の危険要因です。 WGSL の設計では、このようなケースを最小限に抑えることを目指していますが、実現可能性や、 幅広いデバイスで高いパフォーマンスを実現するという目標による制約があります。

動作 要件とは、WGSL プログラムを処理または実行するときに実装が 行う動作です。これは、プログラマーとの契約における 実装の義務を記述します。 仕様では、それらの義務が他の方法では明らかでない場合に、明示的に記述します。

1.2. 構文表記

以下の構文表記は、WGSL の構文文法の規約を記述します:

• 規則の両側にある斜体のテキストは、構文規則を示します。

• 規則の右辺にある、単一引用符（'）で始まり終わる太字の等幅テキストは、 キーワードおよびトークンを示します。

• 通常のテキスト内のコロン（:）は構文規則を登録します。

• 通常のテキスト内の縦棒（|）は選択肢を示します。

• 通常のテキスト内の疑問符（?）は、直前のキーワード、トークン、規則、 またはグループが 0 回または 1 回出現する（任意である）ことを示します。

• 通常のテキスト内のアスタリスク（*）は、直前のキーワード、トークン、規則、またはグループが 0 回以上出現することを示します。

• 通常のテキスト内のプラス（+）は、直前のキーワード、トークン、規則、またはグループが 1 回以上出現することを示します。

• 通常のテキスト内で対応する開き括弧（(）と閉じ括弧 （)）の組は、要素のグループを示します。

1.3. 数学的な用語と表記

角度:

• 慣例により、角度はラジアンで測定します。

• 角度を測定する基準となる半直線は、原点 (0,0) から (+∞,0) に向かう半直線です。

• θ を比較対象の半直線と基準半直線によって張られる角度とします。 このとき、比較対象の半直線が反時計回りに移動するにつれて θ は増加します。

• 円一周は 2 π ラジアンです。

• 例:

  • 角度 0 は原点から右方向、(1,0) に向かいます

  • 角度 2π は原点から右方向、(1,0) に向かいます

  • 角度 π/4 は原点から点 (1,1) に向かいます

  • 角度 π/2 は原点から点 (0,1) に向かいます

  • 角度 π は原点から点 (-1,0) に向かいます

  • 角度 (3/2)π は原点から点 (0,-1) に向かいます

双曲角は、 従来の意味での角度ではなく、無次元の面積です。 具体的には:

• x > 0 に対して、双曲線 x² - y² = 1 を考えます。

• R を、原点から双曲線上のある点 (x, y) への半直線とします。

• a を、R、x 軸、および双曲線そのものの曲線によって 囲まれる面積の 2 倍とします。

• R が x 軸より上にある場合は a を正、 下にある場合は負とします。

このとき、面積 a は双曲角であり、x は a の双曲線余弦、 y は a の双曲線正弦です。

正の無限大は +∞ で表され、すべての実数より厳密に大きい一意の値です。

負の無限大は −∞ で表され、すべての実数より厳密に小さい一意の値です。

拡張実数 （アフィン拡張実数とも呼ばれます）は、実数に +∞ と −∞ を加えた集合です。 コンピューターは浮動小数点型を使用して、 両方の無限大の値を含む拡張実数を近似的に表現します。 § 15.7 浮動小数点評価を参照してください。

区間は、 下限と上限を持つ連続した数の集合です。 コンテキストに応じて、それらは整数、浮動小数点数、実数、または拡張実数の集合です。

• 閉区間 [a,b] は、a ≤ x ≤ b を満たす数 x の集合です。

• 半開区間 [a,b) は、a ≤ x < b を満たす数 x の集合です。

• 半開区間 (a,b] は、a < x ≤ b を満たす数 x の集合です。

床関数式は、 拡張実数 x に対して次のように定義されます:

• ⌊ +∞ ⌋ = +∞

• ⌊ −∞ ⌋ = −∞

• 実数 x に対して、⌊x⌋ = k。ここで k は k ≤ x < k+1 を満たす一意の整数です

天井 関数式は、拡張実数 x に対して次のように定義されます:

• ⌈ +∞ ⌉ = +∞

• ⌈ −∞ ⌉ = −∞

• 実数 x に対して、⌈x⌉ = k。ここで k は k-1 < x ≤ k を満たす一意の整数です

truncate 関数は、 拡張実数 x に対して次のように定義されます:

• truncate(+∞) = +∞

• truncate(−∞) = −∞

• 実数 x に対して、絶対値が x の絶対値以下である最も近い整数を計算します:

  • x ≥ 0 の場合 truncate(x) = ⌊x⌋、また x < 0 の場合は ⌈x⌉。

roundUp 関数は、 正の整数 k および n に対して次のように定義されます:

• roundUp(k, n) = ⌈n ÷ k⌉ × k

roundDown 関数は、 正の整数 k および n に対して次のように定義されます:

• roundDown(k, n) = ⌊n ÷ k⌋ × k

c 列 r 行の行列 A の転置は、 A の行を A^T の列としてコピーすることによって形成される、r 列 c 行の行列 A^T です:

• transpose(A) = A^T

• transpose(A)_i,j = A_j,i

列ベクトルの転置は、列ベクトルを 1 行の行列として解釈することで定義されます。 同様に、行ベクトルの転置は、行ベクトルを 1 列の行列として解釈することで定義されます。

2. WGSL モジュール

WGSL プログラムは単一の WGSL モジュールで構成されます。

モジュールは、任意のディレクティブのシーケンスと、それに続くモジュールスコープの宣言およびアサーションからなります。 モジュールは次の要素で構成されます:

• ディレクティブ。モジュールレベルの動作制御を指定します。

• 関数。実行動作を指定します。

• 文。宣言または実行可能な動作の単位です。

• リテラル。純粋な数学的値のテキスト表現です。

• 変数。それぞれ、値を保持するメモリーに名前を付けます。

• 定数。それぞれ、特定の時点で計算された値に名前を付けます。

• 式。それぞれ、一連の値を組み合わせて結果値を生成します。

• 型。それぞれ、次のものを記述します:

  • 値の集合。

  • サポートされる式に対する制約。

  • それらの式のセマンティクス。

• 属性。オブジェクトを変更して、次のような追加情報を指定します:

  • エントリーポイントへのインターフェイスを指定する。

  • 診断フィルターを指定する。

2.1. シェーダーのライフサイクル

WGSL プログラムおよびそこに含まれる可能性のあるシェーダーのライフサイクルには、4 つの主要なイベントがあります。 最初の 2 つは、WGSL プログラムを実行用に準備するために使用される WebGPU API メソッドに対応します。 最後の 2 つは、シェーダー実行の開始と終了です。

イベントは次のとおりです:

1. シェーダーモジュール 作成

   • これは、 WebGPU createShaderModule() メソッド が呼び出されたときに発生します。 この時点で WGSL プログラムのソーステキストが提供されます。

2. パイプライン 作成

   • これは、 WebGPU createComputePipeline() メソッド または WebGPU createRenderPipeline() メソッド が呼び出されたときに発生します。 これらのメソッドは、以前に作成された 1 つ以上のシェーダーモジュールを、他の 設定情報とともに使用します。

   • パイプライン作成時に考慮されるのは、 GPUProgrammableStage の指定されたエントリーポイントの シェーダーを形成するコードのみです。 つまり、エントリーポイントと無関係なコードは、コンパイル前に事実上除去されます。

   • 注: 各シェーダーステージは個別に コンパイルされるものとみなされるため、 モジュールの異なる部分が含まれる場合があります。

3. シェーダー 実行開始

   • これは、描画コマンドまたはディスパッチコマンドが GPU に発行され、 パイプラインの実行を開始し、 シェーダーステージのエントリーポイント関数を呼び出したときに発生します。

4. シェーダー実行 終了

   • これは、シェーダー内のすべての処理が完了したときに発生します:

     • そのすべての呼び出しが終了し、

     • リソースへのすべてのアクセスが完了し、

     • 出力がある場合、それが下流のパイプラインステージに渡されます。

これらのイベントは、次の理由により順序付けられます:

• データ依存関係: シェーダーの実行にはパイプラインが必要であり、パイプラインにはシェーダーモジュールが必要です。

• 因果関係: シェーダーは実行を終了する前に実行を開始しなければなりません。

2.2. エラー

WebGPU 実装がシェーダーの処理に失敗する理由は 2 つあります:

• シェーダーが WGSL または WebGPU 仕様の要件を満たしていない場合、 プログラムエラー が発生します。

• 未分類 エラーは、 WGSL および WebGPU のすべての要件を満たしている場合でも発生することがあります。 考えられる原因には次のものがあります:

  • シェーダーが複雑すぎて実装の能力を超えているが、 規定された制限では容易に表現できない場合。 シェーダーを単純化することで問題を回避できる場合があります。

  • WebGPU 実装の欠陥。

処理エラーは、シェーダーのライフサイクルにおける 3 つの段階で発生する可能性があります:

• シェーダー作成 エラーは、 シェーダーモジュール作成時に現実的に検出可能なエラーです。 検出は、WGSL モジュールのソーステキストと createShaderModule API メソッドで利用可能なその他の情報のみに依存します。 この仕様でプログラムが何かを実行しなければならない（must）と記述している文は、 一般に、それらの要件に違反した場合にシェーダー作成エラーを発生させます。

• パイプライン作成エラーは、 パイプライン作成時に検出可能なエラーです。 検出は、WGSL モジュールのソーステキストと 特定のパイプライン作成 API メソッドで利用可能なその他の情報に依存します。 これらのエラーは、GPUProgrammableStage 向けにコンパイルされているエントリーポイントの シェーダー内に存在するコードに対してのみ発生します。

• 動的エラーは、 シェーダー実行中に発生するエラーです。 これらのエラーは検出できる場合も、できない場合もあります。

注: たとえば、データ競合は検出できない場合があります。

各要件は、可能な限り早い機会に検査されます。 つまり:

• シェーダー作成時に検出可能な要件を満たさない場合、シェーダー作成エラーになります。

• パイプライン作成時には検出できるが、それ以前には検出できない要件を満たさない場合、 パイプライン作成エラーになります。

コンテキストから明確でない場合、この仕様では、 特定の要件を満たさないことが シェーダー作成エラー、パイプライン作成エラー、または動的エラーのいずれを引き起こすかを示します。

エラーの結果は次のとおりです:

• シェーダー作成エラーまたはパイプライン作成エラー がある WGSL モジュールは、 パイプラインに組み込まれず、 したがって実行されません。

• 検出可能なエラーは診断をトリガーします。

• 動的エラー が発生した場合:

  • メモリーアクセス は次のものに制限されます:

    • シェーダーステージ入力、

    • シェーダーステージ出力、

    • WGSL モジュール内の変数にバインドされたリソースの任意の部分、および

    • WGSL モジュール内で宣言されたその他の変数。

  • それ以外では、プログラムはこの仕様の残りの部分で記述されているとおりに動作しない場合があります。 注: これらの影響は非局所的である場合があります。

2.3. 診断

実装は、シェーダーモジュール作成中またはパイプライン作成中に 診断を生成できます。 診断は、 アプリケーション作者のために実装によって生成されるメッセージです。

診断は、トリガーされるとき、すなわち トリガールールと呼ばれる特定の条件が満たされたときに作成されます。 条件が満たされたソーステキスト内の場所を、ソーステキスト内の点または範囲として表したものを、 トリガー位置と呼びます。

診断には次の プロパティがあります:

• 重大度。

• トリガールール。

• トリガー位置。

診断の重大度は次のいずれかであり、高いものから低いものの順に並んでいます:

error

診断はエラーです。 これはシェーダー作成エラーまたはパイプライン作成 エラーに対応します。

warning

診断は、アプリケーション開発者が注意を払うに値する異常を記述しますが、 エラーではありません。

info

診断は、アプリケーション開発者が注意を払うに値する特筆すべき状態を記述しますが、 エラーでも警告でもありません。

off

診断は無効です。アプリケーションには通知されません。

トリガールールの名前は、次のいずれかです:

• diagnostic_name_token、または

• ピリオド '.' (U+002E) で区切られた 2 つのdiagnostic_name_token トークン。

2.3.1. 診断処理

トリガーされた診断は次のように処理されます:

1. 各診断 D について、D のトリガー 位置を含み、同じトリガールールを持つ、 影響範囲が最も小さい診断フィルターを見つけます。

   • そのようなフィルターが存在する場合、それを D に適用し、D の重大度を更新します。

   • それ以外の場合、D は変更されません。

2. 重大度がoff の診断を破棄します。

3. 残っている診断の少なくとも 1 つ DI の重大度がinfo である場合:

   • 同じトリガールールを持つ他のinfo 診断を破棄し、元の診断 DI のみを残してもかまいません。

4. 残っている診断の少なくとも 1 つ DW の重大度がwarning である場合:

   • 同じトリガールールを持つ他のinfo または warning 診断を破棄し、元の診断 DW のみを残してもかまいません。

5. 残っている診断の少なくとも 1 つがerror 重大度を持つ場合:

   • error 重大度を持つ他の診断を含め、他の診断を破棄してもかまいません。

   • プログラムエラー が生成されます。

     • 診断がシェーダーモジュール作成時にトリガーされた場合、 エラーはシェーダー作成エラーです。

     • 診断がパイプライン作成時にトリガーされた場合、 エラーはパイプライン作成エラーです。

6. シェーダーモジュール作成時の処理の場合、残りの診断は WebGPU GPUCompilationInfo オブジェクトの messages メンバーに格納されます。

7. パイプライン作成中の処理の場合、error 診断は、 GPUProgrammableStage を検証するときに WebGPU 検証失敗を引き起こします。

注: 規則では、エラーが検出されるとすぐに実装が WGSL モジュールの処理を停止することが許可されています。 さらに、特定の種類の警告の解析は最初の警告で停止でき、 特定の種類の info 診断の解析は最初の発生で停止できます。 WGSL は、異なる種類の解析を実行する順序や、単一の解析内での順序を指定しません。 したがって、同じ WGSL モジュールに対して、異なる実装が 同じ重大度を持つ診断の異なるインスタンスを報告する場合があります。

2.3.2. フィルター可能なトリガールール

ほとんどの診断は無条件に WebGPU アプリケーションへ報告されます。 一部の種類の診断は、そのトリガールールを指定することなどによって、 フィルターできます。 次の表は、フィルター可能なトリガールールの標準セットを示します。

単一の診断 name-tokenで構成される、 認識されないトリガールールを使用した場合、ユーザーエージェントは警告を トリガーするべきです。

実装は、ここで指定されていないトリガールールをサポートしてもかまいません。 ただし、それらはdiagnostic_rule_name の複数トークン形式を使用して記述されなければなりません。 複数トークン形式で記述された認識されないトリガールールの使用自体が、 診断をトリガーしてもかまいません。

この仕様の将来のバージョンでは、特定の規則を削除したり、その既定の重大度を弱めたり （つまり、現在の既定値をより低い重大度の既定値に置き換えたり）しても、後方互換性を 満たしていると見なされる場合があります。 たとえば、WGSL の将来のバージョンでは、derivative_uniformity の既定の重大度を error から warning または info に変更する可能性があります。 仕様にそのような変更が加えられた後も、それ以前に有効だったプログラムは有効なままです。

2.3.3. 診断フィルタリング

フィルター可能なトリガールールを持つ診断が トリガーされると、WGSL は 診断を破棄したり、その重大度を変更したりするための機構を提供します。

診断フィルター DF には 3 つのパラメーターがあります:

• AR: 影響範囲と呼ばれるソーステキストの範囲

• NS: 新しい重大度

• TR: トリガールール

診断フィルター DF(AR,NS,TR) を診断 D に適用すると、 次の効果があります:

• D のトリガー位置が AR 内にあり、 D のトリガールールが TR の場合、 D の重大度プロパティを NS に設定します。

• それ以外の場合、D は変更されません。

範囲診断 フィルターは、診断フィルターのうち、その影響範囲が指定された ソーステキストの範囲であるものです。 範囲診断フィルターは、次の表に示すように、影響を受けるソース範囲の先頭に @diagnostic 属性として指定します。 @diagnostic 属性は、それ以外の場所に現れてはなりません。

注: 次のものも複合文です: 関数本体、case 節、default-alone 節、 while およびfor ループの本体、 ならびにif_clause、else_if_clause、およびelse_clause の本体。

var<private> d: f32;
fn helper() -> vec4<f32> {
  // 「if」の本体で derivative_uniformity 診断を
  // 無効にする。
  if (d < 0.5) @diagnostic(off,derivative_uniformity) {
    return textureSample(t,s,vec2(0,0));
  }
  return vec4(0.0);
}

グローバル 診断フィルターを使用すると、WGSL モジュール全体に診断フィルターを適用できます。

diagnostic(off,derivative_uniformity);
var<private> d: f32;
fn helper() -> vec4<f32> {
  if (d < 0.5) {
    // ここではグローバル診断フィルターによって
    // derivative_uniformity 診断が無効にされている。
    return textureSample(t,s,vec2(0,0));
  } else {
    // derivative_uniformity 診断は 'warning' 重大度に設定される。
    @diagnostic(warning,derivative_uniformity) {
      return textureSample(t,s,vec2(0,0));
    }
  }
  return vec4(0.0);
}

2 つの診断フィルター DF(AR1,NS1,TR1) と DF(AR2,NS2,TR2) は、 次の場合に競合します:

• (AR1 = AR2)、かつ

• (TR1 = TR2)、かつ

• (NS1 ≠ NS2)。

診断フィルターは競合してはなりません。

注: 複数のグローバル診断フィルターは、 競合しない場合に許可されます。

WGSL の診断フィルターは、その影響範囲が完全に入れ子になるよう設計されています。 DF1 の影響範囲と DF2 の影響範囲が重なる場合、 DF1 の影響範囲が DF2 の影響範囲に完全に含まれるか、その逆のいずれかです。

ソース位置 L およびトリガールール TR に対する 最も近い包含診断フィルターは、存在する場合、 次を満たす診断フィルター DF(AR,NS,TR) です:

• L が影響範囲 AR 内にあり、かつ

• L が AR' 内にある別のフィルター DF'(AR',NS',TR) が存在する場合、 AR が AR' に含まれる。

影響範囲は入れ子になるため、最も近い包含診断フィルターは:

• 一意の範囲診断フィルターであるか、

• そうでなければ、重複する（競合しない）グローバル診断 フィルターの集合の 1 つであるか、

• そうでなければ存在しません。

2.4. 制限

WGSL 実装は、次の制限を満たすシェーダーをサポートします。 WGSL 実装は、指定された制限を超えるシェーダーをサポートしてもかまいません。

注: WGSL 実装は、 指定された制限を超えるシェーダーをサポートしない場合、エラーを発行するべきです。

3. テキスト構造

text/wgsl メディアタイプは、コンテンツを WGSL モジュールとして識別するために使用されます。 付録 A: text/wgsl メディアタイプを参照してください。

WGSL モジュールは、バイトオーダーマーク (BOM) を持たない UTF-8 エンコーディングの Unicode テキストです。

WGSL モジュールのテキストは、Unicode コードポイントのシーケンスからなり、連続した空でない集合として次のものにグループ化されます:

• コメント

• トークン

• 空白

プログラムテキストには null コードポイント (U+0000) を含めてはなりません。

3.1. 構文解析

WGSL モジュールを構文解析するには:

1. コメントを削除します:

   • 最初のコメントを空白コードポイント (U+0020) に置き換えます。

   • コメントがなくなるまで繰り返します。

2. § 3.9 テンプレートリストにあるアルゴリズムを使用して、 テンプレートリストを見つけます。 この手順では、'<' (U+003C) および '>' (U+003E) コードポイントを テンプレートリストの区切りとして使用する場合と、比較演算子など他の用途で使用する場合を曖昧さなく区別します。

3. テキスト全体を構文解析し、translation_unit 文法規則との一致を試みます。 構文解析では、次のカスタマイズを施した LALR(1) パーサー（1 トークン先読み）[DeRemer1969]を使用します:

   • トークン化は構文解析と交互に行われ、コンテキストを認識します。 パーサーが次のトークンを要求したとき:

     • 先頭の空白コードポイントのシーケンスを消費して無視します。

     • 次のコードポイントがテンプレートリストの開始である場合、 それを消費して_template_args_start を返します。

     • 次のコードポイントがテンプレートリストの終了である場合、 それを消費して_template_args_end を返します。

     • それ以外の場合:

       • トークン候補とは、まだ消費されていないコードポイントの 空でない接頭辞から形成される任意の WGSL トークンです。

       • 返されるトークンは、現在のパーサー状態に対して有効な先読みトークンでもある、 最長のトークン 候補です。[VanWyk2007]

次の場合、シェーダー作成 エラーになります:

• ソーステキスト全体を有限個の有効なトークンのシーケンスに変換できない場合、または

• translation_unit 文法規則が トークンシーケンス全体と一致しない場合。

3.2. 空白と改行

空白は、 Unicode Pattern_White_Space プロパティに含まれる 1 つ以上のコードポイントの任意の組み合わせです。 次はPattern_White_Space のコードポイント集合です:

• 空白 (U+0020)

• 水平タブ (U+0009)

• 改行 (U+000A)

• 垂直タブ (U+000B)

• フォームフィード (U+000C)

• キャリッジリターン (U+000D)

• 次行 (U+0085)

• 左から右へのマーク (U+200E)

• 右から左へのマーク (U+200F)

• 行区切り (U+2028)

• 段落区切り (U+2029)

改行は、 行の終端を示す空白コードポイントの連続したシーケンスです。 これは、 UAX14 第 6.1 節「カスタマイズ不可能な改行規則」のLB4 およびLB5 で定義される「強制改行」を示す空白として定義されます。 つまり、改行は次のいずれかです:

• 改行 (U+000A)

• 垂直タブ (U+000B)

• フォームフィード (U+000C)

• 後に改行 (U+000A) が続かないキャリッジリターン (U+000D)

• 後に改行 (U+000A) が続くキャリッジリターン (U+000D)

• 次行 (U+0085)

• 行区切り (U+2028)

• 段落区切り (U+2029)

注: 行番号を単位としてソーステキスト位置を報告する診断では、 行数を数えるために改行を 使用するべきです。

3.3. コメント

コメントは、 コメントによってトークンを分離できることを除き、 WGSL プログラムの妥当性や意味に影響を与えないテキスト範囲です。 シェーダー作者はコメントを使用してプログラムを文書化できます。

行末 コメントは、コメントの一種で、 2 つのコードポイント //（U+002F に続いて U+002F）と、それに続く コードポイントからなり、次の直前まで続きます:

• 次の改行、または

• プログラムの終端。

ブロックコメントは、 コメントの一種で、次の要素からなります:

• 2 つのコードポイント /*（U+002F に続いて U+002A）

• 次に、以下の任意のシーケンス:

  • ブロックコメント、 または

  • */（U+002A に続いて U+002F） も /*（U+002F に続いて U+002A） も含まないテキスト

• 次に 2 つのコードポイント */（U+002A に続いて U+002F）

注: ブロックコメントはネストできます。 ブロックコメントには対応する開始テキストシーケンスと終了テキストシーケンスが必要であり、任意のネストが許可されるため、 ブロックコメントを正規表現で認識することはできません。 これは正規言語のポンピング補題から導かれる結果です。

const f = 1.5; // これは行末コメント。
const g = 2.5; /* これは複数行にまたがる
                ブロックコメント。
                /* ブロックコメントはネストできる。
                 */
                ただし、すべてのブロックコメントは終了しなければならない。
               */

3.4. トークン

トークンは、 次のいずれかを形成する連続したコードポイントのシーケンスです:

• リテラル。

• キーワード。

• 予約語。

• 構文トークン。

• 識別子。

• コンテキスト依存名。

3.5. リテラル

リテラルは次のいずれかです:

• ブール リテラル: true または false。

• 数値 リテラル: 整数リテラルまたは浮動小数点リテラルのいずれかで、 数値を表現するために使用されます。

3.5.1. ブールリテラル

const a = true;
const b = false;

3.5.2. 数値リテラル

数値 リテラルの形式は、パターンマッチングによって定義されます。

整数リテラルは:

• 次のいずれかとして指定される整数:

  • 0

  • 最初の桁が 0 ではない 10 進数字のシーケンス。

  • 0x または 0X に続く 16 進数字のシーケンス。

• その後に、任意の i または u 接尾辞。

注: 0 ではない整数リテラルの先頭に 0 を付けること（例: 012）は禁止されています。 これは、他の言語で「先頭の 0 は 8 進数を意味する」という表記との混同を避けるためです。

const a = 1u;
const b = 123;
const c = 0;
const d = 0i;

const a = 0x123;
const b = 0X123u;
const c = 0x3f;

浮動小数点 リテラルは、10 進浮動小数点リテラル または16 進浮動小数点リテラルのいずれかです。

浮動小数点 リテラルには、論理的に 2 つの部分があります。小数を表す仮数部と、任意の指数部です。 おおまかには、リテラルの値は、仮数部に底の値を指定された指数で累乗したものを掛けた値です。 仮数部の桁が 0 ではない場合、またはその左右の仮数部の桁の両方が 0 ではない場合、 その桁を有効とします。 有効桁は左から右に数えます。N 番目の有効桁の左側には N-1 個の 有効桁があります。

10 進 浮動小数点リテラルは:

• 数字のシーケンスとして指定され、その途中に任意の小数点 (.) を持つ仮数部。 仮数部は 10 進表記の小数を表します。

• その後に、次から構成される任意の指数接尾辞:

  • e または E。

  • 次に、任意の先頭符号（+ または -）を持つ 10 進数として指定される指数。

  • 次に、任意の f または h 接尾辞。

• 小数点、指数、または f あるいは h 接尾辞の少なくとも 1 つが 存在しなければなりません。 いずれも存在しない場合、そのトークンは代わりに整数リテラルです。

const a = 0.e+4f;
const b = 01.;
const c = .01;
const d = 12.34;
const f = .0f;
const g = 0h;
const h = 1e-3;

注: 10 進の仮数部は 20 桁の 10 進数字の後で切り捨てられ、小数部にはおよそ log(10)/log(2)×20 ≈ 66.4 有効ビットが保持されます。

16 進浮動小数点リテラルは:

• 0x または 0X 接頭辞

• 次に、16 進数字のシーケンスとして指定され、その途中に任意の 16 進小数点 (.) を持つ仮数部。 仮数部は 16 進表記の小数を表します。

• 次に、次から構成される任意の指数接尾辞:

  • p または P

  • 次に、任意の先頭符号（+ または -）を持つ 10 進数として指定される指数。

  • 次に、任意の f または h 接尾辞。

• 16 進小数点または指数の少なくとも 1 つが存在しなければなりません。 どちらも存在しない場合、そのトークンは代わりに整数リテラルです。

const a = 0xa.fp+2;
const b = 0x1P+4f;
const c = 0X.3;
const d = 0x3p+2h;
const e = 0X1.fp-4;
const f = 0x3.2p+2h;

注: 16 進の仮数部は 16 桁の 16 進数字の後で切り捨てられ、小数部にはおよそ 4 ×16 = 64 有効ビットが保持されます。

数値リテラルに 接尾辞がある場合、そのリテラルは特定の具象スカラー型の値を表します。 それ以外の場合、リテラルは以下で定義される抽象数値型のいずれかの値を表します。 どちらの場合も、リテラルが表す値は、その数学的な値を対象型に変換した後の値であり、 § 15.7.6 浮動小数点変換の規則に従います。

次の場合、シェーダー作成 エラーになります:

• i または u 接尾辞を持つ整数リテラル を対象型で表現できない場合。

• f または h 接尾辞を持つ16 進浮動小数点リテラルが オーバーフローするか、対象型で正確に表現できない場合。

• f または h 接尾辞を持つ10 進浮動小数点リテラルが 対象型をオーバーフローする場合。

• f16 拡張が有効ではないときに、h 接尾辞を持つ浮動小数点リテラルを 使用した場合。

注: 16 進浮動小数点値 0x1.00000001p0 を 正確に表現するには 33 個の仮数ビットが必要ですが、 f32 には明示的な仮数ビットが 23 個しかありません。

注: f 接尾辞を使用して 16 進浮動小数点リテラルの型を強制したい場合、そのリテラルでは 2 進指数も使用する必要があります。たとえば 0x1p0f と記述します。これに対して、0x1f は 16 進整数リテラルです。

3.6. キーワード

キーワードは、 事前定義された言語概念を参照するトークンです。 WGSL キーワードの一覧については、§ 16.1 キーワード概要を参照してください。

3.7. 識別子

識別子は、 名前として使用されるトークンの一種です。 § 5 宣言とスコープを参照してください。

WGSL は、ユースケースを分離するために 2 つの文法非終端記号を使用します:

• ident は、 宣言されたオブジェクトに名前を付けるために使用されます。

• member_ident は、構造体型のメンバーに名前を付けるために使用されます。

識別子の形式は、 Unicode 標準附属書 #31の Unicode バージョン 14.0.0 に基づき、次の詳細化を加えています。

識別子は、UAX31 文法の観点から記述される 次のプロファイルを使用します:

<Identifier> := <Start> <Continue>* (<Medial> <Continue>+)*

<Start> := XID_Start + U+005F
<Continue> := <Start> + XID_Continue
<Medial> :=

これは、次のような非 ASCII コードポイントを含む識別子が 有効であることを意味します: Δέλτα、réflexion、Кызыл、𐰓𐰏𐰇、 朝焼け、سلام、검정、שָׁלוֹם、गुलाबी、 փիրուզ。

ただし、次の例外があります:

• 識別子は、キーワードまたは予約語と 同じ綴りであってはなりません。

• 識別子は _（単一のアンダースコア、 U+005F）であってはなりません。

• 識別子は __（2 つのアンダースコア、 U+005F に続いて U+005F）で始まってはなりません。

Unicode バージョン 14.0.0 の Unicode 文字データベースには、 XID_Start と XID_Continue の両方について、 すべての有効なコードポイントを含む非規範的な一覧があります。

注: 一部の組み込み関数の戻り値の型は、その名前を WGSL ソースで使用できない構造体型です。 それらの構造体型は、2 つのアンダースコアで始まる名前を持つ事前宣言済みであるかのように記述されます。 結果値は、型推論を使用して新たに宣言された let または var に保存するか、 そのメンバーの 1 つを名前によって直ちに 取り出すことができます。使用例については frexp および modf の説明を参照してください。

3.7.1. 識別子の比較

2 つの WGSL 識別子は、同じコードポイントのシーケンスから構成される場合に限り、同一です。

注: この仕様では、比較を目的とした値の Unicode 正規化を許可しません。 視覚的および意味的に同一でも、異なる Unicode 文字シーケンスを使用する値は 一致しません。 コンテンツ作者には、同じエンコーディングシーケンスを一貫して使用するか、 値を選択するときに問題を引き起こす可能性のある 文字を避けることを推奨します。詳細については、[CHARMOD-NORM]を参照してください。

注: 識別子のすべてのインスタンスを その識別子の同形異義文字のいずれかで置き換えた場合に WGSL モジュールの意味が変わるなら、 ユーザーエージェントは開発者から見える警告を発行するべきです。 （同形異義文字とは、別のコードポイントシーケンスと読者には同じように見える可能性があるコードポイントのシーケンスです。 同形異義文字を検出するための対応付けの例としては、前段落で述べた 変換、対応付け、およびマッチングアルゴリズムがあります。 識別子が部分シーケンスをその同形異義文字で繰り返し置き換えることで一方を他方に 変換できる場合、2 つのコードポイントシーケンスは同形異義です。）

3.8. コンテキスト依存名

コンテキスト依存 名は、概念に名前を付けるために使用されるトークンですが、 特定の文法的コンテキストでのみ使用されます。 トークンの綴りは識別子と同じ場合がありますが、そのトークンは宣言されたオブジェクトに解決されません。 この節では、コンテキスト依存名として使用されるトークンを列挙します。 トークンはキーワードまたは予約語であってはなりません。

3.8.1. 属性名

§ 12 属性を参照してください。

属性名は次のとおりです:

• 'align'

• 'binding'

• 'builtin'

• 'compute'

• 'const'

• 'diagnostic'

• 'fragment'

• 'group'

• 'id'

• 'interpolate'

• 'invariant'

• 'location'

• 'blend_src'

• 'must_use'

• 'size'

• 'vertex'

• 'workgroup_size'

• 'subgroup_size'

3.8.2. 組み込み値名

組み込み値 name-tokenは、組み込み値の 名前で使用されるトークンです。

§ 13.3.1.1 組み込み入力と出力を参照してください。

組み込み値名は 次のとおりです:

• 'vertex_index'

• 'instance_index'

• 'position'

• 'front_facing'

• 'frag_depth'

• 'sample_index'

• 'sample_mask'

• 'local_invocation_id'

• 'local_invocation_index'

• 'global_invocation_id'

• 'global_invocation_index'

• 'workgroup_id'

• 'workgroup_index'

• 'num_workgroups'

• 'subgroup_invocation_id'

• 'subgroup_size'

• 'primitive_index'

• 'subgroup_id'

• 'num_subgroups'

3.8.3. 診断ルール名

診断 name-tokenは、診断のトリガールールの名前で使用されるトークンです。

§ 2.3 診断を参照してください。

事前定義された診断ルール名は次のとおりです:

• 'derivative_uniformity'

• 'subgroup_uniformity'

3.8.4. 診断重大度制御名

有効な診断 フィルターの重大度制御名は§ 2.3 診断に列挙されていますが、 識別子と同じ形式を持ちます:

診断フィルターの 重大度制御名は次のとおりです:

• 'error'

• 'warning'

• 'info'

• 'off'

3.8.5. 拡張名

有効なenable-extension 名は§ 4.1.1 Enable 拡張に列挙されていますが、一般には 識別子と同じ形式を持ちます:

enable-extension 名は次のとおりです:

• 'f16'

• 'clip_distances'

• 'dual_source_blending'

• 'subgroups'

• 'primitive_index'

• 'subgroup_size_control'

有効な言語 拡張名は§ 4.1.2 言語拡張に列挙されていますが、 一般には識別子と同じ形式を持ちます:

言語 拡張名は次のとおりです:

• 'readonly_and_readwrite_storage_textures'

• 'packed_4x8_integer_dot_product'

• 'unrestricted_pointer_parameters'

• 'pointer_composite_access'

• 'uniform_buffer_standard_layout'

• 'subgroup_id'

• 'subgroup_uniformity'

• 'texture_and_sampler_let'

• 'texture_formats_tier1'

• 'linear_indexing'

• 'immediate_address_space'

• 'fragment_depth'

• 'buffer_view'

• 'swizzle_assignment'

3.8.6. 補間型名

補間型 name-tokenは、interpolate_type_name の補間型の名前で使用されるトークンです。

§ 13.3.1.4 補間を参照してください。

補間 型名は次のとおりです:

• 'perspective'

• 'linear'

• 'flat'

3.8.7. 補間サンプリング名

補間サンプリング name-tokenは、補間 サンプリングの名前で使用されるトークンです。

§ 13.3.1.4 補間を参照してください。

補間 サンプリング名は次のとおりです:

• 'center'

• 'centroid'

• 'sample'

• 'first'

• 'either'

3.8.8. スウィズル名

スウィズル名は、ベクトルアクセス式およびスウィズルビュー 式で使用されます:

3.8.9. 深度モード名

深度モード name-tokenは、fragment_depth 機能がサポートされている場合に、 frag_depth 組み込みで使用されるトークンです。

• 'less'

• 'greater'

3.9. テンプレートリスト

テンプレート パラメーター化は、一般的な概念を変更するパラメーターを指定する方法です。 テンプレートパラメーター化を記述するには、一般的な概念に続けてテンプレートリストを記述します。

コメントと空白を無視すると、テンプレートリストは次のものです:

• 最初の '<' (U+003C) コードポイント、続いて

• 1 つ以上のテンプレートパラメーターをカンマで区切った リスト、続いて

• 任意の末尾のカンマ、続いて

• 終端の '>' (U+003E) コードポイント。

テンプレート パラメーターの形式は、以下のテンプレートリスト検出アルゴリズムによって暗黙的に定義されます。 一般に、それらは名前、式、または型です。

注: たとえば、vec3<f32> という句は テンプレートパラメーター化であり、vec3 は変更される一般的な概念で、 <f32> は 1 つのパラメーターであるf32 型を含むテンプレートリストです。 これらを合わせた vec3<f32> は、特定のベクトル型を表します。

注: たとえば、 var<storage,read_write> という句は、一般的な var 概念を テンプレートパラメーター storage および read_write で変更します。

テンプレートリストを区切る '<' (U+003C) および '>' (U+003E) コードポイントは、 次の記述にも使用されます:

• relational_expression 内の比較演算子。

• shift_expression 内のシフト演算子。

• シフト操作の後に代入を行うためのcompound_assignment_operator。

構文上の曖昧さは、テンプレートリストを優先して解決されます:

• テンプレートリストは、宣言、式、文が構文解析される前の、構文解析の早い段階で検出されます。

• 後の段階でのトークン化中に、 テンプレートリストの最初の '<' (U+003C) は_template_args_start トークンに対応付けられ、 テンプレートリストの終端の '>' (U+003E) は_template_args_end トークンに対応付けられます。

テンプレート リスト検出アルゴリズムを以下に示します。 これは次の仮定および特性を使用します:

1. テンプレート パラメーターは式であるため、 '<' (U+003C) または '=' (U+003D) コードポイントのいずれでも始まりません。

2. 式には、';' (U+003B)、'{' (U+007B)、または ':' (U+003A) コードポイントは含まれません。

3. 式には代入は含まれません。

4. '=' (U+003D) コードポイントが現れるのは、比較操作の一部として、 つまり次のいずれかの場合のみです: '<='、 '>='、 '=='、または '!='。 それ以外の場合、'=' (U+003D) コードポイントは代入の一部として現れます。

5. テンプレートリスト区切りは、括弧 '(...)' および配列インデックス '[...]' によって形成されるネストした式を考慮します。 テンプレートリストの開始と終了は同じネストレベルに現れなければなりません。

アルゴリズム: テンプレートリスト検出

入力: プログラムのソーステキスト。

レコード型:

UnclosedCandidate を、次を含むレコード型とします:

• position、ソーステキスト内の位置

• depth、整数、position における式のネスト深度

TemplateList を、次を含むレコード型とします:

• start_position、このテンプレートリストを開始する '<' (U+003C) コードポイントの ソース位置。

• end_position、このテンプレートリストを終了する '>' (U+003E) コードポイントの ソース位置。

出力: DiscoveredTemplateLists、TemplateList レコードのリスト。

手順:

• DiscoveredTemplateLists を空のリストに初期化します。

• Pending 変数を、UnclosedCandidate レコードの空のスタックとして初期化します。

• CurrentPosition 整数変数を 0 に初期化します。 これは、現在調べているコードポイントの位置を、ソーステキストの先頭から後にある コードポイント数として符号化します。

  • この変数は、アルゴリズムの実行中にテキスト内を前方に進みます。 テキストの終端に達したら、直ちにアルゴリズムを終了し、 DiscoveredTemplateLists を返します。

• NestingDepth 整数変数を 0 に初期化します。

• 次の手順を繰り返します:

  • CurrentPosition を、空白、コメント、およびリテラルの後まで進めます。

  • ident_pattern_token が CurrentPosition のテキストに一致する場合:

    • CurrentPosition をident_pattern_token の後まで進めます。

    • 存在する場合、CurrentPosition を空白とコメントの後まで進めます。

    • CurrentPosition に '<' (U+003C) が現れる場合:

      • 注: このコード ポイントはテンプレートリストの開始候補です。 入力の後方に現れる終端の '>' (U+003E) と照合できるよう、十分な状態を保存します。

      • UnclosedCandidate(position=CurrentPosition,depth=NestingDepth) を Pending スタックにプッシュします。

      • CurrentPosition を次のコードポイントへ進めます。

      • CurrentPosition に '<' (U+003C) が現れる場合:

        • 注: 仮定 1 により、テンプレートパラメーターは '<' (U+003C) で始まらないため、直前のコードポイントは テンプレートリストの開始ではあり得ません。 したがって、現在のコードポイントと直前のコードポイントは'<<' 演算子でなければなりません。

        • Pending スタックの先頭エントリーをポップします。

        • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

      • CurrentPosition に '=' (U+003D) が現れる場合:

        • 注: 仮定 1 により、テンプレートパラメーターは '=' (U+003C) で始まらないため、直前のコードポイントはテンプレートリストの開始ではあり得ません。 現在のコードポイントと直前のコードポイントが'<=' 比較演算子を形成すると仮定します。 後の手順で代入と誤認しないよう、'=' (U+003D) コードポイントを読み飛ばします。

        • Pending スタックの先頭エントリーをポップします。

        • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

    • ループの次の反復を開始します。

  • CurrentPosition に '>' (U+003E) が現れる場合:

    • 注: このコードポイントは テンプレートリストの終了候補です。

    • Pending が空でない場合、その先頭エントリーを T とし、 T.depth が NestingDepth と等しい場合:

      • 注: このコード ポイントは、開始位置が T に記録された現在のテンプレートリストを終了します。

      • TemplateList(start_position=T.position, end_position=CurrentPosition) を DiscoveredTemplateLists に追加します。

      • T を Pending スタックからポップします。

      • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

    • それ以外の場合、このコードポイントはテンプレートリストを終了しません:

      • CurrentPosition をこのコードポイントの後まで進めます。

      • CurrentPosition に '=' (U+003D) が現れる場合:

        • 注: 現在のコードポイントと直前のコードポイントが'>=' 比較演算子を形成すると仮定します。 後の手順で代入と誤認しないよう、'=' (U+003D) コードポイントを読み飛ばします。

        • CurrentPosition をこのコードポイントの後まで進めます。

      • ループの次の反復を開始します。

  • CurrentPosition に '(' (U+0028) または '[' (U+005B) が現れる場合:

    • 注: ネストした 式に入ります。

    • NestingDepth に 1 を加えます。

    • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

  • CurrentPosition に ')' (U+0029) または ']' (U+005D) が現れる場合:

    • 注: ネストした 式から出ます。

    • Pending スタックが空になるか、その先頭エントリーの depth が NestingDepth より小さくなるまでエントリーをポップします。

    • NestingDepth を、0 または NestingDepth − 1 の いずれか大きい方に設定します。

    • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

  • CurrentPosition に '!' (U+0021) が現れる場合:

    • CurrentPosition をこのコードポイントの後まで進めます。

    • CurrentPosition に '=' (U+003D) が現れる場合:

      • 注: 現在のコードポイントと直前のコードポイントが'!=' 比較 演算子を形成すると仮定します。 後の手順で代入と誤認しないよう、'=' (U+003D) コードポイントを読み飛ばします。

      • CurrentPosition をこのコードポイントの後まで進めます。

    • ループの次の反復を開始します。

  • CurrentPosition に '=' (U+003D) が現れる場合:

    • CurrentPosition をこのコードポイントの後まで進めます。

    • CurrentPosition に '=' (U+003D) が現れる場合:

      • 注: 現在のコードポイントと直前のコードポイントが'==' 比較演算子を形成すると仮定します。 後の手順で代入と誤認しないよう、'=' (U+003D) コードポイントを読み飛ばします。

      • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

    • 注: このコードポイントは 代入の一部であると仮定します。代入は式の一部として現れることができないため、 テンプレートリスト内にも現れません。 保留中の未閉鎖候補を消去します。

    • NestingDepth を 0 に設定します。

    • Pending スタックからすべてのエントリーを削除します。

    • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

  • CurrentPosition に ';' (U+003B)、'{' (U+007B)、または ':' (U+003A) が現れる場合:

    • 注: これらは 式の途中に現れることができないため、テンプレートリスト内にも現れません。 保留中の未閉鎖候補を消去します。

    • NestingDepth を 0 に設定します。

    • Pending スタックからすべてのエントリーを削除します。

    • CurrentPosition をこのコードポイントの後まで進め、 ループの次の反復を開始します。

  • '&&' または'||' が CurrentPosition のテキストと一致する場合:

    • 注: これらは比較より 優先順位の低い演算子です。現在の式レベルにある保留中の未閉鎖候補を拒否します。

    • 注: この規則により、 プログラム断片 a<b || c>d ではテンプレートリストは検出されません。 代わりに、2 つの比較の短絡論理和として認識されます。

    • Pending スタックが空になるか、その先頭エントリーの depth が NestingDepth より小さくなるまでエントリーをポップします。

    • CurrentPosition を 2 つのコードポイントの後まで進め、 ループの次の反復を開始します。

  • CurrentPosition を現在のコードポイントの後まで進めます。

注: 一部の数値リテラルは、たとえば 1.0f のように 文字で終わるため、アルゴリズムはリテラルを明示的に読み飛ばします。 終端の f をident_pattern_token の開始と誤認してはなりません。

注: A ( B < C, D > ( E ) ) という句では、< C, D > の部分はテンプレートリストです。

注: アルゴリズムは式のネストを考慮します。特定の テンプレートリストの開始と終了が異なる式ネストレベルに現れることはできません。 たとえば、array<i32,select(2,3,a>b)> ではテンプレートリストに 3 つのパラメーターがあり、 最後のパラメーターは select(2,3,a>b) です。 a>b 内の '>' は、select 関数を呼び出す式の括弧付き部分内に あるため、テンプレートリストを終了しません。

注: テンプレートリストの両端は、同じ インデックス 式内に現れなければなりません。たとえば a[b<d]>() には有効なテンプレートリストは含まれません。

注: A<B<<C> という句では、 B<<C という句は B、左シフト演算子'<<'、C の順として構文解析されます。 テンプレート検出アルゴリズムは B、続いて '<' (U+003C) を調べ始めますが、 次の '<' (U+003C) コードポイントがテンプレート引数を開始できないことを確認するため、 B の直後の '<' はテンプレートリストの開始ではありません。 最初の '<' と最後の '>' のみがテンプレートリストの区切りであり、 テンプレートパラメーターは B<<C です。

注: A<B<=C> という句も 前の注記と同様に解析されるため、B<=C という句は B、小なりイコール演算子'<='、C の順として構文解析されます。 テンプレート検出アルゴリズムは B、続いて '<' (U+003C) を調べ始めますが、 次の '=' (U+003D) コードポイントがテンプレート引数を開始できないことを確認するため、 B の直後の '<' はテンプレートリストの開始ではありません。 最初の '<' と最後の '>' のみがテンプレートリストの区切りであり、 テンプレートパラメーターは B<=C です。

注: A<(B>=C)> という句を調べると、最初の '<' (U+003C) コードポイントから始まり、最後の '>' (U+003E) コードポイントで終わる 1 つのテンプレートリストがあり、 テンプレート引数は B>=C です。 最初の '>' (U+003C) コードポイント（B の後）を調べた後、 '=' (U+003D) コードポイントが代入の一部であるとみなされないよう、特別に認識する必要があります。

注: A<(B!=C)> という句を調べると、最初の '<'' (U+003C) コードポイントから始まり、最後の '>' (U+003E) コードポイントで終わる 1 つのテンプレートリストがあり、 テンプレート引数は B!=C です。 '!' (U+0021) コードポイント（'B' の後）を調べた後、 '=' (U+003D) コードポイントが代入の一部であるとみなされないよう、特別に認識する必要があります。

注: A<(B==C)> という句を調べると、最初の '<' (U+003C) コードポイントから始まり、最後の '>' (U+003E) コードポイントで終わる 1 つのテンプレートリストがあり、 テンプレート引数は B==C です。 最初の '=' (U+003D) コードポイント（'B' の後）を調べた後、2 番目の '=' (U+003D) コードポイントを特別に認識し、どちらも代入の一部であるとみなされないようにする必要があります。

テンプレート リスト検出が完了した後、 構文解析は、各テンプレートリストをtemplate_list 文法規則と一致させようとします。

4. ディレクティブ

ディレクティブは、 WebGPU 実装による WGSL プログラムの処理方法を変更するトークンシーケンスです。

ディレクティブは任意です。 存在する場合、すべてのディレクティブは、あらゆる宣言またはconst アサーションより前に現れなければなりません。

4.1. 拡張

WGSL は時間とともに進化することが想定されています。

拡張は、 WGSL 仕様に対する一貫した変更の集合に名前を付けてグループ化したものであり、 次の任意の組み合わせからなります:

• 新しい構文による新しい概念および動作の追加。これには次が含まれます:

  • 宣言、文、属性、および組み込み関数。

• 現在の仕様または以前に公開された拡張における制限の削除。

• 許容される動作の集合を減らすための構文。

• プログラムの一部で利用可能な機能を制限するための構文。

• 拡張が既存の仕様、および任意で他の拡張とどのように相互作用するかの説明。

仮に、拡張では次のことが可能です:

• 異なるビット幅の整数など、新しい数値スカラー型を追加する。

• 浮動小数点の丸めモードを制約する構文を追加する。

• シェーダーがアトミック型を使用しないことを示す構文を追加する。

• 新しい種類の文を追加する。

• 新しい組み込み関数を追加する。

• シェーダー呼び出しの実行方法を制約する構文を追加する。

• 新しいシェーダーステージを追加する。

拡張には 2 種類あります: enable-extension と言語拡張です。

4.1.1. 拡張の有効化

enable-extensionは、 次のすべてを満たす場合にのみ機能を利用できる拡張です:

• 実装がそれをサポートし、かつ

• シェーダーがenable ディレクティブによって明示的に要求し、かつ

• 対応する WebGPU GPUFeatureName が、GPUDevice の作成時に要求された必須機能の 1 つであった場合。

Enable-extensionは、 普遍的にはサポートされていないハードウェア機能を公開することを目的としています。

enable ディレクティブは、 1 つ以上の enable-extension のサポートを有効にするディレクティブです。 実装が列挙された enable-extension のすべてをサポートしていない場合、シェーダー作成 エラーになります。

他のディレクティブと同様に、enable ディレクティブが存在する場合、それはすべての宣言およびconst アサーションより前に現れなければなりません。 拡張名は識別子ではありません: 宣言に解決されません。

有効なenable-extensionを次の表に示します。

// 任意精度浮動小数点型のための仮想的な拡張を有効にする。
enable arbitrary_precision_float;
enable arbitrary_precision_float; // 冗長な enable ディレクティブでも問題ない。

// 丸めモードを制御するための仮想的な拡張を有効にする。
enable rounding_mode;

// arbitrary_precision_float により次の使用が有効になると仮定する:
//    - 型 f<E,M>
//    - 関数の戻り値、仮引数、および let 宣言における型として
//    - AbstractFloat からの値コンストラクターとして
//    - 除算演算子 / のオペランドとして
// @rounding_mode 属性が rounding_mode enable ディレクティブによって有効になると仮定する。
@rounding_mode(round_to_even)
fn halve_it(x : f<8, 7>) -> f<8, 7> {
  let two = f<8, 7>(2);
  return x / 2; // 最近接偶数への丸めモードを使用する。
}

4.1.2. 言語拡張

言語拡張は、 実装がサポートしている場合に自動的に利用可能になる拡張です。 プログラムが明示的に要求する必要はありません。

言語拡張は、 どの WebGPU 実装でも合理的にサポートできる機能を具体化します。 機能が普遍的にサポートされていない場合、それは一部の WebGPU 実装がまだ実装していないためです。

注: たとえば、do-while ループは言語 拡張になり得ます。

WebGPU GPU オブジェクトの wgslLanguageFeatures メンバーは、実装がサポートする 言語拡張の集合を列挙します。

requires-directiveは、 プログラムによる 1 つ以上の言語拡張の使用を文書化するディレクティブです。 実装が公開する機能は変更しません。 実装が必須拡張のいずれかをサポートしていない場合、シェーダー作成 エラーになります。

WGSL モジュールは、requires-directive を使用して、移植性がない可能性を示し、 移植性に対する意図された最低条件を示すことができます。

注: WebGPU 実装の外部にあるツールは、 プログラムで使用されるすべての言語拡張が、そのプログラム内のrequires-directiveによって網羅されているかどうかを確認できます。

他のディレクティブと同様に、requires-directive が存在する場合、それはすべての宣言およびconst アサーションより前に現れなければなりません。 拡張名は識別子ではありません: 宣言に解決されません。

注: 意図としては、時間の経過とともに WGSL が、 その時点で一般的にサポートされている言語拡張内のすべての機能を具体化する言語拡張を定義することです。 requires-directive では、これらは そのような一般的な機能をすべて列挙するための短縮記法として機能します。 それらは段階的に増加する機能の集合を表し、ある意味では言語バージョンと考えることができます。

4.2. グローバル診断フィルター

グローバル診断 フィルターは、影響範囲が WGSL モジュール全体である診断フィルターです。 これはディレクティブであるため、 あらゆるモジュールスコープ 宣言より前に現れます。 属性形式と同じように記述しますが、先頭の @ (U+0040) コードポイントはなく、 末尾にセミコロンが付きます。

5. 宣言とスコープ

宣言は、 識別子を次のいずれかの種類の オブジェクトに関連付けます:

• 型

• 型ジェネレーター

• 値

• 変数

• 関数

• 仮 引数

• 列挙子

言い換えると、宣言はオブジェクトの名前を導入します。

宣言がプログラムソース内に現れるものの、他の宣言のテキストの外側にある場合、 その宣言はモジュール スコープにあります。

関数宣言は モジュールスコープに現れます。 関数宣言は、仮引数がある場合にはその宣言を含み、 またその本体内に 変数および値の宣言を含む場合があります。 したがって、それらの包含された宣言はモジュールスコープにはありません。

注: 別の宣言を含む唯一の種類の宣言は、 関数宣言です。

特定のオブジェクトは WebGPU 実装によって提供され、 WGSL モジュールのソース開始前に宣言されているかのように扱われます。 このようなオブジェクトを事前宣言済みと呼びます。 たとえば、WGSL は次を事前宣言します:

• 組み込み 関数、

• i32 やf32 などの組み込み型、

• array、ptr、 texture_2d などの組み込み型ジェネレーター、および

• read_write、 workgroup、rgba8unorm などの列挙子。

宣言のスコープとは、 宣言された識別子が関連付けられたオブジェクトを表し得る プログラムソース位置の集合です。 それらのソース位置では、その識別子が宣言のスコープ内にあるといいます。

宣言が現れる場所によって、そのスコープが決まります:

• 事前宣言済みオブジェクト、およびモジュールスコープで宣言されたオブジェクトは、 プログラムソース全体でスコープ内にあります。

• ユーザー宣言関数の各仮 引数は、対応する関数本体全体でスコープ内にあります。 § 11.1 ユーザー定義関数の宣言を参照してください。

• それ以外の場合、スコープは宣言の終了直後から始まるテキスト範囲です。 詳細については、§ 7 変数と値の宣言を参照してください。

同じ WGSL ソースプログラム内の 2 つの宣言は、同時に次を満たしてはなりません:

• 同じ識別子名を導入し、かつ

• 同じスコープ終了位置を持つ。

注: 事前宣言済みオブジェクトには WGSL ソース内の 宣言がありません。 したがって、モジュールスコープまたは関数内のユーザー指定宣言は、事前宣言済み オブジェクトと同じ名前を持つことができます。

識別子は、文法的コンテキストによって区別され、次のように使用されます:

• ident 文法要素に一致するトークンは:

  • 宣言内で、宣言されるオブジェクトの名前として使用されるか、

  • 別の場所で宣言されたオブジェクトを表す名前として使用されます。これは一般的なケースです。

• member_ident 文法要素に一致するトークンは:

  • 構造体 型の宣言内で、 メンバーの名前として使用されるか、

  • 構造体値のメンバーを表す名前、または構造体のメンバーへの参照を表す名前として使用されます。 § 8.5.4 構造体アクセス式を参照してください。

別の場所で宣言されたオブジェクトを表す名前としてident トークンが現れる場合、それはいずれかの宣言についてスコープ内になければなりません。 識別子トークンが表すオブジェクトは次のように決定されます:

• トークンが少なくとも 1 つの非モジュールスコープ宣言についてスコープ内にある場合、 トークンはそれらのうち最も近い宣言に関連付けられたオブジェクトを表します。

  注: そのような最も近い宣言は、 識別子トークンより前に現れます。

• それ以外の場合、その名前を持つモジュールスコープ宣言があれば、トークンは その宣言されたオブジェクトを表します。

  注: モジュールスコープ宣言は、 識別子トークンの前または後に現れる場合があります。

• それ以外の場合、その名前を持つ事前宣言済みオブジェクトがあれば、トークンはその オブジェクトを表します。

上記のアルゴリズムを使用して識別子を宣言に対応付ける場合、その識別子がその宣言に 解決されるといいます。 同様に、識別子が宣言されたオブジェクトに解決されるともいいます。

いずれかのモジュールスコープ宣言が再帰的である場合、 シェーダー作成エラーになります。 つまり、宣言間に循環が存在してはなりません:

次の有向グラフを考えます:

• 各ノードは宣言 D に対応します。

• D の定義が T に解決される識別子に言及する場合、宣言 D から宣言 T への辺があります。

このグラフに循環があってはなりません。

注: 関数本体は関数宣言の一部であるため、 関数は直接的にも間接的にも再帰的であってはなりません。

注: 非モジュールスコープの識別子宣言は、 テキスト内でその使用より前に置かれなければなりません。

// 有効。ユーザー定義変数は組み込み関数と同じ名前を持つことができる。
var<private> modf: f32 = 0.0;

// 有効。foo_1 はプログラム全体でスコープ内にある。
var<private> foo: f32 = 0.0; // foo_1

// 有効。bar_1 はプログラム全体でスコープ内にある。
var<private> bar: u32 = 0u; // bar_1

// 有効。my_func_1 はプログラム全体でスコープ内にある。
// 有効。foo_2 は関数の終端までスコープ内にある。
fn my_func(foo: f32) { // my_func_1, foo_2
  // 'foo' へのあらゆる参照は関数パラメーターに解決される。

  // 無効。modf はモジュールスコープ変数に解決される。
  let res = modf(foo);

  // 無効。foo_2 のスコープは関数の終端で終了する。
  var foo: f32; // foo_3

  // 有効。bar_2 は関数の終端までスコープ内にある。
  var bar: u32; // bar_2
  // 'bar' への参照は bar_2 に解決される。
  {
    // 有効。foo_4 は複合文の終端までスコープ内にある。
    var foo : f32; // foo_4

    // 有効。bar_3 は複合文の終端までスコープ内にある。
    var bar: u32; // bar_3
    // 'bar' への参照は bar_3 に解決される。
    
    // 無効。bar_4 は bar_3 と同じスコープ終了位置を持つ。
    var bar: i32; // bar_4

    // 有効。i_1 は for ループの終端までスコープ内にある。
    for ( var i: i32 = 0; i < 10; i++ ) { // i_1
      // 無効。i_2 は i_1 と同じスコープ終了位置を持つ。
      var i: i32 = 1; // i_2.
    }
  }

  // 無効。bar_5 は bar_2 と同じスコープ終了位置を持つ。
  var bar: u32; // bar_5

  // 有効。モジュールスコープ宣言 later_def はプログラム全体でスコープ内にある。
  var early_use : i32 = later_def;
}

// 無効。bar_6 は bar_1 と同じスコープを持つ。
var<private> bar: u32 = 1u; // bar_6

// 無効。my_func_2 は my_func_1 と同じスコープ終了位置を持つ。
fn my_func() { } // my_func_2

// 有効。my_foo_1 はプログラム全体でスコープ内にある。
fn my_foo( //my_foo_1
  // 有効。my_foo_2 は関数の終端までスコープ内にある。
  my_foo: i32 // my_foo_2
) { }

var<private> later_def : i32 = 1;

// この宣言は、事前宣言済みの 'min' 組み込み関数を隠す。
// この宣言はモジュールスコープにあるため、ソース全体でスコープ内にある。
// 組み込み関数にはアクセスできなくなる。
fn min() -> u32 { return 0; }

const rgba8unorm = 12; // これは事前宣言済みの 'rgba8unorm' 列挙子をシャドーイングする。

6. 型

プログラムは値を計算します。

WGSL において、型は 値の集合であり、各値は正確に 1 つの型に属します。 値の型は、その値に対して実行できる操作の構文とセマンティクスを決定します。

たとえば、数学的な数 1 は、WGSL では次の異なる値に対応します:

• 32 ビット符号付き整数値 1i、

• 32 ビット符号なし整数値 1u、

• 32 ビット浮動小数点値 1.0f、

• f16 拡張が有効な場合の 16 ビット浮動小数点値 1.0h、

• AbstractInt 値 1、 および

• AbstractFloat 値 1.0

WGSL は、これらのマシン表現と操作が異なるため、別々のものとして扱います。

一部の型は不透明です。 これは、その値が実装定義の表現を持ち、 WGSL プログラム内で直接検査、分解、または操作できないことを意味します。

型は事前宣言済みであるか、 WGSL ソース内の宣言によって作成されます。

一部の型はテンプレートパラメーター化として表現されます。 型ジェネレーターは、 テンプレート リストでパラメーター化されたときに型を表す事前宣言済みオブジェクトです。 たとえば、型 atomic<u32> は、型ジェネレーター atomic と テンプレートリスト <u32> を組み合わせています。

型という概念と、その型を表す WGSL の構文とを区別します。 多くの場合、この仕様での型の綴りは WGSL 構文と同じです。 たとえば:

• 32 ビット符号なし整数値の集合は、この仕様でも WGSL モジュールでも u32 と記述されます。

• 構造体型、または構造体を含む型では、綴りが異なります。

一部の WGSL 型は、ソースプログラムの解析と プログラムの実行時動作の決定にのみ使用されます。 この仕様ではそのような型について説明しますが、WGSL ソーステキストには現れません。

注: 参照型は WGSL モジュールには記述されません。§ 6.5.3 参照型とポインター型を参照してください。

注: スウィズルビュー型は WGSL モジュールには記述されません。§ 6.5.4 スウィズルビュー型を参照してください。

6.1. 型チェック

WGSL の値は、式を評価することによって計算されます。 式は、 名前が「expression」で終わる WGSL 文法規則のいずれかとして 構文解析されたソーステキストの一部分です。 式 E は、外側の式 E に適切に含まれる式である 部分式を含むことができます。 トップレベル 式は、それ自体が部分式ではない式です。 § 8.19 式文法の概要を参照してください。

式の評価によって生成される具体的な値は、次に依存します:

• 静的コンテキスト: 式を取り囲むソーステキスト、および

• 動的 コンテキスト: 式を評価する呼び出しの状態、 およびその呼び出しが実行されている実行コンテキスト。

特定の式を評価した結果として得られる可能性のある値は、常に特定の WGSL 型に属します。これを式の静的型 と呼びます。 WGSL の規則は、式の静的型が式の静的 コンテキストのみに依存するように設計されています。

型アサーションは、 ある WGSL ソース式から WGSL 型への対応付けです。 次の表記

e : T

は、T が WGSL 式 e の静的型であることを意味する型アサーションです。

注: 型アサーションは、プログラムの テキストに関する事実の記述です。 実行時チェックではありません。

文では式が頻繁に使用され、それらの式の静的型に要件が課される場合があります。 たとえば:

• if 文の条件式はbool 型でなければなりません。

• 明示的な型が指定された let 宣言では、初期化子 式はその型に評価されなければなりません。

正常に構文解析された WGSL モジュールの型チェックとは、 各式をその静的型に対応付け、 各文の型要件が満たされていることを検証するプロセスです。 型チェックが失敗すると、シェーダー作成エラーの特殊な場合である 型エラーになります。

型チェックは、構文句に型規則を再帰的に適用することで実行できます。ここで、 構文句とは、式または文のいずれかです。 型規則は、 構文句の静的コンテキストが、 その句に含まれる式の静的型をどのように決定するかを記述します。 型規則には 2 つの部分があります:

• 結論。

  • 句が式の場合、結論はその式に対する型アサーションです。

  • 句が文の場合、結論は型アサーションの集合であり、 文のトップレベル式ごとに 1 つあります。

  • どちらの場合も、構文句は模式的に指定され、 部分式やその他の構文的に決定されるパラメーターを表すために斜体の名前を使用します。

• 前提条件。次から構成されます:

  • 式の場合:

    • 部分式を持つ場合、部分式に対する型アサーション。 それぞれは直接満たされるか、実行可能な自動変換（§ 6.1.2 変換ランクで 定義）によって満たすことができます。 WGSL は、プログラムの静的型を決定するために必要なすべてのconst 式を評価します。

    • 文で式がどのように使用されているか。

  • 文の場合:

    • 文の構文形式、および

    • 文内のトップレベル式に対する型アサーション。

  • その他の模式的パラメーターがある場合、それらに対する条件。

  • 任意で、その他の静的コンテキスト。

型規則は、その前提条件と結論に型 パラメーターを持つ場合があります。 型規則の結論または前提条件に型パラメーターが含まれる場合、その規則は パラメーター化されているといいます。 含まれない場合、その規則は完全に具体化されているといいます。 パラメーター化された型規則から、各型パラメーターに型を代入することで、 完全に具体化された型規則を作成できます。 このとき、規則内の同じパラメーターのすべての出現箇所には同じ型を使用します。 規則の型パラメーターへの型の割り当てを 代入と呼びます。

たとえば、次は論理否定 （!e 形式の式）の型規則です:

これは型パラメーター T を含むため、パラメーター化された規則です。 T は次の 4 つの型のいずれかを表すことができます: bool、vec2<bool>、 vec3<bool>、または vec4<bool>。 T を vec3<bool> に対応付ける代入を適用すると、 次の完全に具体化された型規則が生成されます:

規則のその他の条件を満たす何らかの代入を適用することで パラメーター化された規則から生成できる各完全に具体化された規則を、 そのパラメーター化された規則のオーバーロードと呼びます。 たとえば、ブール否定規則には 4 つのオーバーロードがあります。 これは、その型パラメーター T に型を割り当てる方法が 4 通りあるためです。

注: 言い換えると、パラメーター化された型規則は、 完全に具体化された型規則の集合に対するパターンを提供し、 それぞれはパラメーター化された規則に異なる代入を適用して生成されます。

型規則が構文句に適用されるのは、次の場合です:

• 規則の結論が構文句の有効な構文解析と一致し、かつ

• 規則の前提条件が満たされている。

パラメーター化された型規則が 式に適用されるのは、その式に適用される完全に 具体化された型規則を生成する代入が 存在する場合です。

式 1u+2u を考えます。 これには 1u と 2u という 2 つのリテラル部分式があり、 どちらも型は u32 です。 トップレベル 式は加算です。 § 8.8 算術式の規則を参照すると、 加算の型規則は次の理由でこの式に適用されます:

• 1u+2u は e1+e2 形式の構文解析に一致し、 e1 は 1u、e2 は 2u を表し、かつ

• e1 の型は u32 であり、かつ

• e2 の型は u32 であり、かつ

• 型規則の型パラメーター T に u32 を代入でき、 式全体に適用される完全に 具体化された規則が得られる。

構文 句を解析するとき、次の 3 つの場合があります:

• 式に適用される型規則がない。この場合、型エラーになります。

• ちょうど 1 つの完全に 具体化された型規則が式に適用される。 この場合、規則の結論が表明され、式の静的型が決定されます。

• 複数の型規則が適用される。つまり、複数のオーバーロードの前提条件が満たされます。 この場合、§ 6.1.3 オーバーロード解決で説明されるタイブレーク手順が使用されます。

  • オーバーロード解決が成功すると、単一のオーバーロードが式に適用されるものとして決定されます。 そのオーバーロードの結論にある型 アサーションが表明され、 したがって構文 句内の式の型が決定されます。

  • オーバーロード解決が失敗すると、型エラーになります。

上記の例を続けると、式 1u+2u に適用される型規則は 1 つだけであるため、 型チェックはその型規則の結論、すなわち 1u+2u の型が u32 であることを受け入れます。

WGSL ソースプログラムが適切に型付けされているのは、次の場合です:

• 型規則を適用することで、プログラム内の各式の静的型を決定でき、かつ

• 各文の型要件が満たされている。

それ以外の場合、型エラーがあり、 ソースプログラムは有効な WGSL モジュールではありません。

WGSL が静的型付け言語であるのは、 WGSL モジュールの型チェックが、プログラムのソーステキストを調べるだけで、 成功するか型エラーを検出するかのいずれかになるためです。

6.1.1. 型規則表

式に対する WGSL の型規則は、 型規則 表として編成され、型規則ごとに 1 行があります。

式の セマンティクスとは、その式を評価することによる効果であり、 主として結果値の生成です。 式に適用される型規則の説明列で、その式の セマンティクスを指定します。 セマンティクスは通常、部分式の想定値を含む、 型規則パラメーターの値に依存します。 式のセマンティクスには、部分式の非結果値効果など、 結果値を生成する以外の効果が含まれる場合もあります。

fn foo(p : ptr<function, i32>) -> i32 {
  let x = *p;
  *p += 1;
  return x;
}

fn bar() {
  var a: i32;
  let x = foo(&a); // foo の呼び出しは値を返し
                   // a の値を更新する
}

6.1.2. 変換ランク

型アサーション e:T が型規則の前提条件として使用される場合、 次のとき満たされます:

• e がすでに型 T である場合、または

• e が型 S であり、以下で定義されるように型 S が 型 T に自動的に 変換可能である場合。

この規則は、以下の表で定義される型の組に対するConversionRank 関数によって形式化されます。 ConversionRank 関数は、 ある型（Src）の値を別の型（Dest）に自動変換する際の 優先度と実行可能性を表します。 ランクが低いほど望ましいものです。

実行可能な 自動変換は、値を型 Src から型 Dest に変換するもので、 ConversionRank(Src,Dest) が有限の場合に許可されます。 このような変換は、§ 15.7 浮動小数点評価で説明される制限の範囲内で、値を保持します。

注: 自動変換が発生する状況は 2 種類だけです。 1 つ目は、const 式を、GPU で使用できる対応する型付き数値に変換する場合です。 2 つ目は、メモリーへの参照からロードが発生し、そのメモリーに格納された値が得られる場合です。

注: 無限ランクの変換は実行不可能、すなわち 許可されません。

注: 変換が行われない場合、変換ランクは 0 です。

1. これらの変換には、unrestricted_pointer_parameters およびbuffer_view 言語機能が必要です。

型 T が型 S の具象化であるのは、次の場合です:

• T が具象であり、かつ

• T が参照型ではなく、かつ

• T がスウィズルビュー型ではなく、かつ

• ConversionRank(S, T) が有限であり、かつ

• 上記の基準を満たす他の任意の型 T2 について、ConversionRank(S, T2) > ConversionRank(S, T)。

型 T の値 e の値の 具象化とは、e に対して、T を T の 具象化へ対応付ける実行可能な変換を適用して得られる値です。

注: f32 への変換は常にf16 より優先されるため、 自動変換がf16 値を生成するのは、 モジュールでf16 拡張が有効な場合だけです。

6.1.3. オーバーロード解決

複数の型規則が構文句に適用される場合、 どれを有効にするべきかを決定するために タイブレーク手順を使用します。 この手順をオーバーロード解決と呼び、 部分式の静的型を見つける型チェックが すでに成功しているものと仮定します。

構文 句 P と、P に適用されるすべての型規則を考えます。 オーバーロード解決アルゴリズムでは、これらの型規則をオーバーロード候補と呼びます。 各候補について:

• その前提条件は、直接または自動 変換によって満たされています。

• その結論には次のものがあります:

  • P の有効な構文解析と一致する構文形式、および

  • P 内の各トップレベル式に対応する型 アサーション。

P に対するオーバーロード解決は、単一の最も望ましいオーバーロード 候補を見つけることを目的として、次のように進めます:

1. 各候補 C について、構文句内の部分式の変換ランクを列挙します。 候補の前提条件は満たされているため、P 内の i 番目の部分式について:

   • その静的型は計算済みです。

   • 式の静的型から、前提条件内の対応する型アサーションが要求する型への 実行可能な自動変換があります。 C.R(i) を、その変換のConversionRank とします。

2. 実行可能な自動変換後、その部分式の 1 つが抽象型に解決される一方で、 候補の別の部分式がconst 式ではない候補を除外します。

   注: その結果、句内のいずれかの部分式が const 式でない場合、 句内のすべての部分式は具象型を持たなければなりません。

3. 候補を順位付けします: 2 つのオーバーロード候補 C1 と C2 があるとき、 次の場合、C1 は C2 より優先されます:

   • P 内の各式位置 i について、C1.R(i) ≤ C2.R(i)。

     • つまり、C1 を P に適用するために必要な各式変換が、 C2 を P に適用するために必要な対応する式変換と 少なくとも同程度に望ましいということです。

   • C1.R(i) < C2.R(i) となる式位置 i が少なくとも 1 つ存在する。

     • つまり、C1 を適用するために必要な式変換のうち少なくとも 1 つが、 C2 を適用するために必要な対応する変換より厳密に望ましいということです。

4. 他のすべてより優先される候補 C が 1 つだけ存在する場合、オーバーロード解決は成功し、候補型規則 C が得られます。 それ以外の場合、オーバーロード解決は失敗します。

6.2. プレーン型

プレーン型は、ブール値、数値、ベクトル、 行列、またはそのような値の集合体のマシン表現のための型です。

プレーン型は、スカラー型、アトミック 型、または複合型のいずれかです。

注: WGSL のプレーン型は C++ の Plain-Old-Data 型に似ていますが、 アトミック型および抽象数値型も含みます。

6.2.1. 抽象数値型

これらの型は WGSL ソースで記述できません。型チェックでのみ使用されます。

特定の式はシェーダー作成時に評価され、 GPU が直接実装する範囲より大きな数値範囲および精度を持つ場合があります。

WGSL は、これらの評価のために 2 つの抽象数値型を定義します:

• AbstractInt 型は、符号ビットを最上位ビット位置に持つ 64 ビット 2 の補数形式で 表現可能な整数の集合です。

• AbstractFloat 型は、IEEE-754 のbinary64（倍精度）形式で 表現可能な有限浮動小数点数の集合です。

これらの型のいずれかでの式の評価は、オーバーフローしたり、無限大または NaN 値を生成したりしてはなりません。

型が抽象数値型であるか、抽象数値型を含む場合、その型は抽象です。 型が抽象ではない場合、その型は具象です。

接尾辞のない数値リテラルは、 抽象数値型の値を表します:

• i または u 接尾辞のない整数リテラルは AbstractInt 値を表します。

• f または h 接尾辞のない浮動小数点リテラルはAbstractFloat 値を表します。

例: 式 log2(32) は次のように解析されます:

• log2(32) は、オペランドにAbstractInt 値 32 を持つ log2 組み込み関数への関数呼び出しとして構文解析されます。

• 整数スカラー仮引数を持つ log2 のオーバーロードはありません。

• 代わりにオーバーロード 解決が適用され、3 つの可能なオーバーロードと実行可能な 自動変換が検討されます:

  • AbstractInt からAbstractFloat。（変換ランク 4）

  • AbstractInt からf32。（変換ランク 5）

  • AbstractInt からf16。（変換ランク 6）

• 結果の計算はAbstractFloat として行われます（例: log2(32.0)）。

例: 式 1 + 2.5 は次のように解析されます:

• 1 + 2.5 は、部分式としてAbstractInt 値 1 とAbstractFloat 値 2.5 を持つ加算演算として構文解析されます。

• e が整数型で f が浮動小数点型である e+f のオーバーロードはありません。

• ただし、実行可能な自動変換を使用すると、3 つの潜在的なオーバーロードがあります:

  • 1 はAbstractFloat 値 1.0 に変換され（ランク 4）、 2.5 はAbstractFloat のままです（ランク 0）。

  • 1 はf32 値 1.0f に変換され（ランク 5）、2.5 はf32 値 2.5f に変換されます（ランク 1）。

  • 1 はf16 値 1.0f に変換され（ランク 6）、2.5 はf16 値 2.5h に変換されます（ランク 2）。

• 最初のオーバーロードが望ましい候補であり、型チェックは成功します。

• 結果の計算はAbstractFloat として 1.0 + 2.5 の形で行われます。

例: let x = 1 + 2.5;

• この例は上記と似ていますが、x は抽象数値型に解決できない点が異なります。

• したがって、実行可能なオーバーロード候補は 2 つあります。f32 またはf16 を使用する加算です。

• 望ましい候補はf32 を使用します。

• 宣言の効果は、let x : f32 = 1.0f + 2.5f; と記述した場合と同じです。

例: 1u + 2.5 はシェーダー作成エラーになります:

• 1u 項はu32 型の式です。

• 2.5 項はAbstractFloat 型の式です。

• 有効なオーバーロード候補はありません:

  • GPU 上で具体化された整数スカラー型から 浮動小数点型への実行可能な自動変換はありません。

  • e が整数 スカラー型で、f が浮動小数点型である e+f に一致する型規則はありません。

例: -1 * i32(-2147483648) はシェーダー作成エラーにはなりません:

• -1 項はAbstractInt 型の式です。

• i32(-2147483648) 項はi32 型の式です。

• これら 2 つの型に対する乗算演算子のオーバーロードはなく、i32 項をAbstractInt に昇格することもできません。

• 実行可能な自動変換は、AbstractInt をi32 に変換することだけなので:

  • 結果の計算はi32 間の乗算演算になります。

  • i32 演算は 2 の補数 算術を使用し、オーバーフロー動作が定義されています。

  • ラップアラウンドが発生します。

// 明示的に型指定された符号なし整数リテラル。
var u32_1 = 1u; // 変数は u32 を保持する

// 明示的に型指定された符号付き整数リテラル。
var i32_1 = 1i; // 変数は i32 を保持する

// 明示的に型指定された浮動小数点リテラル。
var f32_1 = 1f; // 変数は f32 を保持する

// 明示的に型指定された符号なし整数リテラルは否定できない。
var u32_neg = -1u; // 無効: 単項マイナスは u32 をサポートしない

// 具象型が必要だが、文または式のどの部分も
// 特定の具象型を強制しない場合、整数リテラルは
// i32 値として解釈される:
//   let 宣言の初期化子は構築可能（またはポインター）でなければならない。
//   AbstractInt から構築可能型への最も優先される自動変換は
//   AbstractInt から i32 で、変換ランクは 2。したがって '1' は i32 と推論される。
let some_i32 = 1; // let some_i32: i32 = 1i; と同様

// 宣言型から推論される。
var i32_from_type : i32 = 1; // 変数は i32 を保持する。AbstractInt から i32、変換ランク 2
var u32_from_type : u32 = 1; // 変数は u32 を保持する。AbstractInt から u32、変換ランク 3

// 接尾辞のない整数リテラルは、必要に応じて浮動小数点へ変換できる:
//   AbstractInt を f32 に自動変換する。変換ランクは 5。
var f32_promotion : f32 = 1; // 変数は f32 を保持する

// 無効: 浮動小数点から整数への実行可能な変換はない
var i32_demotion : i32 = 1.0; // 無効

// 式から推論される。
var u32_from_expr = 1 + u32_1; // 変数は u32 を保持する
var i32_from_expr = 1 + i32_1; // 変数は i32 を保持する

// 値は表現可能でなければならない。
let u32_too_large   : u32 = 1234567890123456890; // 無効、オーバーフロー
let i32_too_large   : i32 = 1234567890123456890; // 無効、オーバーフロー
let u32_large : u32 = 2147483649; // 有効
let i32_large : i32 = 2147483649; // 無効、オーバーフロー
let f32_out_of_range1 = 0x1p500; // 無効、範囲外
let f32_hex_lost_bits = 0x1.0000000001p0; // 無効、f32 で正確に表現できない

// 最小整数: AbstractInt に単項否定を適用し、その後 i32 を推論する。
// AbstractInt から構築可能型への最も優先される変換（最小の
// 変換ランク）は AbstractInt から i32。
let i32_min = -2147483648;  // 型は i32

// 無効。上記と同様に AbstractInt から i32 を選択するが、値が
// 範囲外であり、シェーダー作成エラーになる。
let i32_too_large_2 = 2147483648; // 無効。

// 部分式は AbstractInt および AbstractFloat に解決できる。
// 次の例はすべて有効で、変数の値は 6u。
var u32_expr1 = (1 + (1 + (1 + (1 + 1)))) + 1u;
var u32_expr2 = 1u + (1 + (1 + (1 + (1 + 1))));
var u32_expr3 = (1 + (1 + (1 + (1u + 1)))) + 1;
var u32_expr4 = 1 + (1 + (1 + (1 + (1u + 1))));

// 組み込み関数のパラメーターに基づく推論。

// 最優先候補は clamp(i32,i32,i32)->i32
let i32_clamp = clamp(1, -5, 5);
// 最優先候補は clamp(u32,u32,u32)。
// リテラルは AbstractInt から u32 への自動変換を使用する。
let u32_clamp = clamp(5, 0, u32_from_expr);
// 最優先候補は clamp(f32,f32,f32)->f32
// リテラルは AbstractInt から f32 への自動変換を使用する。
let f32_clamp = clamp(0, f32_1, 1);

// 次の例はすべて、初期値 10f を持つ f32 に昇格する。
let f32_promotion1 = 1.0 + 2 + 3 + 4;
let f32_promotion2 = 2 + 1.0 + 3 + 4;
let f32_promotion3 = 1f + ((2 + 3) + 4);
let f32_promotion4 = ((2 + (3 + 1f)) + 4);

// 型規則違反。

// 無効。初期化子は f32 にしか解決できない:
// AbstractFloat から u32 への実行可能な自動変換はない。
let mismatch : u32 = 1.0;

// 無効。符号の異なるパラメーターを許可する clamp のオーバーロードはない。
let ambiguous_clamp = clamp(1u, 0, 1i);

// 推論は文レベルで完了する。

// let 宣言の初期化子は構築可能（またはポインター）でなければならない。
// AbstractInt から構築可能型への最も優先される自動変換は
// AbstractInt から i32 で、変換ランクは 2。したがって '1' は i32 と推論される。
let some_i32 = 1; // let some_i32: i32 = 1i; と同様

let some_f32 : f32 = some_i32; // 型エラー: i32 は f32 に代入できない

// 別のオーバーフローの例
let overflow_u32 = (1 -2) + 1u; // 無効、-1 は u32 の範囲外

// 理想値は 32 ビットの範囲外だが、その後範囲内に戻される
let out_and_in_again = (0x1ffffffff / 8);

// 同様だが無効
let out_of_range = (0x1ffffffff / 8u); // 計算を 32 ビットで行う必要があり、
                                       // 0x1ffffffff が範囲外になる。

6.2.2. ブール型

bool 型は、 true および false の値を含みます。

6.2.3. 整数型

u32 型は、32 ビット 符号なし整数の集合です。

i32 型は、32 ビット 符号付き整数の集合です。 符号ビットを最上位ビット位置に持つ 2 の補数表現を使用します。

オーバーフローする具象整数型に対する式は、 2^bitwidth を法とする結果を生成します

注: AbstractInt も整数型です。

6.2.4. 浮動小数点型

f32 型は、 IEEE-754 のbinary32（単精度）形式の 32 ビット浮動小数点値の集合です。 詳細については§ 15.7 浮動小数点評価を参照してください。

f16 型は、 IEEE-754 のbinary16（半精度）形式の 16 ビット浮動小数点値の集合です。プログラムに f16 拡張を有効にする enable f16; ディレクティブが含まれていない状態でf16 型を使用すると、シェーダー作成 エラーになります。詳細については§ 15.7 浮動小数点評価を参照してください。

次の表は、浮動小数点型の特定の極値を示します。 それぞれに対応する負の値があります。

注: AbstractFloat も浮動小数点型です。

6.2.5. スカラー型

スカラー型は、bool、AbstractInt、 AbstractFloat、i32、u32、f32、およびf16 です。

数値スカラー型は、 AbstractInt、 AbstractFloat、i32、u32、f32、およびf16 です。

整数スカラー型は、 AbstractInt、i32、および u32 です。

スカラー変換は、 あるスカラー型の値を別のスカラー型の値へ対応付けます。 一般に、結果値は、変換先の型の制限の範囲内で元の値に近いものになります。 スカラー変換は次のいずれかによって発生します:

• 値コンストラクターを明示的に呼び出すことによって、または

• 抽象数値型のconst 式を、 実行可能な自動変換によって 別の型へ変換するとき。

6.2.6. ベクトル型

ベクトルは、 2、3、または 4 個のスカラー 成分をグループ化したシーケンスです。

ベクトルの成分型が数値スカラーである場合、そのベクトルは数値 ベクトルです。

ベクトルの主な用途には次のものがあります:

• 方向と大きさの両方を表現する。

• 空間内の位置を表現する。

• ある色空間の色を表現する。 たとえば、成分は赤、緑、青の強度を表し、 4 番目の成分はアルファ（不透明度）値とすることができます。

ベクトル（および行列）に対する多くの演算は、 成分ごとに動作します。つまり、 各スカラー成分を独立して演算することで結果が形成されます。

vec2<f32>  // 2 つの f32 からなるベクトル。

let x : vec3<f32> = a + b; // a と b は vec3<f32>
// x[0] = a[0] + b[0]
// x[1] = a[1] + b[1]
// x[2] = a[2] + b[2]

WGSL は、次の型エイリアスも 事前宣言します:

6.2.7. 行列型

行列は、 2、3、または 4 個の浮動小数点ベクトルをグループ化したシーケンスです。

行列の主な用途は、線形変換を表現することです。 この解釈では、行列のベクトルは列ベクトルとして扱われます。

積演算子（*）は、次のいずれかに使用されます:

• スカラーの大きさによって変換をスケーリングする。

• 変換をベクトルに適用する。

• 変換を別の行列と組み合わせる。

§ 8.8 算術式を参照してください。

mat2x3<f32>  // これは 32 ビット浮動小数点数の 2 列 3 行の行列。
             // 同等に、型 vec3<f32> の 2 つの列ベクトル。

WGSL は、次の型エイリアスも 事前宣言します:

6.2.8. アトミック型

アトミック型は、 具象整数 型をカプセル化し、次を満たします:

• アトミックオブジェクトは、並行する観測者に対して特定の保証を提供し、

• アトミックオブジェクトに対する有効な操作はアトミック 組み込み関数のみです。

式はアトミック型に評価されてはなりません。

アトミック型は、workgroup アドレス空間内の変数、またはread_write アクセスモードを持つストレージ バッファー変数によってのみインスタンス化できます。 格納型が atomic<vec2<u32>> であるかこれを含む場合、その変数はstorage アドレス空間になければなりません。

その型に対する操作のメモリースコープは、 インスタンス化されたアドレス空間によって決まります。 workgroup アドレス空間内のアトミック型は Workgroup のメモリースコープを持ち、 storage アドレス空間内のものは QueueFamily のメモリースコープを持ちます。

アトミック 変更とは、アトミックオブジェクトの内容を設定する、そのオブジェクトに対する任意の 操作です。 新しい値がオブジェクトの既存値と同じであっても、その操作は変更として数えられます。

WGSL では、アトミック変更はオブジェクトごとに相互に順序付けられます。 つまり、シェーダーステージの実行中、各アトミックオブジェクト A について、すべての エージェントは A に適用される変更操作の同じ順序を観測します。 異なるアトミックオブジェクトの順序付けは互いに関連しない場合があり、 因果関係は暗示されません。 workgroup 空間の変数は ワークグループ内では共有されますが、 異なるワークグループ間では共有されないことに注意してください。

6.2.9. 配列型

配列は、 インデックス付け可能な要素値のシーケンスです。

配列の最初の要素はインデックス 0 にあり、後続の各要素は次の整数インデックスにあります。 § 8.5.3 配列アクセス式を参照してください。

式はランタイムサイズ配列型に評価されてはなりません。

固定サイズ配列の要素数式 N には、次の制約が適用されます:

• これはoverride 式でなければなりません。

• これは具象整数スカラーに評価されなければなりません。

• N が 0 より大きくない場合:

  • N がconst 式である場合はシェーダー作成エラーです。

  • それ以外の場合はパイプライン作成エラーです。

注: N がoverride 宣言に依存する場合、要素数の値はパイプライン作成時に完全に決定され、 それ以外の場合はシェーダーモジュール作成時に完全に決定されます。

注: 型の同値性を満たすには、const 式ではない override 式は識別子でなければなりません。 オーバーライド可能な定数によってサイズ指定される ワークグループ変数を参照してください

ランタイムサイズ配列の要素数は、 対応するストレージバッファー変数に関連付けられた バッファーバインディングのサイズによって決まります。 § 13.3.4 バッファーバインディングによる ランタイムサイズ配列の要素数の決定を参照してください。

配列の要素型は次のいずれかでなければなりません:

• スカラー型

• ベクトル型

• 行列型

• アトミック型

• 作成時固定フットプリントを持つ配列型

• 作成時固定フットプリントを持つ構造体型。

注: 配列の要素型をbuffer 型にすることはできません。

注: 要素型はプレーン型でなければなりません。

2 つの配列型が同一であるのは、次のすべてが真である場合に限ります:

• 要素型が同じである。

• 要素数の指定が一致する。つまり、次のいずれかが真である:

  • 両方ともランタイムサイズである。

  • 両方とも作成時固定フットプリントを持つ固定サイズであり、 一方が符号付きで他方が符号なしであっても、要素数の値が等しい。 （要素数は常に正であるため、この場合は符号付き値と符号なし値を 比較できます。）

  • 両方とも、同じパイプラインオーバーライド可能定数の宣言に解決される識別子として 要素数が指定された固定サイズである。

// array<f32,8> と array<i32,8> は異なる型:
// 要素型が異なる
var<private> a: array<f32,8>;
var<private> b: array<i32,8>;
var<private> c: array<i32,8u>;  // array<i32,8> と array<i32,8u> は同じ型

const width = 8;
const height = 8;

// array<i32,8>、array<i32,8u>、array<i32,width> は同じ型。
// それらの要素数は 8 に評価される。
var<private> d: array<i32,width>;

// array<i32,height> と array<i32,width> は同じ型。
var<private> e: array<i32,width>;
var<private> f: array<i32,height>;

注: オーバーライド可能な定数によってサイズ指定された配列型の 唯一の有効な使用法は、workgroup アドレス空間内のメモリービューとしての使用です。 これには、workgroup 変数の格納型が含まれます。 § 7 変数と値の宣言を参照してください。

override blockSize = 16;

var<workgroup> odds: array<i32,blockSize>;
var<workgroup> evens: array<i32,blockSize>; // 同じ型

// 次のどれも 'odds' および 'evens' と同じ型ではない。

// 異なる型: 識別子 'blockSize' ではない
var<workgroup> evens_0: array<i32,16>;
// 異なる型: 要素数を表すために算術を使用している。
var<workgroup> evens_1: array<i32,(blockSize * 2 / 2)>;
// 異なる型: 単なる識別子ではなく括弧を使用している。
var<workgroup> evens_2: array<i32,(blockSize)>;

// 無効な例。オーバーライド可能な要素数は最外層にのみ現れることができるため。
// var<workgroup> both: array<array<i32,blockSize>,2>;

// 無効な例。オーバーライド可能な要素数は
// workgroup 変数に対してのみ有効であるため。
// var<private> bad_address_space: array<i32,blockSize>;

6.2.10. 構造体型

構造体は、 名前付きのメンバー値を 名前付きでグループ化したものです。

構造体型はモジュールスコープで宣言されます。 プログラムソースの他の場所では、構造体型はその識別子名によって表されます。 § 5 宣言とスコープを参照してください。

2 つの構造体型が同一であるのは、それらが同じ名前を持つ場合に限ります。

構造体メンバー型は次のいずれかでなければなりません:

• スカラー型

• ベクトル型

• 行列型

• アトミック型

• 作成時固定フットプリントを持つ固定サイズ 配列型

• ランタイムサイズ配列 型。ただし構造体の最後のメンバーである場合のみ

• 作成時固定フットプリントを持つ構造体型

注: 構造体メンバー型をbuffer 型にすることはできません。

注: ユーザー宣言されたすべての構造体型は具象です。

注: 各メンバー型はプレーン型でなければなりません。

構造体メンバー型および配列要素型に対する制限から生じる結果には次のものがあります:

• ポインター、テクスチャー、またはサンプラーは、配列または構造体内の どのネストレベルにも現れてはなりません。

• ランタイムサイズ 配列がより大きな型の一部である場合、それは構造体の最後の要素としてのみ現れることができ、 その構造体自体を外側の配列または構造体の一部にすることはできません。

// 3 つのメンバーを持つ構造体。
struct Data {
  a: i32,
  b: vec2<f32>,
  c: array<i32,10>, // 最後のカンマは任意
}

// Data 型の値を格納する変数を宣言する。
var<private> some_data: Data;

構造体メンバーには次の属性を適用できます:

• align

• builtin

• location

• blend_src

• interpolate

• invariant

• size

属性 builtin、location、blend_src、 interpolate、 およびinvariant は IO 属性です。 構造体 S のメンバーにあるIO 属性が効果を持つのは、 S がエントリーポイントの仮引数または戻り値の型として使用される場合だけです。 § 13.3.1 ステージ間入力および出力インターフェイスを参照してください。

属性 align およびsize はレイアウト属性であり、 構造体型をuniform バッファーまたは storage バッファーの定義に使用する場合に 必要になることがあります。 § 14.4 メモリーレイアウトを参照してください。

struct my_struct {
  a: f32,
  b: vec4<f32>
}

// ランタイム配列
alias RTArr = array<vec4<f32>>;
struct S {
  a: f32,
  b: f32,
  data: RTArr
}
@group(0) @binding(0) var<storage> buffer: S;

6.2.11. 複合型

他の型の合成として表現される内部構造を持つ場合、その型は複合です。 内部部分は重ならず、成分と呼ばれます。 複合値はその成分に分解できます。§ 8.5 複合値分解式を参照してください。

複合型は次のとおりです:

• ベクトル型

• 行列型

• 配列型

• 構造体型

複合型 T について、T のネスト深度を NestDepth(T) と表すと、次のようになります:

• ベクトル型の場合は 1

• 行列型の場合は 2

• 要素型 E を持つ配列型の場合は 1 + NestDepth(E)

• T がメンバー型 M₁,...,M_N を持つ構造体型の場合は 1 + max(NestDepth(M₁),..., NestDepth(M_N))

6.2.12. 構築可能型

多くの種類の値は、作成、ロード、格納、関数への受け渡し、 および関数からの返却が可能です。 これらを構築可能と呼びます。

型が構築可能であるのは、 次のいずれかである場合です:

• スカラー型

• ベクトル型

• 行列型

• 作成時固定フットプリントを持ち、要素型が構築可能である固定サイズ 配列型。

• すべてのメンバーが構築可能である構造体型。

注: すべての構築可能型は作成時固定フットプリントを持ちます。

注: アトミック型およびランタイムサイズ配列型は 構築可能ではありません。 アトミックまたはランタイムサイズ配列を含む複合型も構築可能ではありません。

6.2.13. 固定フットプリント型

変数のメモリーフットプリントは、 変数の内容を格納するために使用されるメモリー位置の数です。 変数のメモリーフットプリントはその格納型に依存し、 シェーダーライフサイクルのある時点で確定します。 ほとんどの変数は非常に早い段階、すなわちシェーダー作成時にサイズが決まります。 一部の変数は後のパイプライン作成時にサイズが決まる場合があり、 また他の変数はシェーダー実行開始時までサイズが決まらない場合があります。

型の具象化のサイズが シェーダー作成時に完全に決定される場合、その型は 作成時固定フットプリントを持ちます。

型のサイズがパイプライン作成 時に完全に決定される場合、その型は固定 フットプリントを持ちます。

注: すべての具象の作成時固定フットプリント型および固定フットプリント型は格納可能です。

注: パイプライン作成はシェーダー作成に依存するため、 作成時固定フットプリントを持つ型は固定フットプリントも持ちます。

作成時固定フットプリントを持つ型は次のとおりです:

• スカラー型

• ベクトル型

• 行列型

• アトミック型

• 次の場合の固定サイズ 配列型:

  • その要素 数がconst 式である。

• すべてのメンバーが作成時固定フットプリントを持つ構造体型。

• 次の場合の固定サイズ バッファー型:

  • そのサイズがconst 式である。

注: 構築可能型は作成時固定フットプリントを持ちます。

固定 フットプリントを持つ型は次のいずれかです:

• 作成時固定フットプリントを持つ型

• その要素数をさらに制約しない固定サイズ 配列型

• そのサイズをさらに制約しない固定サイズ バッファー型

注: 要素数が const 式ではないoverride 式である固定サイズ配列の唯一の有効な使用法は、 workgroup アドレス空間内のメモリービューとしての使用です。 これには workgroup 変数の格納型が含まれます。

注: 固定フットプリント型はアトミック型を直接または 間接的に含むことができますが、構築可能型は含むことができません。

注: 固定フットプリント型にはランタイムサイズ配列、 ランタイムサイズ バッファー、およびランタイムサイズ配列を含む構造体は含まれません。

6.3. バッファー型

これらの型にはbuffer_view 言語機能が必要です。

buffer 値は、 その内容を別のホスト共有可能型として 解釈できるメモリーを表します。 それ以外では、バッファー値は不透明です。 これは、変数内のデータに複数の異なる型としてアクセスする必要がある場合や、 シェーダー作者が 1 つの変数を複数の論理変数に分割したい場合に役立ちます。

バッファー型は構築可能ではなく、 buffer_view 組み込み関数を介して データを不透明でないホスト共有可能型として解釈した後にのみアクセスできます。

式は バッファー型に評価されてはなりません。

固定サイズバッファーは、storage、uniform、およびworkgroup アドレス空間の変数によってのみインスタンス化できます。 storage および uniform アドレス空間では、サイズはconst 式でなければなりません。 ランタイムサイズバッファーは storage アドレス空間の変数によってのみインスタンス化できます。

N が 0 より大きくない場合:

• N がconst 式である場合はシェーダー作成エラーです。

• それ以外の場合はパイプライン作成エラーです。

f16 拡張が有効で、 N が 2 で割り切れない場合:

• N がconst 式である場合はシェーダー作成エラーです。

• それ以外の場合はパイプライン作成エラーです。

f16 拡張が 有効ではなく、N が 4 で割り切れない場合:

• N がconst 式である場合はシェーダー作成エラーです。

• それ以外の場合はパイプライン作成エラーです。

2 つのバッファー型が同一であるのは、次の条件のいずれかを満たす場合に限ります:

• 両方ともランタイムサイズである。

• 両方とも作成時固定フットプリントを持つ固定サイズであり、 一方が符号付き、他方が符号なしであってもサイズの値が等しい。（サイズは常に 正であるため、この場合は符号付き値と 符号なし値を比較できます。）

• 両方とも、同じパイプラインオーバーライド可能定数の宣言に解決される 識別子としてサイズが指定された固定サイズである。

6.4. 列挙型

列挙型は、 名前付きの値からなる限定された集合です。 列挙は、有効なテクセルフォーマットの集合など、 特定の概念について可能な選択肢の集合を区別するために使用されます。

列挙子は、 列挙内の名前付き値の 1 つです。 各列挙子は他のすべての列挙子とは異なり、 また他のすべての種類の値とも異なります。

WGSL ソースで新しい列挙子または新しい列挙型を宣言する機構はありません。

注: 列挙子はテンプレートパラメーターとして使用されます。

6.4.1. 事前宣言済み列挙子

次の表は、WGSL の列挙型、それらの事前宣言済み列挙子、および列挙型に必要な言語拡張を示します。 列挙型は存在しますが、WGSL ソースでは記述できません。

6.5. メモリービュー

プレーン 値を使用した計算に加えて、WGSL プログラムは メモリーアクセス操作を介して、 メモリーから値を読み取ったり、メモリーに値を書き込んだりすることもよくあります。 各メモリーアクセスはメモリービューを介して実行されます。

メモリービューは次のものから構成されます:

• 特定のアドレス空間内のメモリー 位置の集合、

• メモリーモデル参照、

• それらの位置の内容を WGSL 型として解釈したもの。これは格納型と呼ばれます、および

• アクセスモード。

メモリービューのアクセスモードは、アドレス空間でサポートされていなければなりません。§ 7 変数と値の宣言を参照してください。

6.5.1. 格納可能型

変数に含まれる値は、 格納可能型でなければなりません。 格納可能型は、§ 14.4.4 値の内部レイアウトで説明されるように、 WGSL によって定義された明示的な表現を持つ場合もあれば、テクスチャーやサンプラーのように 不透明である場合もあります。

型が格納可能であるのは、 具象であり、かつ次のいずれかである場合です:

• スカラー型

• ベクトル型

• 行列型

• アトミック型

• 配列型

• 構造体型

• テクスチャー型

• サンプラー型

• バッファー型

注: つまり、格納可能型は具象のプレーン型、 テクスチャー型、サンプラー型、およびバッファー型です。

6.5.2. ホスト共有可能型

ホスト共有可能型は、ホストと GPU の間で共有されるバッファー、 またはフォーマット変換なしでホストと GPU の間でコピーされるバッファーの内容を記述するために使用されます。 この目的で使用する場合、§ 14.4 メモリーレイアウトで説明されるように、 型にはさらにレイアウト属性を適用できます。 § 7.3 var 宣言で説明されるように、uniform バッファーおよびstorage バッファー 変数の格納型は ホスト共有可能でなければなりません。

型がホスト共有可能であるのは、 具象であり、かつ 次のいずれかである場合です:

• 数値スカラー型

• 数値ベクトル型

• 行列型

• アトミック型

• 作成時固定フットプリントを持ち、その要素型が ホスト共有可能である固定サイズ 配列型

• その要素型がホスト共有可能であるランタイムサイズ配列 型

• すべてのメンバーがホスト共有可能である構造体型

• バッファー型

注: ステージ間の入力および出力の型に対する制限は、 § 13.3.1 ステージ間入力および出力インターフェイスおよび それ以降の節で説明されています。 これらの型にもサイズがありますが、数え方が異なります。

注: テクスチャーと サンプラーもホストと GPU の間で共有できますが、その内容は不透明です。 この節のホスト共有可能型は、特にstorage およびuniform バッファーで使用するためのものです。

6.5.3. 参照型とポインター型

WGSL には、メモリービューを表現するための 2 種類の型があります: 参照型とポインター型です。

2 つのポインター型が同一であるのは、アドレス空間、格納型、およびアクセスモードが同じ場合に限ります。

WGSL モジュールを解析するとき、参照型とポインター型は、 アドレス空間、格納可能型、およびアクセスモードによって完全にパラメーター化されます。 この仕様のコード例では、コメントにこの完全にパラメーター化された形式を示します。

ただし、WGSL ソーステキストでは:

• 参照型は現れてはなりません。

• ポインター型は現れることができます。

  • ポインター型は、次のものによってパラメーター化して記述されます:

    • アドレス 空間、

    • 格納型、 および

    • 場合によっては、§ 14.3 アドレス空間で指定されるアクセスモード。

  • ポインター型がプログラムソースに現れる場合、 そのポインター型のアドレス 空間、格納 型、およびアクセスモードを使用して、 プログラム内のどこかで変数を宣言するか、 関数呼び出しから その型を返すことも有効でなければなりません。

    注: この制限により、 実行時に決して使用できない特定の型エイリアスおよび関数の仮引数の宣言が禁止されます。 この制限がなければ、ポインター型へのエイリアスを宣言することは有効でも、 その型のポインター値を作成することは決してできない場合があります。 同様に、ポインター型の仮引数を持つ関数を宣言することは有効でも、 その関数を呼び出すことは決してできない場合があります。

fn my_function(
  /* 'ptr<function,i32,read_write>' は、'function' アドレス空間のメモリー位置を使用して
     'i32' 値を保持するためのメモリーを参照するポインター値の型。
     ここでは 'i32' が格納型。
     暗黙のアクセスモードは 'read_write'。
     既定値については「アドレス空間」節を参照。 */
  ptr_int: ptr<function,i32>,

  // 'ptr<private,array<f32,50>,read_write>' は、'private' アドレス空間の
  // メモリー位置を使用して、型 'f32' の 50 個の要素からなる配列を保持する
  // メモリーを参照するポインター値の型。
  // ここで格納型は 'array<f32,50>'。
  // 暗黙のアクセスモードは 'read_write'。
  // 既定値については「アドレス空間」節を参照。
  ptr_array: ptr<private, array<f32, 50>>
) { }

参照型とポインター型は、どちらもメモリービューの集合です: 特定のメモリービューには、一意の参照値と一意のポインター値の両方が関連付けられます:

型 ptr<AS,T,AM> の各ポインター値 p は、 型 ref<AS,T,AM> の一意の参照値 r に対応し、 その逆も同様です。 ここで p と r は同じメモリービューを記述します。

6.5.4. スウィズルビュー型

スウィズルビューを使用すると、 メモリー内のベクトルの成分の部分集合にアクセスし、 読み取り時にはそれらを並べ替え、書き込み時にはそれらを散在させることができます。

スウィズルビューはswizzle_assignment 言語拡張によって有効になります。

アルゴリズム: SwizzleViewRead

ポインター p を介して読み取ったベクトルの成分を選択し、任意で並べ替えて形成したベクトルを生成します。

注: スウィズルビューを介した読み取りでは、 一部のベクトル成分が結果に現れない場合でも、ベクトル全体をメモリーから読み取ります。

入力: 型 swizzle<AS,S,N,K> のスウィズル ビュー (p, IndexList)

出力: (p, IndexList) を読み取った結果である、 型 vecK<S> のベクトル値 v。

手順:

1. 慣例により IndexList = « Idx₀, ..., Idx_K−1 »

2. p が参照するメモリーからベクトル u を読み取ります。u の型は vecN<S> です。

3. 次を満たす型 vecK<S> の値 v を形成します:

   • v の成分 i は、 u の成分 Idx_i の値に等しい。 インデックス構文を使用すると: v[i] = u[Idx_i]

アルゴリズム: SwizzleViewWrite

update ベクトルの成分を、ポインター p が参照するメモリー内の ベクトルの指定された成分へ書き込みます。

注: スウィズルビューを介した書き込みでは、処理によって成分値が変更されない場合でも、 メモリー内のベクトル全体を書き込みます。

スウィズルビューへの書き込みが一様性解析および メモリーモデルとどのように相互作用するかについては、以下の説明を参照してください。

入力:

• 型 swizzle<AS,S,N,K> のスウィズルビュー (p, IndexList)

• 型 vecN<S> のベクトル old_contents

  • old_contents は、過去のある時点で p が参照するメモリーに格納されていたベクトルであることを意図しています。

• 型 vecK<S> のベクトル update

IndexList 内のすべてのインデックスは互いに異なっていなければなりません。 この要件から K ≤ N が導かれます。

変数:

• 型 vecN<S> の new

手順:

1. 慣例により IndexList = « Idx₀, ..., Idx_K−1 »

2. 変数 new を old_contents で初期化します。

3. 0..K-1 の各 i について:

   • new の成分 Idx_i を、 update の成分 i の値に設定します。 インデックス構文を使用すると: new[Idx_i] = update[i]。

4. new を p が参照するメモリーに書き込みます。

注: すべての Idx_i を 異なるものとすることで、 手順 3 が new の各成分を最大 1 回しか更新しないことが保証されます。

注: 手順 2 と 3 は、次の宣言的な規則に置き換えることができます:

型 vecN<S> のベクトル new を形成します。ここで new の成分 i は:

• 0...K−1 のある j について i = Idx_j の場合、 update の成分 j。

• それ以外の場合、old_contents の成分 i。

スウィズルビューにはスウィズルビュー型があります:

スウィズルビューは次の方法で形成されます:

• read_write アクセスモードを持ち、 格納型がベクトルである参照値またはポインター値に、 複数文字のスウィズル 名を適用する、または

• スウィズルビューに複数文字のスウィズル 名を適用する。

これらの規則により、スウィズルビューの IndexList 部分は シェーダー 作成時に完全に決定されます。

詳細については§ 8.6 スウィズルビュー式を参照してください。

fn swizzle_read_and_write() {
  var v: vec4u;

  // 次と同じ:  v.y = 1; v.z = 2;
  // ただし、ベクトルはメモリーから正確に 1 回読み取られ、メモリーへ正確に 1 回書き込まれる。
  v.yz = vec2u(1,2);

  // 次と同じ: let u = vec3(v.a, v.g, v.b);
  // 式 v.agb は型 swizzle<function,u32,4,3> のスウィズルビュー。
  // let 宣言の実効値型は具象の構築可能型
  // またはポインター型でなければならない。そのためスウィズルビューは、
  // スウィズルビューロード規則を呼び出すことで自動的に vec3u に変換される。
  let u = v.agb;

  // スウィズルは連鎖できる。
  // 次と同じ: v.zy = vec2(99,100);
  v.yz.yx = vec2(99,100);

  // スウィズルビューに単一文字を適用すると、基礎となるベクトルの
  // スカラー成分への参照が得られる。
  // 次と同じ: v.z = 99;
  v.yz.y = 99;

  // スウィズルビューをインデックス付けすると、ベクトルの成分の 1 つへの参照が得られる。
  // 単一文字のベクトルアクセスを形成するかのように動作するが、
  // インデックスによって、スウィズル名からどの文字を使用するかが選択される。
  v.zy[1] = 50; // v.y = 50; と同じ
}

// この例では、swizzle_assignment に加えて pointer_composite_access のサポートも前提とする。
requires pointer_composite_access;

fn swizzle_read_and_write_via_pointer(p: ptr<function,vec4u>) {
  // 次と同じ:  (*p).y = 1; (*p).z = 2;
  p.yz = vec2u(1,2);

  // 次と同じ: let u = vec3((*p).a, (*p).g, (*p).b);
  let u = p.agb;

  // スウィズルは連鎖できる。
  // 次と同じ: p.zy = vec2(99,100);
  p.yz.yx = vec2(99,100);

  // スウィズルビューはポインターから構築できる。
  // スウィズルビューに単一文字を適用すると、基礎となるベクトルの
  // スカラー成分への参照が得られる。
  // 次と同じ: (*p).z = 99;
  p.yz.y = 99;

  // スウィズルビューをインデックス付けすると、ベクトルの成分の 1 つへの参照が得られる。
  // 単一文字のベクトルアクセスを形成するかのように動作するが、
  // インデックスによって、スウィズル名からどの文字を使用するかが選択される。
  p.zy[1] = 50; // (*p).y = 50; と同じ
}