WebGPUシェーディング言語

1. はじめに

WebGPUシェーディング言語（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プログラムはエントリーポイントを必須としません。ただし、エントリーポイントがないプログラムはAPIによって実行できません。なぜなら、エントリーポイントがGPUProgrammableStageの作成に必要だからです。

シェーダーステージを実行する際、実装は次の処理を行います：

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

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

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

• エントリーポイントに形式的パラメータがある場合、シェーダーステージの入力値でそれらを初期化します。

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

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

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) 方向への光線です。

• θ を比較光線と基準光線のなす角度とします。 比較光線が反時計回りに動くと θ は増加します。

• 1周（円）は 2π ラジアンです。

• 例：

  • 角度0は原点から右 (1,0) 方向

  • 角度2πも原点から右 (1,0) 方向

  • 角度π/4は原点から (1,1) 方向

  • 角度π/2は原点から (0,1) 方向

  • 角度πは原点から (-1,0) 方向

  • 角度(3/2)πは原点から (0,-1) 方向

双曲線角は、従来の角度ではなく単位のない面積です。具体的には：

• x² - y² = 1 かつ x > 0 の双曲線を考えます。

• R を原点から双曲線上の点 (x, y) への光線とします。

• a を、R、x軸、双曲線の曲線で囲まれる面積の2倍とします。

• R がx軸より上なら a を正、下なら負とします。

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

正の無限大（+∞）はすべての実数より大きい固有の値です。

負の無限大（−∞）はすべての実数より小さい固有の値です。

拡張実数（またはアフィン拡張実数）は実数全体に+∞と−∞を加えた集合です。 コンピューターは浮動小数点型で拡張実数を近似表現し、両方の無限大値も含みます。 § 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 ≤ x < k+1 となる唯一の整数 k）

天井関数は、拡張実数 x に対して：

• ⌈ +∞ ⌉ = +∞

• ⌈ −∞ ⌉ = −∞

• 実数 x について、⌈x⌉ = k（k-1 < x ≤ k となる唯一の整数 k）

切り捨て関数は、拡張実数 x に対して：

• truncate(+∞) = +∞

• truncate(−∞) = −∞

• 実数 x について、絶対値が x 以下となる最近接の整数値：

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

roundUp関数は、正の整数 k と n に対し：

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

転置は、c列 r行の行列 Aを、r列 c行の行列 A^Tとして Aの行を 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メソッドで得られる他の情報のみが利用されます。 この仕様でプログラムが必須とされていることに違反した場合、通常はシェーダー作成エラーとなります。

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

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

注意: 例えば、データレースは検出できない場合があります。

各要件willは最も早く検出可能なタイミングでチェックされます。 つまり：

• シェーダー作成時に検出できる要件違反はシェーダー作成エラーとなります。

• パイプライン作成時にのみ検出できる要件違反はパイプライン作成エラーとなります。

文脈からはっきりしない場合、この仕様書は 特定の要件の違反が シェーダー作成エラー・パイプライン作成エラー・動的エラーのどれにつながるかを明示します。

エラーの結果は以下の通りです：

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

• 検出可能なエラーはwillトリガーされ、診断が出力されます。

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

  • メモリアクセスは 以下に制限されます：

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

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

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

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

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

2.3. 診断

実装はシェーダーモジュール作成またはパイプライン作成の間に診断を生成できます。 診断とは、実装がアプリケーション制作者のために出力するメッセージです。

診断は、特定の条件が満たされたときに作成またはトリガーされます。 この条件をトリガールールと呼びます。 条件が満たされたソーステキスト中の場所（点または範囲）は トリガー位置と呼ばれます。

診断は次のプロパティを持ちます：

• 重大度

• トリガールール

• トリガー位置

診断の重大度は以下のいずれかであり、高い順に並んでいます：

error（エラー）

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

warning（警告）

診断は開発者の注意を促す異常を記述しますが、エラーではありません。

info（情報）

診断は開発者の注意を促す注目すべき状態を記述しますが、エラーでも警告でもありません。

off（無効）

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

トリガールールの名前は、以下のいずれかです：

• diagnostic_name_token（単独）

• 2つのdiagnostic_name_tokenをピリオド'.'（U+002E）で区切ったもの

2.3.1. 診断処理

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

1. 各診断Dについて、Dのトリガー位置を含み、かつ同じトリガールールを持つ最小の診断フィルターを探す。

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

   • 存在しなければDは変更されません。

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

3. 残った診断のうち、少なくとも1つDIが重大度infoなら：

   • 同じトリガールールを持つ他のinfo診断は破棄されることがあり、元の診断DIのみが残ります。

4. 残った診断のうち、少なくとも1つDWが重大度warningなら：

   • 同じトリガールールを持つ他のinfoやwarning診断は破棄されることがあり、元の診断DWのみが残ります。

5. 残った診断にerror重大度がある場合：

   • 他の診断（error重大度を含む）も破棄されることがあります。

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

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

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

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

7. パイプライン作成時処理の場合、error診断はWebGPUの検証失敗につながります（GPUProgrammableStageの検証時）。

注意: これらの規則により、実装はエラーを検出した時点でWGSLモジュールの処理を打ち切ることができます。 また、ある種の警告分析は最初の警告で打ち切ることができ、ある種の情報診断分析も最初の発生で打ち切ることができます。 WGSLはさまざまな分析手順の順序や、単一分析内の順序を規定しません。 したがって、同じWGSLモジュールでも、異なる実装は同じ重大度の診断のインスタンスを異なる数だけ報告する場合があります。

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

ほとんどの診断は無条件でWebGPUアプリケーションに報告されます。 一部の診断は、フィルタ可能であり、トリガールールの名前付けにより部分的に制御できます。 以下の表はフィルタ可能な標準的トリガールールを示します。

1つのdiagnostic 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に含まれること

• 他にDF'(AR',NS',TR)がありLがAR'に含まれる場合、ARはAR'に含まれること

影響範囲が入れ子になるため、最も近い囲み診断：

• 唯一の範囲診断フィルターとなる、

• または重複した（競合しない）グローバル診断フィルターのひとつ、

• または存在しない。

2.4. 制限

WGSLの実装は以下の制限を満たすシェーダーをサポートするwill必要があります。 WGSLの実装は、指定された制限を超えるシェーダーもサポートしてよいです。

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

3. テキスト構造

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

WGSLモジュールはUTF-8でエンコードされたUnicodeテキストで、バイトオーダーマーク（BOM）はありません。

WGSLモジュールのテキストは、Unicodeのコードポイントの並びであり、連続した空でない集合にグループ化されて以下を形成します：

• コメント

• トークン

• 空白

プログラムテキストはヌルコードポイント（U+0000）を含んではなりません。

3.1. 構文解析

WGSLモジュールを構文解析する手順：

1. コメントを削除：

   • 最初のコメントをスペースコードポイント（U+0020）に置換。

   • コメントが残っている限り繰り返す。

2. テンプレートリストを、 § 3.9 テンプレートリストのアルゴリズムで検出。

