Add Pipeline Overrides for workgroup_size#6635
Conversation
cwfitzgerald
left a comment
There was a problem hiding this comment.
Some nits on the integration tests, looks good on the wgpu side.
cwfitzgerald
left a comment
There was a problem hiding this comment.
Tests look good :) Need a naga review still
|
Don't forget to re-request review when things are addressed! |
|
I think this PR causes crash in servo when testing webgpu:shader,execution,limits:workgroup_array_byte_size_override:* (not reproducible in ff yet because it's using too old wgpu version). I suspect that the problem is that new expr is not put into expr arena. |
|
Yeah, replace is documented to panic:
We don't handle the case where the new value already exists in the arena in the 2 new places we use replace. In the compaction phase we are already building a new type arena so we could move the check and do the work there avoiding @kentslaney could you take a look? |
|
I've just run debugger and the problem is indeed that expr already exists in the arena, with that in mind I was able to change existing test to expose this panic: diff --git a/tests/tests/shader/array_size_overrides.rs b/tests/tests/shader/array_size_overrides.rs
index 7f1d32425..476e414c7 100644
--- a/tests/tests/shader/array_size_overrides.rs
+++ b/tests/tests/shader/array_size_overrides.rs
@@ -9,7 +9,7 @@ const SHADER: &str = r#"
var<workgroup> arr: array<u32, n - 2>;
@group(0) @binding(0)
- var<storage, read_write> output: array<u32>;
+ var<storage, read_write> output: array<u32, 14 - 2>;
@compute @workgroup_size(1) fn main() {
// 1d spiral
@@ -34,9 +34,9 @@ const SHADER: &str = r#"
static ARRAY_SIZE_OVERRIDES: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(TestParameters::default().limits(wgpu::Limits::default()))
.run_async(move |ctx| async move {
- array_size_overrides(&ctx, None, &[534], false).await;
+ //array_size_overrides(&ctx, None, &[534], false).await;
array_size_overrides(&ctx, Some(14), &[286480122], false).await;
- array_size_overrides(&ctx, Some(1), &[0], true).await;
+ //array_size_overrides(&ctx, Some(1), &[0], true).await;
});
async fn array_size_overrides(
|
|
Could you file this as a new issue so it doesn't get lost. |
Connections
resolves #4450
Description
The WebGPU spec allows for pipeline overrides of
workgroup_size. This is achieved by adding aworkgroup_size_overridesfield tocrate::EntryPointand registering any override expressions inworkgroup_sizeas an override with the name__workgroup_size_{[012]}since identifiers aren't allowed to start with two underscores in a row. From there, they are resolved to constants with the other overrides and anySomefields inworkgroup_size_overridesreplace the correspondingworkgroup_size.Testing
This change adds an integration test via
wgpu_test, which checks that the workgroup size equals the initialization value by default and that it changes with pipeline overrides.Checklist
cargo fmt.taplo format.cargo clippy. If applicable, add:--target wasm32-unknown-unknown--target wasm32-unknown-emscriptencargo xtask testto run tests.CHANGELOG.md. See simple instructions inside file.