Document WebGPU subgroup-size-control capability

docs/user-manual/graphics/shaders/wgsl-capabilities.md

@@ -66,6 +66,11 @@ At device creation, the engine reads `navigator.gpu.wgslLanguageFeatures` and ad
   - **Preprocessor define:** `CAPS_SUBGROUPS`
   - **Shader stages:** fragment and compute
   - **Details:** Subgroup builtins (`subgroupBroadcast`, `subgroupAdd`, …). `device.supportsSubgroupUniformity` does not add a separate `requires` or `enable` line; the driver uses it together with the subgroups feature
+- **`device.supportsSubgroupSizeControl`**
+  - **Engine injects:** *(none)* — `enable subgroups;` is already added when `device.supportsSubgroups` (a dependency of this feature); add `enable subgroup_size_control;` in your own WGSL when you use `@subgroup_size`
+  - **Preprocessor define:** `CAPS_SUBGROUP_SIZE_CONTROL`
+  - **Shader stages:** compute
+  - **Details:** Pin a compute entry point to a fixed subgroup size with the `@subgroup_size(N)` attribute, where `N` is a power of two within the `device.minSubgroupSize`…`device.maxSubgroupSize` range (both read-only). Without it, the driver picks the size and it is not known ahead of time. The engine only advertises the capability; the `enable` line is author-written so usage stays explicit. See the [WebGPU 151-152](https://developer.chrome.com/blog/new-in-webgpu-151-152#subgroup-size) overview
 - **`device.supportsSubgroupId`**
   - **Engine injects:** `requires subgroup_id;`
   - **Preprocessor define:** `CAPS_SUBGROUP_ID`
@@ -121,3 +126,25 @@ fn main(
     _ = wg;
 }
 ```
+
+Example (compute) — pin the subgroup size when `CAPS_SUBGROUP_SIZE_CONTROL` is set, otherwise let the driver choose:
+
+```wgsl
+#ifdef CAPS_SUBGROUP_SIZE_CONTROL
+enable subgroup_size_control; // enable subgroups; is already injected by the engine
+#endif
+
+@compute @workgroup_size(64)
+#ifdef CAPS_SUBGROUP_SIZE_CONTROL
+@subgroup_size(32)            // power of two within device.minSubgroupSize..maxSubgroupSize
+#endif
+fn main(@builtin(global_invocation_id) global_id: vec3u) {
+    _ = global_id;
+}
+```
+
+:::note
+
+All `enable …;` directives must precede every module-scope declaration. The engine prepends its own directives (such as `enable subgroups;`) ahead of your shader source, so place your `enable subgroup_size_control;` at the very top of the chunk — before any `const`, `struct`, `var`, or `fn`. The order among the `enable` directives themselves does not matter.
+
+:::

i18n/ja/docusaurus-plugin-content-docs/current/user-manual/graphics/shaders/wgsl-capabilities.md

@@ -66,6 +66,11 @@ output.color = vec4f(vec3f(result), 1.0);
   - **プリプロセッサ定義:** `CAPS_SUBGROUPS`
   - **シェーダー段階:** フラグメントとコンピュート
   - **説明:** サブグループ組み込み（`subgroupBroadcast` 等）。`device.supportsSubgroupUniformity` 専用の `requires` / `enable` は追加されず、subgroups 機能と合わせて扱う
+- **`device.supportsSubgroupSizeControl`**
+  - **エンジンが注入:** *なし* — `enable subgroups;` は `device.supportsSubgroups` のとき既に付与されます（本機能はこれに依存）。`@subgroup_size` を使う場合は `enable subgroup_size_control;` を自前の WGSL に記述してください
+  - **プリプロセッサ定義:** `CAPS_SUBGROUP_SIZE_CONTROL`
+  - **シェーダー段階:** コンピュート
+  - **説明:** `@subgroup_size(N)` 属性でコンピュートのエントリポイントをサブグループサイズに固定します。`N` は `device.minSubgroupSize`…`device.maxSubgroupSize` の範囲内（いずれも読み取り専用）の2のべき乗です。指定しない場合はドライバがサイズを選び、事前には分かりません。エンジンは能力の宣言のみで、`enable` は筆者が書く想定です。解説: [WebGPU 151-152](https://developer.chrome.com/blog/new-in-webgpu-151-152#subgroup-size)
 - **`device.supportsSubgroupId`**
   - **エンジンが注入:** `requires subgroup_id;`
   - **プリプロセッサ定義:** `CAPS_SUBGROUP_ID`
@@ -121,3 +126,25 @@ fn main(
     _ = wg;
 }
 ```
+
+使用例（コンピュート）— `CAPS_SUBGROUP_SIZE_CONTROL` があるときはサブグループサイズを固定し、なければドライバに選ばせます。
+
+```wgsl
+#ifdef CAPS_SUBGROUP_SIZE_CONTROL
+enable subgroup_size_control; // enable subgroups; はエンジンが既に注入
+#endif
+
+@compute @workgroup_size(64)
+#ifdef CAPS_SUBGROUP_SIZE_CONTROL
+@subgroup_size(32)            // device.minSubgroupSize..maxSubgroupSize 内の2のべき乗
+#endif
+fn main(@builtin(global_invocation_id) global_id: vec3u) {
+    _ = global_id;
+}
+```
+
+:::note
+
+すべての `enable …;` ディレクティブは、モジュールスコープのあらゆる宣言より前に置く必要があります。エンジンは自身のディレクティブ（`enable subgroups;` など）をシェーダーソースの前に付加するため、`enable subgroup_size_control;` はチャンクの先頭、つまり `const`・`struct`・`var`・`fn` のいずれよりも前に記述してください。`enable` ディレクティブどうしの順序は問いません。
+
+:::