3. 全文を、translation_unit構文規則にマッチさせて解析する。 解析にはLALR(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 Section 6.1 Non-tailorable Line Breaking Rules、LB4、LB5で規定される「必須改行」です。 すなわち、改行となるのは以下のいずれかです：

• 改行（U+000A）

• 垂直タブ（U+000B）

• フォームフィード（U+000C）

• 復帰（U+000D）で、その後に改行（U+000A）が続かない場合

• 復帰（U+000D）の直後に改行（U+000A）が続く場合

• 次行（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サフィックス

注意: 非ゼロ整数リテラルの先頭ゼロ（例：012）は禁止です。他言語の先頭ゼロ＝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つの論理部分を持ちます。 リテラルの値は概ね、仮数部を基数^指数の値で乗算したものです。 仮数部の数字は、ゼロでない場合、または左と右にゼロでない数字が両方ある場合、有効桁とみなされます。 有効桁は左から右へ数えます。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;

注意: 小数部は20桁以降切り捨て、約log(10)/log(2)×20 ≈ 66.4ビットの有効精度が保持されます。

16進浮動小数点リテラルは以下の通り：

• 0xまたは0Xプリフィックス

• その後、16進数字の並びで、どこかに16進小数点（.）があってもよい 仮数部は16進表記の分数を表します。

• その後、オプションで指数部（以下の並び）：

  • pまたはP

  • 符号（+または-）を先頭に付けてもよい10進数の指数

  • その後、オプションでfまたはhサフィックス

• 16進小数点または指数のいずれかが必須です。 どちらもない場合、そのトークンは整数リテラルです。

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桁以降切り捨て、約4×16=64ビットの有効精度が保持されます。

数値リテラルにサフィックスがある場合、そのリテラルは特定の具象スカラー型の値を表します。 サフィックスがない場合、リテラルは下記の抽象数値型のいずれかを表します。 いずれの場合も、リテラルの値は型変換後の数学的値であり、 § 15.7.6 浮動小数点変換の規則に従います。

次の場合はシェーダー作成エラーとなります：

• 整数リテラルにiやuサフィックスが付き、対象型で表せない場合

• 16進浮動小数点リテラルにfやhサフィックスが付き、型であふれるか、正確に表現できない場合

• 10進浮動小数点リテラルにfやhサフィックスが付き、型であふれる場合

• 浮動小数点リテラルにhサフィックスが付き、f16拡張が有効化されていない場合

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

注意: 16進floatリテラルを型指定したい場合、fサフィックスを使うなら2進指数も必要です。例：0x1p0f。0x1fは16進整数リテラルです。

3.6. キーワード

キーワードは、あらかじめ定義された言語概念を指すトークンです。 WGSLキーワード一覧は§ 16.1 キーワード一覧を参照してください。

3.7. 識別子

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

WGSLは用途ごとに2つの文法非終端記号を使い分けています：

• identは宣言されたオブジェクトの名前に使われます。

• member_identは構造体型のメンバーの名前に使われます。

識別子の形式は、Unicode Standard Annex #31 (Unicode Version 14.0.0)に基づき、 以下の補足があります。

識別子はUAX31文法で次のプロファイルを使います：

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

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

つまり、非ASCIIコードポイントを含む識別子も有効です：Δέλτα、réflexion、Кызыл、𐰓𐰏𐰇、朝焼け、سلام、検定、שָׁלוֹם、गुलाबी、փիրուզなど。

ただし以下は例外です：

• 識別子はキーワードや予約語と同じ綴りを持ってはなりません。

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

• 識別子は__（アンダースコア2文字、U+005F×2）で始まってはなりません。

Unicode Character Database for Unicode Version 14.0.0には XID_Startおよび XID_Continueの全有効コードポイントが非規範的に列挙されています。

注意: 一部の組み込み関数の戻り値型は、名前がWGSLソースで使用できない構造体型です。 それらの構造体型は、名前がアンダースコア2文字で始まるものとして事前宣言されたものとみなされます。 結果値は新しく宣言したletやvarに型推論で保存するか、そのメンバーを直接名前指定で抽出できます。 使用例は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'

3.8.2. 組み込み値名

組み込み値名トークンは、組み込み値の名前に使われるトークンです。

§ 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'

• 'workgroup_id'

• 'num_workgroups'

• 'subgroup_invocation_id'

• 'subgroup_size'

3.8.3. 診断ルール名

診断名トークンは、診断トリガールールの名前に使われるトークンです。

§ 2.3 診断参照。

あらかじめ定義された診断ルール名：

• 'derivative_uniformity'

• 'subgroup_uniformity'

3.8.4. 診断重大度制御名

有効な診断フィルター重大度制御名は§ 2.3 診断に記載されていますが、形式は識別子と同じです：

診断フィルター重大度制御名：

• 'error'

• 'warning'

• 'info'

• 'off'

3.8.5. 拡張名

有効な拡張有効化名は§ 4.1.1 拡張の有効化に記載されていますが、一般的には識別子と同じ形式です：

拡張有効化名：

• 'f16'

• 'clip_distances'

• 'dual_source_blending'

• 'subgroups'

• 'primitive_index'

有効な言語拡張名は§ 4.1.2 言語拡張に記載されていますが、一般的には識別子と同じ形式です：

言語拡張名：

• 'readonly_and_readwrite_storage_textures'

• 'packed_4x8_integer_dot_product'

• 'unrestricted_pointer_parameters'

• 'pointer_composite_access'

3.8.6. 補間型名

補間型名トークンは、補間型の名前に使われるトークンです。 interpolate_type_name用。

§ 13.3.1.4 補間参照。

補間型名：

• 'perspective'

• 'linear'

• 'flat'

3.8.7. 補間サンプリング名

補間サンプリング名トークンは、補間サンプリングの名前に使われるトークンです。

§ 13.3.1.4 補間参照。

補間サンプリング名：

• 'center'

• 'centroid'

• 'sample'

• 'first'

• 'either'

3.8.8. スウィズル名

スウィズル名はベクトルアクセス式で使用されます：

3.9. テンプレートリスト

テンプレートパラメータ化は、一般的な概念を修飾するためのパラメータを指定する方法です。 テンプレートパラメータ化を書くには、一般概念の後にテンプレートリストを書きます。

コメントや空白を無視すると、テンプレートリストは：

• 最初に'<'（U+003C）コードポイント、その後

• カンマ区切りで1つ以上のテンプレートパラメータ、その後

• オプションの末尾カンマ、その後

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

テンプレートパラメータの形式は、下記のテンプレートリスト発見アルゴリズムで暗黙的に定義されます。 一般的には名前、式、または型です。

注: 例えば、vec3<f32>はテンプレートパラメータ化で、vec3が修飾される一般概念、 <f32>がテンプレートリスト（パラメータは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）が現れるのは比較演算（ '<=', '>=', '==', '!=' ）の場合のみ。他は代入の一部。

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をさらに進める。

    • '<'（U+003C）がCurrentPositionに現れたら：

      • 注: このコードポイントはテンプレートリスト開始の候補。 後で終端'>'（U+003E）と対応付けできるよう状態を保存。

      • UnclosedCandidate(position=CurrentPosition,depth=NestingDepth) をPendingスタックに積む。

      • CurrentPositionを次のコードポイントに進める。

      • '<'（U+003C）がCurrentPositionに現れたら：

        • 注: 前提1より、いかなるテンプレートパラメータも'<'（U+003C）で始まらないので、直前のコードポイントはテンプレートリスト開始ではない。 よって現在・前のコードポイントは'<<'演算子。

        • Pendingスタックからトップをpop。

        • CurrentPositionを進めて次のループへ。

      • '='（U+003D）がCurrentPositionに現れたら：

        • 注: 前提1より、いかなるテンプレートパラメータも'='（U+003C）で始まらないので、直前のコードポイントはテンプレートリスト開始ではない。 よって現在・前のコードポイントは'<='比較演算子。 '='（U+003D）を後の代入と誤認しないようスキップ。

        • Pendingスタックからトップをpop。

        • CurrentPositionを進めて次のループへ。

    • 次のループへ。

  • '>'（U+003E）がCurrentPositionに現れたら：

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

    • Pendingが空でなければトップTを取得し、T.depthがNestingDepthと等しければ：

      • 注: このコードポイントはT記録のテンプレートリストの終端。

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

      • PendingスタックからTをpop。

      • CurrentPositionを進めて次のループへ。

    • それ以外の場合、このコードポイントはテンプレートリスト終端ではない：

      • CurrentPositionを進める。

      • '='（U+003D）がCurrentPositionに現れたら：

        • 注: 現在・前のコードポイントは'>='比較演算子。 '='（U+003D）を後の代入と誤認しないようスキップ。

        • CurrentPositionを進める。

      • 次のループへ。

  • '('（U+0028）または'['（U+005B）がCurrentPositionに現れたら：

    • 注: 入れ子式の開始。

    • NestingDepthに1加算。

    • CurrentPositionを進めて次のループへ。

  • ')'（U+0029）または']'（U+005D）がCurrentPositionに現れたら：

    • 注: 入れ子式の終了。

    • Pendingスタックから、空になるかトップのdepthがNestingDepthより小さいものになるまでpop。

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

    • CurrentPositionを進めて次のループへ。

  • '!'（U+0021）がCurrentPositionに現れたら：

    • CurrentPositionを進める。

    • '='（U+003D）がCurrentPositionに現れたら：

      • 注: 現在・前のコードポイントは'!='比較演算子。 '='（U+003D）を後の代入と誤認しないようスキップ。

      • CurrentPositionを進める。

    • 次のループへ。

  • '='（U+003D）がCurrentPositionに現れたら：

    • CurrentPositionを進める。

    • '='（U+003D）がCurrentPositionに現れたら：

      • 注: 現在・前のコードポイントは'=='比較演算子。 '='（U+003D）を後の代入と誤認しないようスキップ。

      • CurrentPositionを進めて次のループへ。

    • 注: これは代入の一部で、式中やテンプレートリスト中では現れない。 未閉候補をクリア。

    • NestingDepthを0に設定。

    • Pendingスタックを空に。

    • CurrentPositionを進めて次のループへ。

  • ';'（U+003B）、'{'（U+007B）、':'（U+003A）がCurrentPositionに現れたら：

    • 注: これらは式中やテンプレートリスト中では現れない。未閉候補をクリア。

    • NestingDepthを0に設定。

    • Pendingスタックを空に。

    • CurrentPositionを進めて次のループへ。

  • '&&'または'||'がCurrentPositionのテキストにマッチした場合：

    • 注: これらは比較より優先度の低い演算子。現在の式レベルの未閉候補を破棄。

    • 注: この規則により、a<b || c>dのような断片にテンプレートリストは見つからず、2つの比較の短絡ORと認識される。

    • Pendingスタックから、空になるかトップのdepthがNestingDepthより小さいものになるまでpop。

    • 2コードポイント分CurrentPositionを進めて次のループへ。

  • 現在のコードポイント分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）までがテンプレートリストで、引数はB>=Cです。 B後の'>'（U+003C）コードポイント（誤記？）の後、'='（U+003D）は代入と誤認しないよう特別扱いします。

注: A<(B!=C)>では、最初の'<'（U+003C）から最後の'>'（U+003E）までがテンプレートリストで、引数はB!=Cです。 B後の'!'（U+0021）コードポイントの後、'='（U+003D）は代入と誤認しないよう特別扱いします。

注: A<(B==C)>では、最初の'<'（U+003C）から最後の'>'（U+003E）までがテンプレートリストで、引数はB==Cです。 B後の最初の'='（U+003D）コードポイント、次の'='（U+003D）コードポイントもどちらも代入と誤認しないよう特別扱いします。

テンプレートリスト発見完了後、 構文解析は各テンプレートリストをtemplate_list構文規則にマッチさせます。

4. ディレクティブ

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

ディレクティブは省略可能です。 指定されている場合は、すべてのディレクティブが宣言やconstアサーションよりも前に現れる必要があります。

4.1. 拡張

WGSLは今後も進化することが期待されています。

拡張は、WGSL仕様の一貫性のある変更をまとめた名前付きグループであり、以下の任意の組み合わせで構成されます：

• 新しい構文による新しい概念や動作の追加、例：

  • 宣言、文、属性、組み込み関数など。

• 現行標準や既存拡張での制限の除去。

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

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

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

仮に拡張は以下のようなことが可能です：

• 異なるビット幅の整数などの数値スカラー型の追加。

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

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

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

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

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

• 新しいシェーダーステージの追加。

拡張には2種類あります：有効化拡張と言語拡張です。

4.1.1. 有効化拡張

有効化拡張は、以下の条件を満たす場合のみ利用可能な拡張です：

• 実装がそれをサポートしていること

• シェーダーがenableディレクティブで明示的に要求していること

• 対応するWebGPU GPUFeatureName がGPUDevice 作成時に要求された必須機能のひとつであること

有効化拡張は、ハードウェア固有の機能を公開するために意図されています。

enableディレクティブは、ひとつ以上の有効化拡張を有効にするディレクティブです。 実装が全ての指定された有効化拡張をサポートしていない場合、シェーダー生成エラーとなります。

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

有効な有効化拡張は次の表に示されています。

// 任意精度浮動小数点型の仮想拡張を有効化
enable arbitrary_precision_float;
enable arbitrary_precision_float; // 冗長なenableディレクティブも許容されます。

// 丸めモード制御用の仮想拡張を有効化
enable rounding_mode;

// 仮にarbitrary_precision_float拡張により以下が利用可能になる：
//    - 型f<E,M>
//    - 関数の戻り値、仮引数、let宣言の型として
//    - AbstractFloatからの値コンストラクター
//    - 除算演算子: / のオペランドとして
// 仮にrounding_mode拡張で@rounding_mode属性が利用可能になる。
@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ループは言語拡張となり得ます。

wgslLanguageFeatures メンバーはWebGPU GPU オブジェクトにおいて、 実装がサポートする言語拡張のセットを一覧します。

requiresディレクティブ は、プログラムが利用する言語拡張を記述するディレクティブです。 実装が提供する機能を変更するものではありません。 実装が必要とされる拡張のいずれかをサポートしていない場合、シェーダー生成エラーとなります。

WGSLモジュールはrequiresディレクティブを使って、非移植性の可能性や、 意図した移植性の最低限の基準を示すことができます。

注: WebGPU実装外のツールは、プログラムで利用されるすべての言語拡張がプログラム内のrequiresディレクティブで網羅されているか確認できます。

他のディレクティブ同様、requiresディレクティブが存在する場合は、すべての宣言やconstアサーションよりも前に現れる必要があります。 拡張名は識別子ではありません。 宣言に解決されません。

注: 将来的には、WGSLは現時点で広くサポートされている言語拡張の機能すべてを含める言語拡張を定義する意図があります。 requiresディレクティブにおいて、これらは全ての共通機能をリストアップするための略記となります。 機能の集合は徐々に増加し、ある種の言語バージョンと考えることもできます。

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ソース内で宣言によって作成されます。

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

型の概念とWGSLでその型を示す構文は区別されます。 多くの場合、この仕様における型の綴りはWGSL構文と同じです。 例えば：

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

• 構造体型や構造体を含む型の場合は綴りが異なります。

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

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

6.1. 型検査

WGSL値は式の評価によって算出されます。 式とは、WGSL文法規則のうち「expression」で終わる名前を持つものとして解析されるソーステキストの区間です。 式Eは、外側の式Eに真に内包された式である部分式を含むことがあります。 トップレベル式は、それ自身が部分式でない式です。 § 8.18 式文法のまとめを参照。

式の評価で生成される値は次に依存します：

• 静的コンテキスト：式を囲むソーステキスト

• 動的コンテキスト：式を評価している呼び出しの状態や実行コンテキスト

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

型アサーションは、WGSLソース式とWGSL型の対応付けです。 次の記法

e : T

は「eの静的型がTである」という型アサーションです。

注: 型アサーションはプログラムテキストについての事実の記述です。実行時検査ではありません。

文はしばしば式を用い、その式の静的型に要件を課す場合があります。例：

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

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

型検査とは、正しく構文解析されたWGSLモジュールに対して、 各式を静的型に対応付け、各文の型要件が満たされていることを検証する過程です。 型検査に失敗すると、シェーダー生成エラーの一種である型エラーとなります。

型検査は、型規則を構文的フレーズに再帰的に適用することによって行えます。 構文的フレーズは、式または文です。 型規則は、 構文的フレーズの静的コンテキストが、 その句に含まれる式の静的型をどのように決定するかを記述します。 型規則は2つの部分を持ちます：

• 結論

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

  • 句が文の場合、結論はその文の各トップレベル式ごとの型アサーションの集合です。

  • いずれの場合も、構文的フレーズは、部分式や他の構文的パラメータを示すためにイタリック体の名前で模式的に指定されます。

• 前提条件（以下を含む）：

  • 式の場合：

    • 部分式があれば、その部分式の型アサーション。 これは直接満たしてもよいし、実現可能な自動変換によって満たしてもよい（§ 6.1.2 変換ランク参照）。 WGSLは定数式をすべて評価し、プログラムの静的型を決定します。

    • その式が文でどのように使われるか。

  • 文の場合：

    • 文の構文形式、および

    • 文中のトップレベル式の型アサーション。

  • その他の模式パラメータに対する条件（もしあれば）。

  • オプションでその他の静的コンテキスト。

型規則は、前提条件や結論に型パラメータを持つことがあります。 型規則の結論や前提条件に型パラメータが含まれる場合、その規則はパラメータ化されているといいます。 含まれない場合は完全展開されているといいます。 パラメータ化された型規則に対し、各型パラメータに型を代入することで完全展開された型規則を作ることができます。 規則の型パラメータへの型の割り当てを代入と呼びます。

例えば、論理否定（!e形式の式）の型規則は次のとおりです：

これは型パラメータTを含むためパラメータ化された規則です。 Tは4つの型、bool、vec2<bool>、vec3<bool>、vec4<bool>のいずれかになり得ます。 例えばTにvec3<bool>を割り当てると、完全展開された型規則は次のようになります：

パラメータ化された規則に対し、条件を満たす代入によって得られる完全展開規則をそれぞれオーバーロードと呼びます。 例えば、この論理否定規則は型パラメータTの割り当て方法が4通りなので、オーバーロードは4つあります。

注: つまり、パラメータ化された型規則は、異なる代入によって生成される完全展開型規則の集合のパターンを提供します。

型規則が構文的フレーズに適用されるのは、次の場合です：

• 規則の結論が構文的フレーズの有効な構文解析に一致し、

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

パラメータ化された型規則は、ある式に対し、 代入によって完全展開型規則を生成でき、 その型規則がその式に適用できる場合に適用されます。

例えば、式1u+2uを考えます。 これは2つのリテラル部分式（1uと2u、どちらもu32型）を持ちます。 トップレベル式は加算です。 § 8.7 算術式の規則を参照すると、加算の型規則がこの式に適用されます。なぜなら：

• 1u+2uはe1+e2という形式にパースされ、e1が1u、e2が2uとなり、

• e1はu32型であり、

• e2もu32型であり、

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

構文的フレーズを解析する際、以下の3つの場合があります：

• どの型規則もその式に適用できない→型エラーとなる。

• 完全展開された型規則がちょうど1つだけその式に適用できる→その規則の結論が主張され、式の静的型が決定される。

• 複数の型規則が適用できる（複数のオーバーロードが前提条件を満たす）→§ 6.1.3 オーバーロード解決で競合解決を行う。

  • オーバーロード解決が成功した場合、その式に適用されるオーバーロードが1つに決定される。 そのオーバーロードの型アサーションが主張され、 構文的フレーズ内の式の型が決定される。

  • オーバーロード解決が失敗した場合、型エラーとなる。

上記の例の通り、式1u+2uには型規則が一つだけ適用できるため、型検査はその型規則の結論（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つ目は定数式をGPUで利用できる型付き数値値に変換する場合。 2つ目はメモリ参照からロードが発生し、そのメモリに格納された値が得られる場合です。

注: ランクが無限大の変換は実現不可能、つまり許可されません。

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

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

• Tが具象型であり

• Tが参照型でない

• ConversionRank(S, T)が有限である

• 他の非参照型T2について、ConversionRank(S, T2) > ConversionRank(S, T)

型Tの値eの具象化は、 TからTの具象化への実現可能な変換をeに適用した結果です。

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

6.1.3. オーバーロード解決

複数の型規則が構文的フレーズに適用できる場合、 どれを適用するかを決定するための競合解決手順が用いられます。 この手順をオーバーロード解決と呼びます。 型検査によって部分式の静的型がすでに決定されていることを前提とします。

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

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

• その結論は：

  • 構文形式がPの有効な構文解析に一致し、

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

Pのオーバーロード解決は以下の手順で進み、最も望ましいオーバーロード候補を1つ見つけることを目的とします：

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

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

   • その式の静的型から前提条件の型アサーションが要求する型への実現可能な自動変換が存在する。 C.R(i)をその変換のConversionRankとする。

2. 部分式のいずれかが自動変換後に抽象型になるが、 その候補の他の部分式が定数式でない場合、その候補を除外する。

   注: 結果として、フレーズ内のいずれかの部分式が定数式でない場合、フレーズ内のすべての部分式は具象型でなければなりません。

3. 候補を順位付けする：2つのオーバーロード候補C1とC2について、C1が優先されるのは：

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

     • つまり、C1をPに適用するために必要な各式変換が、C2に必要な対応する式変換と同等以上に望ましい。

   • 少なくとも1つの式位置iで、C1.R(i) < C2.R(i)

     • つまり、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)はlog2組み込み関数への関数呼び出しと解釈され、オペランドはAbstractInt値32です。

