[ExecuTorch][WebGPU] Make shader workgroup sizes runtime-configurable#21006
Merged
Conversation
Pull Request resolved: #20954 Instead of hard-coding each compute shader’s workgroup size, make it a WGSL `override wg_size` set at pipeline creation via a host `WGPUConstantEntry` (mirrors the landed `add` op). One WGSL per kernel; each distinct size becomes its own Dawn-cached pipeline — so the size can be tuned per device/shape WITHOUT generating another shader variant. Key changes: - `rms_norm.wgsl` / `sdpa_softmax.wgsl` / `q4gsw_linear_coop4_bicol.wgsl` (the free-knob reduction + decode-GEMV kernels) — `const WG*` -> `override wg_size`; shared arrays, tree-reduction strides, and `@workgroup_size` all use it; handlers wire it from `clamp_workgroup_size(...)`. - `QuantizedLinear.cpp` — split GEMV out of `fixed_wg` so it takes the override; steel (256) + shmem (64) stay fixed. - `q4gsw_linear_gemm_shmem.wgsl` — its `override wg_size` was a false knob (geometry hard-locked to 64, never host-wired) -> plain `const`. - Geometry-lock notes on the kernels that stay fixed (`q4gsw_linear_gemm_steel.wgsl`, `sdpa_fd_decode/sdpa_fd_split.wgsl`, `sdpa_fd_reduce.wgsl`) — their size is bound to tile/head-dim geometry, not a runtime knob. - Regenerated the affected `_wgsl.h` headers. Behavior is identical at the default 64: `clamp_workgroup_size` only reduces toward `maxComputeInvocationsPerWorkgroup` (spec floor 256), so 64 is emitted; dispatch counts are unchanged. Co-authored-with: Claude Code. ghstack-source-id: 403260281 @exported-using-ghexport Differential Revision: [D112060863](https://our.internmc.facebook.com/intern/diff/D112060863/)
🔗 Helpful Links🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/21006
Note: Links to docs will display an error until the docs builds have been completed. ✅ No FailuresAs of commit eb9d8a7 with merge base 5c9cd91 ( This comment was automatically generated by Dr. CI and updates every 15 minutes. |
This PR needs a
|
meta-cla
Bot
added
the
CLA Signed
Pull Request resolved: #20955 `q4gsw_linear_coop4.wgsl` is an unwired dead shader: no handler `#include`s its generated `q4gsw_linear_coop4_wgsl.h`, and it is superseded by `q4gsw_linear_coop4_bicol.wgsl` (the live GEMV decode kernel that `QuantizedLinear.cpp` uses). Remove the shader and its generated header. Key changes: - Deleted `q4gsw_linear_coop4.wgsl` and `q4gsw_linear_coop4_wgsl.h`. - Dropped a stale reference to the removed kernel in `test_quantized_linear.py` and `test_webgpu_native.cpp` comments. No build/CMake change: `WEBGPU_SRCS` lists only `.cpp`, and shaders wire via the `gen_wgsl_headers.py` glob (no static list), so removing both files leaves codegen drift-clean. Co-authored-with: Claude Code. ghstack-source-id: 403097256 @exported-using-ghexport Differential Revision: [D112060303](https://our.internmc.facebook.com/intern/diff/D112060303/)
…rable Pull Request resolved: #20956 Add a native gtest that locks the wg-size configurability guarantee: the same rms_norm WGSL, run at override `wg_size` 64 and 128, must produce the same output. `Module::forward` always uses the handler-clamped size (64), so the test builds the compute pipeline directly and sets `wg_size` via a `WGPUConstantEntry`, mirroring the runtime pipeline setup + `WebGPUGraph::copy_outputs` readback. Key changes: - `test/test_webgpu_native.cpp` — `run_rms_norm_at_wg()` standalone-pipeline helper + `RmsNormWorkgroupSizeConfigurable` test (element-wise agree within abs 1e-4 / rel 1e-3, plus a non-zero sanity guard). Self-contained (no `.pte`/export), so it runs inside `webgpu_native_test` under `test_webgpu_native_ci.sh` with no new wiring. Absolute rms_norm correctness stays covered by the model-driven golden tests; this adds the "128 matches 64" guarantee. Co-authored-with: Claude Code. ghstack-source-id: 403097259 @exported-using-ghexport Differential Revision: [D112060302](https://our.internmc.facebook.com/intern/diff/D112060302/)
JCNTH
approved these changes
Jul 17, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This PR was created by the merge bot to help merge the original PR into the main branch.
ghstack PR number: #20954 by @JCNTH
^ Please use this as the source of truth for the PR details, comments, and reviews
ghstack PR base: https://github.com/pytorch/executorch/tree/gh/JCNTH/86/base
ghstack PR head: https://github.com/pytorch/executorch/tree/gh/JCNTH/86/head
Merge bot PR base: https://github.com/pytorch/executorch/tree/main
Merge bot PR head: https://github.com/pytorch/executorch/tree/gh/JCNTH/86/orig
@diff-train-skip-merge