Fix hardcoded compute

Author: LegNeato

tests/difftests/pipelines/compute/hardcoded/compute.wgsl

@@ -1,7 +1,7 @@
 @group(0) @binding(0)
-var<storage, read_write> data: array<u32>;
+var<storage, read_write> output: array<u32>;
 
-@compute @workgroup_size(64)
-fn main_cs(@builtin(global_invocation_id) id: vec3<u32>) {
-    data[id.x] = id.x;
+@compute @workgroup_size(1)
+fn main_cs() {
+    output[0] = 42u;
 }

tests/difftests/pipelines/compute/hardcoded/src/lib.rs

@@ -1,11 +1,9 @@
 #![no_std]
-use spirv_std::glam;
 use spirv_std::spirv;
 
 #[spirv(compute(threads(64)))]
 pub fn main_cs(
-    #[spirv(global_invocation_id)] id: glam::UVec3,
     #[spirv(storage_buffer, descriptor_set = 0, binding = 0)] data: &mut [u32],
 ) {
-    data[id.x as usize] = id.x;
+    data[0] = 42;
 }

tests/difftests/src/main.rs

@@ -548,12 +548,20 @@ mod runners {
                 compilation_options: Default::default(),
             });
             let buffer_size = config.compute_output.buffer_size();
+            let initial_data = vec![0u8; buffer_size as usize];
             let buffer = device.create_buffer(&wgpu::BufferDescriptor {
                 label: Some("Result Buffer"),
                 size: buffer_size,
-                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
-                mapped_at_creation: false,
+                usage: wgpu::BufferUsages::STORAGE
+                    | wgpu::BufferUsages::COPY_SRC
+                    | wgpu::BufferUsages::COPY_DST,
+                mapped_at_creation: true,
             });
+            {
+                let mut mapping = buffer.slice(..).get_mapped_range_mut();
+                mapping.copy_from_slice(&initial_data);
+            }
+            buffer.unmap();
             let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
                 layout: &pipeline.get_bind_group_layout(0),
                 entries: &[wgpu::BindGroupEntry {