• 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→構築可能型の自動変換で最優先（変換ランク2）はAbstractInt→i32。よって'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);

// 次の例はすべてf32に昇格し、初期値は10f
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。よって'1'はi32型になる。
let some_i32 = 1; // 例：let some_i32: i32 = 1i;

let some_f32 : f32 = some_i32; // 型エラー：i32はf32に代入不可

// もう1つのオーバーフロー例
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ビット浮動小数点値の集合です。シェーダー生成エラー となるのは、プログラムに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です。

スカラー変換は、あるスカラー型の値を別のスカラー型の値へ写像します。 一般に、変換後の値は変換前の値に近く、変換先型の制限内で表現されます。 スカラー変換は次の場合に発生します：

• 値コンストラクター組み込み関数を明示的に呼び出す場合

• 定数式（抽象数値型）を他の型に変換する場合、 自動変換を介して行う。

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.7 算術式を参照してください。

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

WGSLはまた、以下の事前宣言型エイリアスも宣言します：

6.2.8. アトミック型

アトミック型は、 具象整数スカラー型をカプセル化し、以下の特徴を持ちます：

• アトミックオブジェクトはある保証を並行する観測者に提供する

• アトミックオブジェクトに対して有効な操作はatomic組み込み関数のみ

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

アトミック型は、workgroup アドレス空間の変数、またはstorage buffer変数（read_writeアクセスモード）でのみインスタンス化可能です。 型への操作のメモリスコープは、その型がインスタンス化されるアドレス空間で決まります。 workgroup空間のアトミック型はWorkgroupスコープ、 storage空間のアトミック型はQueueFamilyスコープです。

