Boids not working with Vulkan

[baptistemanson]

Only the boid at index 0 is updated.
Here is a minimal shader to reproduce the issue:

[[block]]
struct Particle {
  pos : vec2<f32>;
  vel : vec2<f32>;
};


[[block]]
struct Particles {
  particles : [[stride(16)]] array<Particle>;
};


[[group(0), binding(1)]] var<storage> particlesSrc : [[access(read)]] Particles;
[[group(0), binding(2)]] var<storage> particlesDst : [[access(read_write)]] Particles;

[[builtin(global_invocation_id)]] var gl_GlobalInvocationID : vec3<u32>;

[[stage(compute), workgroup_size(64)]]
 fn main() {
  const index : u32 = gl_GlobalInvocationID.x;
  var vPos : vec2<f32> = particlesSrc.particles[index].pos;
  vPos = vPos + vec2<f32>(0.01, 0.01);
  particlesDst.particles[index].pos = vPos;
}

This shader should increment by 0.01 all the boids position.

I did transpile it from wgsl => spv, spv => glsl and gets this:

#version 310 es

precision highp float;

struct Particle {
    vec2 pos;
    vec2 vel;
};

struct Particles {
    Particle_block_0 {
    vec2 pos;
    vec2 vel;
}[] particles;
};

buffer Particles_block_1 {
    Particle_block_2 {
    vec2 pos;
    vec2 vel;
}[] particles;
} _group_0_binding_1;

buffer Particles_block_3 {
    Particle_block_4 {
    vec2 pos;
    vec2 vel;
}[] particles;
} _group_0_binding_2;

void main() {
    vec2 vPos;
    vPos = _group_0_binding_1.particles[gl_GlobalInvocationID[0]].pos;
    vPos = (vPos + vec2(0.009999999776482582, 0.009999999776482582));
    _group_0_binding_2.particles[gl_GlobalInvocationID[0]].pos = vPos;
    return;
}

The size of the workgroup seems lost in translation. We should probably have some code to say that we align the shader local to the particle buffer?

Inspecting with RenderDoc, the vkCmdDispatch is called with 924, 1, 1), which are the expected values.

Hope that helps someoone!

[kvark]

I was looking at this on Windows/Intel, and the example just crashes (Unknown Vulkan error on submission...). Thus, unable to capture with RenderDoc and have a close look. Validation layers also don't show anything.
I did find a few things that could be related in Naga and addressed them in gfx-rs/naga#398. However, the bug is still there.

[baptistemanson]

The OpAccessChain performed to get a pointer to the output array is slightly different between a handmade minimum example GLSL compiled with glslvalidator and a minimum example WGSL with naga:

OpAccessChain %_ptr_Uniform_v2float gl_GlobalInvocationID %37 0 // GLSL
OpAccessChain %_ptr_Uniform_Particle_0 Particle_Uniform gl_GlobalInvocationID  // Naga

The original Boid shader doesnt crash for me, it just store over and over the result in the first element of the array, I think.

[kvark]

Ok, crazy idea: drivers are written with an assumption that spirv 1.3 is going to be fed. So they expect "StorageBuffer" class for the global particlesDst and pointers into it. We are generating spirv 1.0, and we use the "Uniform" class instead, and it confuses the drivers.

[kvark]

Fixed upstream in gfx-rs/naga@9b9bd5f

[baptistemanson]

Whaoo, had no chances to fix this on my end, was lacking so much background and knowledge! Good job, enjoy your sunday