Skip to content

Commit

Permalink
Allow non struct buffers in wgsl (#2451)
Browse files Browse the repository at this point in the history
* Bump naga

* Update examples
  • Loading branch information
IcanDivideBy0 authored Feb 5, 2022
1 parent 3c3fbe8 commit 3e0305d
Show file tree
Hide file tree
Showing 13 changed files with 47 additions and 72 deletions.
2 changes: 1 addition & 1 deletion Cargo.lock

Some generated files are not rendered by default. Learn more about how customized files appear on GitHub.

9 changes: 3 additions & 6 deletions cts_runner/examples/hello-compute.js
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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({
Expand Down
8 changes: 2 additions & 6 deletions player/tests/data/zero-init-buffer-for-binding.wgsl
Original file line number Diff line number Diff line change
@@ -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;
}
2 changes: 1 addition & 1 deletion wgpu-core/Cargo.toml
Original file line number Diff line number Diff line change
Expand Up @@ -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"]

Expand Down
23 changes: 13 additions & 10 deletions wgpu-core/src/validation.rs
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ struct Resource {
name: Option<String>,
bind: naga::ResourceBinding,
ty: ResourceType,
class: naga::StorageClass,
class: naga::AddressSpace,
}

#[derive(Clone, Copy, Debug)]
Expand Down Expand Up @@ -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),
Expand Down Expand Up @@ -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,
});
Expand Down Expand Up @@ -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),
Expand Down Expand Up @@ -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;
Expand All @@ -915,7 +918,7 @@ impl Interface {
name: var.name.clone(),
bind,
ty,
class: var.class,
class: var.space,
},
Default::default(),
);
Expand Down
4 changes: 2 additions & 2 deletions wgpu-hal/Cargo.toml
Original file line number Diff line number Diff line change
Expand Up @@ -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"]

Expand Down
6 changes: 3 additions & 3 deletions wgpu-hal/src/gles/device.rs
Original file line number Diff line number Diff line change
Expand Up @@ -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,
};

Expand Down
6 changes: 3 additions & 3 deletions wgpu-hal/src/metal/device.rs
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
Expand All @@ -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,
Expand Down
6 changes: 3 additions & 3 deletions wgpu/Cargo.toml
Original file line number Diff line number Diff line change
Expand Up @@ -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"]

Expand Down
20 changes: 8 additions & 12 deletions wgpu/examples/boids/compute.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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;
Expand Down Expand Up @@ -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);
}
8 changes: 2 additions & 6 deletions wgpu/examples/hello-compute/shader.wgsl
Original file line number Diff line number Diff line change
@@ -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
Expand Down Expand Up @@ -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]);
}
17 changes: 4 additions & 13 deletions wgpu/examples/shadow/shader.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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
Expand All @@ -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));
Expand Down
8 changes: 2 additions & 6 deletions wgpu/tests/vertex_indices/draw.vert.wgsl
Original file line number Diff line number Diff line change
@@ -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);
}

Expand Down

0 comments on commit 3e0305d

Please sign in to comment.