アトミック変更とは、 アトミックオブジェクトに対して値を書き換える操作のことです。 新しい値が既存値と同じ場合でも「変更」とみなします。

WGSLでは、アトミック変更は各オブジェクトごとに相互順序付けされます。 つまり、シェーダーステージの実行中、アトミックオブジェクトAごとに、 すべてのエージェントがAに適用された変更操作の順序を同じように観測します。 別個のアトミックオブジェクト間の順序は関連付けられません。因果関係も暗示されません。 workgroup空間の変数は同一workgroup内で共有されますが、 異なるworkgroup間では共有されません。

6.2.9. 配列型

配列は要素値をインデックスでアクセスできる並びです。

配列の最初の要素はインデックス0、以降は整数インデックスでアクセスします。 § 8.5.3 配列アクセス式参照。

式が実行時サイズ配列型に評価されてはなりません。

固定長配列の要素数式Nには以下の制約があります：

• 必ずoverride式であること

• 必ず具象整数スカラーに評価されること

• Nが0以下の場合：

  • Nが定数式ならシェーダー生成エラー

  • それ以外の場合はパイプライン生成エラー

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

注: 型等価性判定のため、定数式でないoverride式は識別子でなければなりません。 オーバーライド可能定数でサイズ指定したworkgroup変数例参照。

