fix metal issue

gfx-rs · May 27, 2024 · e0bd41a · e0bd41a
1 parent 16262d3
commit e0bd41a
Show file tree

Hide file tree

Showing 2 changed files with 9 additions and 4 deletions.
diff --git a/tests/tests/dispatch_workgroups_indirect.rs b/tests/tests/dispatch_workgroups_indirect.rs
@@ -63,12 +63,14 @@ async fn run_test(
 ) -> [u32; 3] {
     const SHADER_SRC: &str = "
         @group(0) @binding(0)
-        var<storage, read_write> out: vec3<u32>;
+        var<storage, read_write> out: array<u32, 3>;
 
         @compute @workgroup_size(1)
         fn main(@builtin(num_workgroups) num_workgroups: vec3<u32>, @builtin(workgroup_id) workgroup_id: vec3<u32>) {
             if (all(workgroup_id == vec3<u32>())) {
-                out = num_workgroups;
+                out[0] = num_workgroups.x;
+                out[1] = num_workgroups.y;
+                out[2] = num_workgroups.z;
             }
         }
     ";

diff --git a/wgpu-core/src/device/global.rs b/wgpu-core/src/device/global.rs
@@ -1799,13 +1799,16 @@ impl Global {
                 @group(0) @binding(0)
                 var<storage, read> src: array<u32>;
                 @group(0) @binding(1)
-                var<storage, read_write> dst: vec3<u32>;
+                var<storage, read_write> dst: array<u32, 3>;
 
                 @compute @workgroup_size(1)
                 fn main() {{
                     let len = arrayLength(&src);
                     let src = vec3(src[len - 3], src[len - 2], src[len - 1]);
-                    dst = select(src, vec3<u32>(), src > vec3({max_compute_workgroups_per_dimension}u));
+                    let res = select(src, vec3<u32>(), src > vec3({max_compute_workgroups_per_dimension}u));
+                    dst[0] = res.x;
+                    dst[1] = res.y;
+                    dst[2] = res.z;
                 }}
             ");