Allow non struct buffers in wgsl

cts_runner/examples/hello-compute.js

@@ -4,13 +4,10 @@ const numbers = [1, 4, 3, 295];
 
 const device = await adapter.requestDevice();
 
-const shaderCode = `@block
-struct PrimeIndices {
-    data: @stride(4) array<u32>;
-}; // this is used as both input and output for convenience
+const shaderCode = `
 @group(0)
 @binding(0)
-var<storage, read_write> v_indices: PrimeIndices;
+var<storage, read_write> v_indices: array<u32>; // this is used as both input and output for convenience
 // The Collatz Conjecture states that for any integer n:
 // If n is even, n = n/2
 // If n is odd, n = 3n+1
@@ -41,7 +38,7 @@ fn collatz_iterations(n_base: u32) -> u32{
 @stage(compute)
 @workgroup_size(1)
 fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
-    v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
+    v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]);
 }`;
 
 const shaderModule = device.createShaderModule({

player/tests/data/zero-init-buffer-for-binding.wgsl

@@ -1,13 +1,9 @@
-struct InOutBuffer {
-    data: @stride(4) array<u32>;
-};
-
 @group(0)
 @binding(0)
-var<storage, read_write> buffer: InOutBuffer;
+var<storage, read_write> buffer: array<u32>;
 
 @stage(compute)
 @workgroup_size(1)
 fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
-    buffer.data[global_id.x] = buffer.data[global_id.x] + global_id.x;
+    buffer[global_id.x] = buffer[global_id.x] + global_id.x;
 }

wgpu-core/Cargo.toml

@@ -38,7 +38,7 @@ thiserror = "1"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "81dc674"
+rev = "0ce98d6"
 #version = "0.8"
 features = ["span", "validate", "wgsl-in"]
 

wgpu-core/src/validation.rs

@@ -25,7 +25,7 @@ struct Resource {
     name: Option<String>,
     bind: naga::ResourceBinding,
     ty: ResourceType,
-    class: naga::StorageClass,
+    class: naga::AddressSpace,
 }
 
 #[derive(Clone, Copy, Debug)]
@@ -181,9 +181,9 @@ pub enum BindingError {
     #[error("type on the shader side does not match the pipeline binding")]
     WrongType,
     #[error("storage class {binding:?} doesn't match the shader {shader:?}")]
-    WrongStorageClass {
-        binding: naga::StorageClass,
-        shader: naga::StorageClass,
+    WrongAddressSpace {
+        binding: naga::AddressSpace,
+        shader: naga::AddressSpace,
     },
     #[error("buffer structure size {0}, added to one element of an unbound array, if it's the last field, ended up greater than the given `min_binding_size`")]
     WrongBufferSize(wgt::BufferSize),
@@ -373,23 +373,23 @@ impl Resource {
                     } => {
                         let (class, global_use) = match ty {
                             wgt::BufferBindingType::Uniform => {
-                                (naga::StorageClass::Uniform, GlobalUse::READ)
+                                (naga::AddressSpace::Uniform, GlobalUse::READ)
                             }
                             wgt::BufferBindingType::Storage { read_only } => {
                                 let mut global_use = GlobalUse::READ | GlobalUse::QUERY;
                                 global_use.set(GlobalUse::WRITE, !read_only);
                                 let mut naga_access = naga::StorageAccess::LOAD;
                                 naga_access.set(naga::StorageAccess::STORE, !read_only);
                                 (
-                                    naga::StorageClass::Storage {
+                                    naga::AddressSpace::Storage {
                                         access: naga_access,
                                     },
                                     global_use,
                                 )
                             }
                         };
                         if self.class != class {
-                            return Err(BindingError::WrongStorageClass {
+                            return Err(BindingError::WrongAddressSpace {
                                 binding: class,
                                 shader: self.class,
                             });
@@ -540,8 +540,8 @@ impl Resource {
         Ok(match self.ty {
             ResourceType::Buffer { size } => BindingType::Buffer {
                 ty: match self.class {
-                    naga::StorageClass::Uniform => wgt::BufferBindingType::Uniform,
-                    naga::StorageClass::Storage { .. } => wgt::BufferBindingType::Storage {
+                    naga::AddressSpace::Uniform => wgt::BufferBindingType::Uniform,
+                    naga::AddressSpace::Storage { .. } => wgt::BufferBindingType::Storage {
                         read_only: !shader_usage.contains(GlobalUse::WRITE),
                     },
                     _ => return Err(BindingError::WrongType),
@@ -905,6 +905,9 @@ impl Interface {
                     class,
                 },
                 naga::TypeInner::Sampler { comparison } => ResourceType::Sampler { comparison },
+                naga::TypeInner::Array { stride, .. } => ResourceType::Buffer {
+                    size: wgt::BufferSize::new(stride as u64).unwrap(),
+                },
                 ref other => {
                     log::error!("Unexpected resource type: {:?}", other);
                     continue;
@@ -915,7 +918,7 @@ impl Interface {
                     name: var.name.clone(),
                     bind,
                     ty,
-                    class: var.class,
+                    class: var.space,
                 },
                 Default::default(),
             );

wgpu-hal/Cargo.toml

@@ -88,14 +88,14 @@ js-sys = { version = "0.3" }
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "81dc674"
+rev = "0ce98d6"
 #version = "0.8"
 
 # DEV dependencies
 
 [dev-dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "81dc674"
+rev = "0ce98d6"
 #version = "0.8"
 features = ["wgsl-in"]
 

wgpu-hal/src/gles/device.rs

@@ -33,9 +33,9 @@ impl CompilationContext<'_> {
             if ep_info[handle].is_empty() {
                 continue;
             }
-            let register = match var.class {
-                naga::StorageClass::Uniform => super::BindingRegister::UniformBuffers,
-                naga::StorageClass::Storage { .. } => super::BindingRegister::StorageBuffers,
+            let register = match var.space {
+                naga::AddressSpace::Uniform => super::BindingRegister::UniformBuffers,
+                naga::AddressSpace::Storage { .. } => super::BindingRegister::StorageBuffers,
                 _ => continue,
             };
 

wgpu-hal/src/metal/device.rs

@@ -116,7 +116,7 @@ impl super::Device {
         let mut sized_bindings = Vec::new();
         let mut immutable_buffer_mask = 0;
         for (var_handle, var) in module.global_variables.iter() {
-            if var.class == naga::StorageClass::WorkGroup {
+            if var.space == naga::AddressSpace::WorkGroup {
                 let size = module.types[var.ty].inner.size(&module.constants);
                 wg_memory_sizes.push(size);
             }
@@ -128,8 +128,8 @@ impl super::Device {
                 };
 
                 if !ep_info[var_handle].is_empty() {
-                    let storage_access_store = match var.class {
-                        naga::StorageClass::Storage { access } => {
+                    let storage_access_store = match var.space {
+                        naga::AddressSpace::Storage { access } => {
                             access.contains(naga::StorageAccess::STORE)
                         }
                         _ => false,

wgpu/Cargo.toml

@@ -137,20 +137,20 @@ env_logger = "0.8"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "81dc674"
+rev = "0ce98d6"
 #version = "0.8"
 optional = true
 
 # used to test all the example shaders
 [dev-dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "81dc674"
+rev = "0ce98d6"
 #version = "0.8"
 features = ["wgsl-in"]
 
 [target.'cfg(target_arch = "wasm32")'.dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "81dc674"
+rev = "0ce98d6"
 #version = "0.8"
 features = ["wgsl-out"]
 

wgpu/examples/boids/compute.wgsl

@@ -13,26 +13,22 @@ struct SimParams {
   rule3Scale : f32;
 };
 
-struct Particles {
-  particles : @stride(16) array<Particle>;
-};
-
 @group(0) @binding(0) var<uniform> params : SimParams;
-@group(0) @binding(1) var<storage, read> particlesSrc : Particles;
-@group(0) @binding(2) var<storage, read_write> particlesDst : Particles;
+@group(0) @binding(1) var<storage, read> particlesSrc : array<Particle>;
+@group(0) @binding(2) var<storage, read_write> particlesDst : array<Particle>;
 
 // https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
 @stage(compute)
 @workgroup_size(64)
 fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
-  let total = arrayLength(&particlesSrc.particles);
+  let total = arrayLength(&particlesSrc);
   let index = global_invocation_id.x;
   if (index >= total) {
     return;
   }
 
-  var vPos : vec2<f32> = particlesSrc.particles[index].pos;
-  var vVel : vec2<f32> = particlesSrc.particles[index].vel;
+  var vPos : vec2<f32> = particlesSrc[index].pos;
+  var vVel : vec2<f32> = particlesSrc[index].vel;
 
   var cMass : vec2<f32> = vec2<f32>(0.0, 0.0);
   var cVel : vec2<f32> = vec2<f32>(0.0, 0.0);
@@ -49,8 +45,8 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
       continue;
     }
 
-    let pos = particlesSrc.particles[i].pos;
-    let vel = particlesSrc.particles[i].vel;
+    let pos = particlesSrc[i].pos;
+    let vel = particlesSrc[i].vel;
 
     if (distance(pos, vPos) < params.rule1Distance) {
       cMass += pos;
@@ -100,5 +96,5 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
   }
 
   // Write back
-  particlesDst.particles[index] = Particle(vPos, vVel);
+  particlesDst[index] = Particle(vPos, vVel);
 }

wgpu/examples/hello-compute/shader.wgsl

@@ -1,10 +1,6 @@
-struct PrimeIndices {
-    data: @stride(4) array<u32>;
-}; // this is used as both input and output for convenience
-
 @group(0)
 @binding(0)
-var<storage, read_write> v_indices: PrimeIndices;
+var<storage, read_write> v_indices: array<u32>; // this is used as both input and output for convenience
 
 // The Collatz Conjecture states that for any integer n:
 // If n is even, n = n/2
@@ -38,5 +34,5 @@ fn collatz_iterations(n_base: u32) -> u32{
 @stage(compute)
 @workgroup_size(1)
 fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
-    v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
+    v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]);
 }

wgpu/examples/shadow/shader.wgsl

@@ -49,21 +49,12 @@ struct Light {
     color: vec4<f32>;
 };
 
-struct Lights {
-    data: @stride(96) array<Light>;
-};
-
-// Used when storage types are not supported
-struct LightsWithoutStorage {
-    data: array<Light, 10>;
-};
-
 @group(0)
 @binding(1)
-var<storage, read> s_lights: Lights;
+var<storage, read> s_lights: array<Light>;
 @group(0)
 @binding(1)
-var<uniform> u_lights: LightsWithoutStorage;
+var<uniform> u_lights: array<Light, 10>; // Used when storage types are not supported
 @group(0)
 @binding(2)
 var t_shadow: texture_depth_2d_array;
@@ -93,7 +84,7 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     // accumulate color
     var color: vec3<f32> = c_ambient;
     for(var i = 0u; i < min(u_globals.num_lights.x, c_max_lights); i += 1u) {
-        let light = s_lights.data[i];
+        let light = s_lights[i];
         // project into the light space
         let shadow = fetch_shadow(i, light.proj * in.world_position);
         // compute Lambertian diffuse term
@@ -114,7 +105,7 @@ fn fs_main_without_storage(in: VertexOutput) -> @location(0) vec4<f32> {
     for(var i = 0u; i < min(u_globals.num_lights.x, c_max_lights); i += 1u) {
         // This line is the only difference from the entrypoint above. It uses the lights
         // uniform instead of the lights storage buffer
-        let light = u_lights.data[i];
+        let light = u_lights[i];
         let shadow = fetch_shadow(i, light.proj * in.world_position);
         let light_dir = normalize(light.pos.xyz - in.world_position.xyz);
         let diffuse = max(0.0, dot(normal, light_dir));

wgpu/tests/vertex_indices/draw.vert.wgsl

@@ -1,14 +1,10 @@
-struct Indices {
-    arr: array<u32>;
-}; // this is used as both input and output for convenience
-
 @group(0) @binding(0)
-var<storage, read_write> indices: Indices;
+var<storage, read_write> indices: array<u32>; // this is used as both input and output for convenience
 
 @stage(vertex)
 fn vs_main(@builtin(instance_index) instance: u32, @builtin(vertex_index) index: u32) -> @builtin(position) vec4<f32> {
     let idx = instance * 3u + index;
-    indices.arr[idx] = idx;
+    indices[idx] = idx;
     return vec4<f32>(0.0, 0.0, 0.0, 1.0);
 }