実行時サイズ配列の要素数は、対応するstorage buffer変数に関連付けられたバッファバインディングのサイズで決まります。 § 13.3.4 バッファバインディングによる実行時サイズ配列の要素数決定参照。

配列要素型は、以下のいずれかでなければなりません：

• スカラー型

• ベクトル型

• 行列型

• アトミック型

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

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

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

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つの構造体型は名前が同じ場合のみ同じ型です。

構造体メンバー型は、以下のいずれかでなければなりません：

• スカラー型

• ベクトル型

• 行列型

• アトミック型

• 固定長配列型（生成時固定フットプリント付き）

• 実行時サイズ配列型（ただし構造体の最後のメンバーのみ）

• 構造体型（生成時固定フットプリント付き）

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

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

構造体メンバー型および配列要素型の制約により生じる結果：

• ポインター、テクスチャ、サンプラーは、いかなる入れ子レベルでも配列や構造体内に現れてはなりません。

• 実行時サイズ配列が複合型の一部の場合、構造体の最後の要素としてのみ現れ、 その構造体自体はさらに配列や構造体に含まれてはなりません。

// 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はレイアウト属性であり、 構造体型がユニフォームバッファやストレージバッファを定義する際に必要となる場合があります。 § 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について、入れ子深度（NestDepth(T)）は：

• ベクトル型なら1

• 行列型なら2

• 要素型Eの配列型なら 1 + NestDepth(E)

• メンバー型M₁,...,M_Nを持つ構造体型Tなら 1 + max(NestDepth(M₁),..., NestDepth(M_N))

6.2.12. 構築可能型

多くの種類の値は生成、ロード、格納、関数への渡し、関数からの返却が可能です。 これらを構築可能と呼びます。

型が次のいずれかであれば、構築可能です：

• スカラー型

• ベクトル型

• 行列型

• 固定長配列型（生成時固定フットプリントがあり、要素型が構築可能な場合）

• すべてのメンバーが構築可能な構造体型

注: すべての構築可能型は生成時固定フットプリントを持ちます。

注: アトミック型や実行時サイズ配列型は構築可能ではありません。 アトミックや実行時サイズ配列を含む複合型も構築可能ではありません。

6.2.13. 固定フットプリント型

変数のメモリフットプリントは、その変数の内容を格納するために使われるメモリ位置の数です。 メモリフットプリントは変数の格納型に依存し、シェーダーライフサイクルのある段階で確定します。 ほとんどの変数は、シェーダー生成時にサイズが決まります。 一部の変数は、パイプライン生成時、 またはシェーダー実行開始時に決まることもあります。

型が生成時固定フットプリントを持つとは、その具象化のサイズがシェーダ生成時に完全に決定される場合を指します。

型が固定フットプリントを持つとは、そのサイズがパイプライン生成時に完全に決まる場合です。

注: すべての具象生成時固定フットプリント型・固定フットプリント型は格納可能型です。

注: パイプライン生成はシェーダー生成に依存するので、生成時固定フットプリントを持つ型は、必ず固定フットプリントも持ちます。

生成時固定フットプリントを持つ型は：

• スカラー型

• ベクトル型

• 行列型

• アトミック型

• 固定長配列型（要素数が定数式の場合）

• すべてのメンバーが生成時固定フットプリントを持つ構造体型

注: 構築可能型は生成時固定フットプリントを持ちます。

プレーン型で固定フットプリントを持つ型は：

• 生成時固定フットプリントを持つ型

• 固定長配列型（要素数への追加制約なし）

注: 定数式でないoverride式による要素数指定の固定長配列型が有効なのは、 workgroupアドレス空間でのメモリビューのみです。 これはworkgroup変数の格納型も含みます。

注: 固定フットプリント型は直接または間接的にアトミック型を含み得ますが、構築可能型は含みません。

注: 固定フットプリント型には実行時サイズ配列、およびそれを含む構造体は含まれません。

6.3. 列挙型

列挙型は、限定された名前付き値の集合です。 列挙型は、例えば有効なテクセルフォーマットの集合など、特定の概念の可能性の集合を区別するために使われます。

列挙値は、列挙型内の名前付き値のひとつです。 各列挙値は他のすべての列挙値、および他の種類の値と区別されます。

WGSLソースで新しい列挙値や列挙型を宣言する方法はありません。

注: 列挙値はテンプレートパラメータとして使われます。

6.3.1. 事前宣言列挙値

次の表はWGSLの列挙型と、その事前宣言列挙値を示します。 列挙型自体はWGSLソースで記述できません。

6.4. メモリビュー

WGSLプログラムはプレーン値の計算に加え、 メモリから値を読み書きすることも多く、これはメモリアクセス操作によって行われます。 各メモリアクセスはメモリビューを介して実行されます。

メモリビューは以下から構成されます：

• 特定のアドレス空間内のメモリ位置の集合

• メモリモデル参照

• それらメモリ位置の内容をWGSL型として解釈したもの（格納型）、

• アクセスモード

メモリビューのアクセスモードは、アドレス空間でサポートされていなければなりません。§ 7 変数と値の宣言参照。

6.4.1. 格納可能型

変数に格納される値は必ず 格納可能型でなければなりません。 格納可能型はWGSLで明示的な表現が定義されている場合もあり（§ 14.4.4 値の内部レイアウト参照）、 テクスチャやサンプラーのように不透明な場合もあります。

型が具象型で、以下のいずれかであれば格納可能型です：

• スカラー型

• ベクトル型

• 行列型

• アトミック型

• 配列型

• 構造体型

• テクスチャ型

• サンプラー型

注: つまり、格納可能型は具象プレーン型、テクスチャ型、サンプラー型です。

6.4.2. ホスト共有型

ホスト共有型は、ホストとGPU間で共有されるバッファの内容を記述する型や、フォーマット変換せずにホストとGPU間でコピーされる型です。 この目的で使う場合、型にはレイアウト属性が適用されることがあります（§ 14.4 メモリレイアウト参照）。 § 7.3 var宣言の通り、ユニフォームバッファやストレージバッファ 変数の格納型は必ずホスト共有型でなければなりません。

型が具象型で、以下のいずれかであればホスト共有型です：

• 数値スカラー型

• 数値ベクトル型

• 行列型

• アトミック型

• 固定長配列型（生成時固定フットプリントがあり、要素型がホスト共有型の場合）

• 実行時サイズ配列型（要素型がホスト共有型の場合）

• すべてのメンバーがホスト共有型の構造体型

注: ステージ間入出力型の制約は§ 13.3.1 ステージ間入出力インターフェースおよびそれ以降の節に記載されています。 それらの型もサイズが決まりますが、カウント方法が異なります。

注: テクスチャやサンプラーもホストとGPU間で共有できますが、その内容は不透明です。 この節のホスト共有型は、特にストレージおよびユニフォームバッファで使います。

6.4.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.4.4. 有効・無効なメモリ参照

参照値は有効または無効です。

参照は§ 6.4.8 参照とポインター値の形成で詳しく説明されているように形成されます。 一般に、有効な参照は次の方法で形成されます：

• 変数を名前で指定する

• 間接参照（単項*）演算を有効なポインターに適用する

• 名前付き構成要素式で、ベースが有効なメモリビューである場合

• 添字式で、ベースが有効なメモリビューであり、かつ範囲内インデックスを使う場合

一般に、無効なメモリ参照は次の方法で形成されます：

• 間接参照演算子を無効なポインターに適用する

• 名前付き構成要素式で、ベースが無効なメモリ参照である場合

• 添字式で、ベースがメモリビューであり、かつ：

  • ベースが無効なメモリ参照である場合

  • インデックスが範囲外の場合

有効なポインターは有効な参照に対応するポインターです。 無効なポインターは無効なメモリ参照に対応するポインターです。

6.4.5. 起源変数

ポインター値の起源変数は、対応する参照値の起源変数として定義されます。

注: 起源変数は動的な概念です。 関数の仮引数の起源変数は、その関数の呼び出し元によります。 呼び出し元が異なれば、異なる起源変数へのポインターが渡されることがあります。

有効な参照は必ず、ある変数のメモリ位置の一部または全部に対応する非空のメモリビューとなります。

struct Particle {
   position: vec2f,
   velocity: vec2f,
   color_index: i32,
}

@group(0) @binding(0)
var<storage,read_write> the_particle: Particle;

fn particle_velocity_component(p: Particle, i: i32) -> f32 {
  return the_particle.velocity[i]; // iが0または1のとき有効な参照
}

6.4.6. 範囲外アクセス

メモリアクセス操作が無効なメモリ参照をアクセスすると、範囲外アクセスとなります。

範囲外アクセスはプログラム不具合です。もしそのまま実行された場合、一般的には：

• 変数の外側のメモリ位置を読み書きする

• それらの位置の内容を誤った格納型として解釈する

• 意図しないデータ競合を発生させる

このため、実装は決してそのままアクセスを実行しません。 範囲外アクセスの実行は動的エラーとなります。

注: 格納型の誤解釈例は 前節の例で発生します。 iが2の場合、the_particle.velocity[i]式は 型ref<storage,f32,read_write>（格納型はf32）ですが、 実際のメモリ位置はcolor_indexメンバー用に割り当てられているため、値はi32型です。

6.4.7. 参照型とポインター型の利用例

参照型とポインター型は用途で区別されます：

• 変数の型は参照型

• アドレス演算（単項&）は参照値を対応するポインター値に変換する

• 間接参照（単項*）はポインター値を対応する参照値に変換する

• let宣言はポインター型にできるが参照型にはできない

• 関数仮引数はポインター型にできるが参照型にはできない

• 単純代入文は、参照型経由でメモリ内容を更新する書き込みアクセスを行う。具体的には：

  • 代入文の左辺は参照型で、アクセスモードはwriteまたはread_writeでなければならない

  • 代入文の右辺は、左辺の格納型に評価されなければならない

• ロード規則：関数内では参照型は型規則を満たすため自動的にデリファレンスされる（値が読み出される）：

  • 関数内で、格納型Tの参照式rが文や式で使われ、

  • rのアクセスモードがreadまたはread_writeで、

  • 一致する型規則がrに型Tの値を要求する場合、

  • その型規則の要求は満たされたものとみなされる

  • その状況でrを評価すると、参照先メモリ位置の値（型T）が得られる。つまり、読み出しアクセスが行われる。

このような参照型定義により、変数の簡便な利用が可能となります：

@compute @workgroup_size(1)
fn main() {
  // 'i'はref<function,i32,read_write>型
  // 'i'のメモリ位置にはi32値0が格納されている。
  var i: i32 = 0;

  // 'i + 1'は'i'部分式がi32型でなければ型規則に一致しない。
  // よって'i + 1'式はi32型。評価時、'i'部分式はその時点でメモリ位置に格納されたi32値となる。
  let one: i32 = i + 1;

  // 'i'のメモリ位置値を2に更新
  i = one + 1;

  // 'i'のメモリ位置値を5に更新。右辺は代入前に評価される。
  i = i + 3;
}

var<private> age: i32;
fn get_age() -> i32 {
  // return文の式の型は'i32'でなければならない（関数の戻り値型に一致）。
  // 'age'式は型ref<private,i32,read_write>。
  // 参照の格納型が要求される型と一致し、他の型規則が適用されないため、Load規則を適用。
  // この文脈での'age'評価はreturn文実行時に参照するメモリ領域からロードされたi32値。
  return age;
}

fn caller() {
  age = 21;
  // copy_age定数にはi32値21が代入される。
  let copy_age: i32 = get_age();
}

ポインター型定義により、2つの主要な用途が可能となります：

• let宣言でポインター型を使い、変数内容の一部の短い名前を作る

• 関数の仮引数にポインター型を使い、呼び出し元関数からアクセス可能な変数のメモリを参照する

  • その関数呼び出しでは、そのオペランドにポインター値を渡す必要があり、通常は&アドレス演算で変数内容へのポインターを得る

struct Particle {
  position: vec3<f32>,
  velocity: vec3<f32>
}
struct System {
  active_index: i32,
  timestep: f32,
  particles: array<Particle,100>
}
@group(0) @binding(0) var<storage,read_write> system: System;

@compute @workgroup_size(1)
fn main() {
  // ストレージメモリ内の特定Particleへのポインターを形成
  let active_particle = &system.particles[system.active_index];

  let delta_position: vec3<f32> = (*active_particle).velocity * system.timestep;
  let current_position: vec3<f32>  = (*active_particle).position;
  (*active_particle).position = delta_position + current_position;
}

fn add_one(x: ptr<function,i32>) {
  /* 'x' の場所に次の大きい整数値を格納するように更新する、
     (もしくは、最大の負の i32 値にラップアラウンドする)。
     左辺では、単項 '*' はポインタを参照に変換し、
     それに代入できるようにする。デフォルトでは read_write アクセスモードを持つ。 
     /* 右辺について:
        - 単項 '*' はポインタを参照に変換し、read_write アクセスモードを持つ。
        - 一致する型ルールは加算 (+) のみで、'*x' が i32 型である必要がある。
          これは '*x' のストア型である。よってロードルールが適用され、
          評価時点で '*x' に格納されている値、すなわち i32 の 0 となる。
        - 0 に 1 を加算し、右辺の最終値として 1 になる。*/
     1 を '*x' のメモリに格納する。*/
  *x = *x + 1;
}

@compute @workgroup_size(1)
fn main() {
  var i: i32 = 0;

  // 'i' の内容を修正し、1 にする。
  // 単項 '&' を使い、'i' のポインタ値を取得する。
  // 呼び出し先の関数が 'i' のメモリにアクセスでき、修正可能であることを明確に示す。
  add_one(&i);
  let one: i32 = i;  // 'one' の値は 1 となる。
}

6.4.8. 参照値とポインター値の生成

参照値は次のいずれかの方法で生成されます：

• 識別子が解決されて、スコープ内の変数vを指す場合、vのメモリへの参照値を表します。

• ポインターに間接参照（単項*）演算を適用する。

• メモリビューを合成型に対して名前付き構成要素式で使う：

  • ストア型がベクトルの場合、1文字のベクトルアクセス句を追加すると、その成分への参照値になります。 § 8.5.1.3 ベクトルメモリビューからの成分参照参照。

  • ストア型が構造体の場合、メンバーアクセス句を追加すると、そのメンバーへの参照値になります。 § 8.5.4 構造体アクセス式参照。

• メモリビューを合成型に対して添字式で使う：

  • ストア型がベクトルの場合、配列インデックスアクセス句を追加すると、その成分への参照値になります。 § 8.5.1.3 ベクトルメモリビューからの成分参照参照。

  • ストア型が行列の場合、配列インデックスアクセス句を追加すると、その列ベクトルへの参照値になります。 § 8.5.2 行列アクセス式参照。

  • ストア型が配列の場合、配列インデックスアクセス句を追加すると、その要素への参照値になります。 § 8.5.3 配列アクセス式参照。

全ての場合で、結果のアクセスモードは元の参照のものと同じです。

struct S {
    age: i32,
    weight: f32
}
var<private> person: S;
// 別の場所で'person'は変数の基底メモリへの参照を表し、型はref<private,S,read_write>となる。

fn f() {
    var uv: vec2<f32>;
    // この関数本体では'uv'はその基底メモリへの参照を表し、型はref<function,vec2<f32>,read_write>となる。

    // 代入の左辺を評価：
    //   'uv.x'を評価して参照値を得る：
    //   1. まず'uv'を評価し、変数'uv'のメモリへの参照（型ref<function,vec2<f32>,read_write>）を得る。
    //   2. '.x'ベクトルアクセス句を適用し、参照値の指すベクトルの第1成分メモリへの参照（型ref<function,f32,read_write>）を得る。
    // 右辺を評価するとf32値1.0となる。1.0をuv.xのストレージメモリ位置に格納。
    uv.x = 1.0;

    // 代入の左辺を評価：
    //   'uv[1]'を評価して参照値を得る：
    //   1. まず'uv'を評価し、変数'uv'のメモリへの参照（型ref<function,vec2<f32>,read_write>）を得る。
    //   2. '[1]'配列インデックス句を適用し、参照値の指すベクトルの第2成分メモリへの参照（型ref<function,f32,read_write>）を得る。
    // 右辺はf32値2.0。2.0をuv[1]のストレージメモリ位置に格納。
    uv[1] = 2.0;

    var m: mat3x2<f32>;
    // 'm[2]'の評価：
    // 1. 'm'を評価し、変数'm'のメモリへの参照（型ref<function,mat3x2<f32>,read_write>）を得る。
    // 2. '[2]'配列インデックス句を適用し、参照値の指す行列の第3列ベクトルメモリへの参照（型ref<function,vec2<f32>,read_write>）を得る。
    // let宣言はvec2<f32>型なので、初期化子はvec2<f32>型でなければならない。
    // ロード規則が適用され、'm[2]'参照位置からロードした値が初期化子となる。
    let p_m_col2: vec2<f32> = m[2];

    var A: array<i32,5>;
    // 'A[4]'の評価：
    // 1. 'A'を評価し、変数'A'のメモリへの参照（型ref<function,array<i32,5>,read_write>）を得る。
    // 2. '[4]'配列インデックス句を適用し、参照値の指す配列の第5要素メモリへの参照（型ref<function,i32,read_write>）を得る。
    // let宣言はi32型なので、右辺はi32型でなければならない。
    // ロード規則が適用され、'A[4]'参照位置からロードした値が初期化子となる。
    let A_4_value: i32 = A[4];

    // 'person.weight'の評価：
    // 1. 'person'を評価し、モジュールスコープ変数のメモリへの参照（型ref<private,S,read_write>）を得る。
    // 2. '.weight'メンバーアクセス句を適用し、参照値の指すメモリの第2メンバー位置への参照（型ref<private,f32,read_write>）を得る。
    // let宣言はf32型なので、右辺はf32型でなければならない。
    // ロード規則が適用され、'person.weight'参照位置からロードした値が初期化子となる。
    let person_weight: f32 = person.weight;

    // また、ポインターからも同じ構文で参照値が生成できる。

    let uv_ptr = &uv;
    // この関数本体では'uv_ptr'は'uv'基底メモリへのポインターで、型はptr<function,vec2<f32>,read_write>となる。

    // 代入左辺を評価：
    //   '*uv_ptr'を評価して参照値を得る：
    //   1. 'uv_ptr'を評価し、'uv'基底メモリへのポインター（型ptr<function,vec2<f32>,read_write>）を得る。
    //   2. 間接参照演算子を適用し、'uv'のメモリへの参照を得る。
    // 右辺はvec2<f32>値(1.0, 2.0)。'uv'のストレージメモリ位置に(1.0, 2.0)を格納。
    *uv_ptr = vec2f(1.0, 2.0);

    // 代入左辺を評価：
    //   'uv_ptr.x'を評価して参照値を得る：
    //   1. 'uv_ptr'を評価し、'uv'基底メモリへのポインター（型ptr<function,vec2<f32>,read_write>）を得る。
    //   2. '.x'ベクトルアクセス句を適用し、参照値の指すベクトルの第1成分メモリへの参照（型ref<function,f32,read_write>）を得る。
    // 右辺はf32値1.0。uv.xのストレージメモリ位置に1.0を格納。
    uv_ptr.x = 1.0;

    // 代入左辺を評価：
    //   'uv_ptr[1]'を評価して参照値を得る：
    //   1. 'uv_ptr'を評価し、'uv'基底メモリへのポインター（型ptr<function,vec2<f32>,read_write>）を得る。
    //   2. '[1]'配列インデックス句を適用し、参照値の指すベクトルの第2成分メモリへの参照（型ref<function,f32,read_write>）を得る。
    // 右辺はf32値2.0。uv[1]のストレージメモリ位置に2.0を格納。
    uv_ptr[1] = 2.0;

    let m_ptr = &m;
    // 'm_ptr[2]'の評価：
    // 1. 'm_ptr'を評価し、'm'基底メモリへのポインター（型ptr<function,mat3x2<f32>,read_write>）を得る。
    // 2. '[2]'配列インデックス句を適用し、参照値の指す行列の第3列ベクトルへの参照（型ref<function,vec2<f32>,read_write>）を得る。
    // let宣言はvec2<f32>型なので、初期化子はvec2<f32>型でなければならない。
    // ロード規則が適用され、'm[2]'参照位置からロードした値が初期化子となる。
    let p_m_col2: vec2<f32> = m_ptr[2];

    let A_ptr = &A;
    // 'A[4]'の評価：
    // 1. 'A_ptr'を評価し、'A'基底メモリへのポインター（型ptr<function,array<i32,5>,read_write>）を得る。
    // 2. '[4]'配列インデックス句を適用し、参照値の指す配列の第5要素メモリへの参照（型ref<function,i32,read_write>）を得る。
    // let宣言はi32型なので、右辺はi32型でなければならない。
    // ロード規則が適用され、'A[4]'参照位置からロードした値が初期化子となる。
    let A_4_value: i32 = A_ptr[4];

    let person_ptr = &person;
    // 'person.weight'の評価：
    // 1. 'person_ptr'を評価し、モジュールスコープ変数のメモリへのポインター（型ptr<private,S,read_write>）を得る。
    // 2. '.weight'メンバーアクセス句を適用し、参照値の指すメモリの第2メンバー位置への参照（型ref<private,f32,read_write>）を得る。
    // let宣言はf32型なので、右辺はf32型でなければならない。
    // ロード規則が適用され、'person.weight'参照位置からロードした値が初期化子となる。
    let person_weight: f32 = person_ptr.weight;
}

ポインター値は次のいずれかの方法で生成されます：

• 参照にアドレス演算（単項&）演算子を適用する。

  • 元の参照が有効の場合に限り、結果は有効なポインターです。

  • 有効な結果の起源変数は元の参照のものです。

• 関数の仮引数がポインター型の場合、関数呼び出し時に、その仮引数の利用は呼び出し元関数内の呼び出し箇所で渡されたポインター値を表します。

  • 実行時、仮引数で表される値が有効なポインターなのは、呼び出し箇所で渡されたポインター値が有効な場合のみです。

  • 有効なポインター仮引数の起源変数は、呼び出し箇所で渡されたポインターオペランドの起源変数です。

全ての場合で、結果のアクセスモードは元のポインターのものと同じです。

// privateアドレス空間にf32値を格納する変数を宣言
var<private> x: f32;

fn f() {
    // functionアドレス空間にi32値を格納する変数を宣言
    var y: i32;

    // 'x'はモジュールスコープ変数'x'に解決され、型はref<private,f32,read_write>となる。
    // 単項'&'演算で参照をポインターに変換。
    // アクセスモードは元変数のものと同じ。完全型はptr<private,f32,read_write>。
    // functionアドレス空間のデフォルトアクセスモードはread_writeなので明記不要。
    let x_ptr: ptr<private,f32> = &x;

    // 'y'は関数スコープ変数'y'に解決され、型はref<private,i32,read_write>となる。
    // 単項'&'演算で参照をポインターに変換。アクセスモードはデフォルトでread_write。
    let y_ptr: ptr<function,i32> = &y;

    // 新しい変数。モジュールスコープ変数と別物。
    var x: u32;

    // ここで'x'は直前に宣言された関数スコープ変数で、型はref<function,u32,read_write>。
    // 単項'&'演算で参照をポインターに変換。アクセスモードはデフォルトでread_write。
    let inner_x_ptr: ptr<function,u32> = &x;
}

6.4.9. 他言語における参照型・ポインター型との比較

この節は参考情報であり、規範的ではありません。

WGSLの参照型・ポインター型は他言語よりも制限されています。特に：

• WGSLでは参照型を別の参照や変数へのエイリアスとして直接宣言できません（変数・仮引数いずれも）。

• WGSLではポインター型・参照型は格納可能型ではありません。 すなわち、WGSLの変数宣言の内容としてポインターや参照を格納できません。

• WGSLでは関数がポインター型や参照型を返してはなりません。

• WGSLでは整数値とポインター値の間で相互変換する方法がありません。

• WGSLではポインター値の型を強制的に別のポインター型に変換する方法がありません。

  • 合成成分参照式は異なります。これは合成値への参照から、合成値内部の構成要素・要素への参照を生成します。 WGSLではこれらは異なる参照とみなされ、実装の抽象度が低い場合は同じ機械アドレスでも異なります。

• WGSLでは参照値の型を強制的に別の参照型に変換する方法がありません。

• WGSLではポインター型や参照型のアクセスモードを変更する方法がありません。

  • 比較として、C++では非constポインターを自動でconstポインターに変換し、const値を非const値に変換するconst_castがあります。

• WGSLでは「ヒープ」から新たなメモリ領域を割り当てる方法がありません。

• WGSLでは変数の明示的な破棄方法がありません。WGSL変数のメモリはスコープを抜けたときのみアクセス不可になります。

注: 上記規則から「ダングリング」ポインター、すなわち「生存」起源変数のメモリを参照しないポインターを生成することはできません。 メモリビューは無効なメモリ参照となる場合がありますが、 決して 起源変数またはバッファに関連しないメモリ位置へアクセスしません。

6.5. テクスチャ型・サンプラー型

テクセルはテクスチャの最小単位として個別にアクセス可能なスカラーまたはベクトルです。 texelはtexture elementの略です。

テクスチャは、レンダリングに有用な特殊な操作をサポートするテクセル集合です。 WGSLではこれらの操作はテクスチャ組み込み関数で呼び出します。 全リストは§ 17.7 テクスチャ組み込み関数参照。

WGSLのテクスチャはWebGPUのGPUTextureに対応します。

テクスチャには以下の特徴があります：

テクセルフォーマット

各テクセルのデータ表現。§ 6.5.1 テクセルフォーマット参照。

次元性

グリッド座標の次元数と座標の解釈方法。次元数は1, 2, 3のいずれかです。 多くのテクスチャは直交座標を使います。 キューブテクスチャは6つの正方形面を持ち、三次元座標は原点からキューブ中心への方向ベクトルとして解釈されます。

WebGPUのGPUTextureViewDimension参照。

サイズ

各次元のグリッド座標の範囲。ミップレベルによって決まります。

ミップレベル数

ミップレベル数はサンプルテクスチャや深度テクスチャでは1以上、ストレージテクスチャでは1です。
ミップレベル0がテクスチャのフルサイズ版を保持し、各次レベルは前レベルをフィルタした半分サイズ（切り上げ）です。
テクスチャをサンプリングする際、明示的または暗黙的な詳細度（LOD）でミップレベルを選択し、フィルタリングで合成値を生成します。

配列化

テクスチャが配列化されているかどうか。

• 非配列テクスチャはテクセルのグリッドです。

• 配列化テクスチャは同種グリッドの配列です。

配列サイズ

配列化テクスチャの場合、グリッド数。

サンプル数

テクスチャがマルチサンプルの場合のサンプル数。

テクスチャ内の各テクセルには一意な論理テクセルアドレスが関連付けられており、これは整数タプルで：

• [0, ミップレベル数)のミップレベル

• 次元性で決まる成分数。それぞれの成分値は[0, S_i)でS_iはその成分のサイズ

• テクスチャが配列化されている場合、[0, 配列サイズ)の配列インデックス。サイズはミップレベルの関数。

• テクスチャがマルチサンプルの場合、[0, サンプル数)のサンプルインデックス。

テクスチャの物理的な構造は描画最適化のために設計されており、データレイアウトや型、内部処理はシェーダ言語から直接指定できません。

このため、シェーダはテクスチャ変数内のテクセルメモリに直接アクセスできません。 代わりに不透明なハンドル経由でアクセスします：

• シェーダ内：

  • モジュールスコープ変数を宣言し、格納型を後述するテクスチャ型のいずれかにする。 変数は基底テクスチャメモリへの不透明ハンドルを格納し、自動的にhandleアドレス空間に配置される。

  • 関数内でテクスチャ組み込み関数を呼び出し、テクスチャ変数や関数仮引数を引数として渡す。

• WebGPUパイプライン構築時、テクスチャ変数の格納型・バインディングは対応するバインドグループレイアウトエントリと互換でなければならない。