Crash using wgpu when trying to create a compute pipeline shader

[birbe]

I'm using wgpu 0.14, on DirectX12 (so Windows). I'm writing a compute shader. I had some validation errors, which I fixed. After doing so, I got this error. (I'm new to writing compute shaders so there's a high chance that it wouldn't work properly, but it should at least not cause a compile issue as I think it's valid wgsl)

thread 'main' panicked at 'wgpu error: Validation Error

Caused by:
    In Device::create_compute_pipeline
      note: label = `Chunk Assembler Pipeline`
    Internal error: D3DCompile error (0x80004005): C:\Users\Birb\wgpu-mc\rust\Chunk Assembler Shader Module(72,42): error X3000: syntax error: unexpected token '{'

The code:

struct Vertex {
    attributes: array<f32, 9>
}

//i32 because this internally starts out as -1 because when it's incremented then the first index we get is 0
@group(0) @binding(0)
var<storage, read_write> vbo_out_index: atomic<i32>;

@group(1) @binding(0)
var<storage> vbo_in: array<Vertex>;

@group(2) @binding(0)
var<storage, read_write> vbo_out: array<Vertex>;

//How many elements are in each chunk. Each vec2 corresponds to a workgroup id
//First component is the beginning index into the vbo_in array of this chunk. Second component is how many elements there are
//Third and fourth are XZ coordinates (MC block space, not chunk coordinates) of the chunk
@group(3) @binding(0)
var<storage> indices: array<vec4<u32>>;

@compute @workgroup_size(10)
fn assemble(
    @builtin(workgroup_id) workgroup_id: vec3<u32>,
    @builtin(local_invocation_index) local_invocation_index: u32
) {
    var chunk_element_info = indices[workgroup_id.x];

    var slice_start: u32 = chunk_element_info.x;
    var element_count: u32 = chunk_element_info.y;

    var elements_per_invocation: u32 = element_count / 10u;
    var working_slice_start: u32 = slice_start + (elements_per_invocation * local_invocation_index);
    var working_slice_end: u32 = working_slice_start + elements_per_invocation;

    var chunk_offset = vec2<u32>(chunk_element_info.z, chunk_element_info.w);

    for(var current_vertex_index = working_slice_start; current_vertex_index < working_slice_end; current_vertex_index = current_vertex_index + 1u) {
        var out_index = u32(atomicAdd(&vbo_out_index, 1));
        var vert: Vertex = vbo_in[current_vertex_index];

        vert.attributes[0] += f32(chunk_offset.x);
        vert.attributes[2] += f32(chunk_offset.y);

        vbo_out[out_index] = vert;
    }
}

And the produced shader:

    struct NagaConstants {
        int base_vertex;
        int base_instance;
        uint other;
    };
    ConstantBuffer<NagaConstants> _NagaConstants: register(b0);

    struct Vertex {
        float attributes[9];
    };

    RWByteAddressBuffer vbo_out_index : register(u0);
    ByteAddressBuffer vbo_in : register(t0);
    RWByteAddressBuffer vbo_out : register(u1);
    ByteAddressBuffer indices : register(t1);

    Vertex ConstructVertex(float arg0[9]) {
        Vertex ret = (Vertex)0;
        ret.attributes = arg0;
        return ret;
    }

    [numthreads(10, 1, 1)]
    void assemble(uint3 workgroup_id : SV_GroupID, uint local_invocation_index : SV_GroupIndex)
    {
        uint4 chunk_element_info = (uint4)0;
        uint slice_start = (uint)0;
        uint element_count = (uint)0;
        uint elements_per_invocation = (uint)0;
        uint working_slice_start = (uint)0;
        uint working_slice_end = (uint)0;
        uint2 chunk_offset = (uint2)0;
        uint current_vertex_index = (uint)0;
        uint out_index = (uint)0;
        Vertex vert = (Vertex)0;

        uint4 _expr8 = asuint(indices.Load4(workgroup_id.x*16));
        chunk_element_info = _expr8;
        uint _expr11 = chunk_element_info.x;
        slice_start = _expr11;
        uint _expr14 = chunk_element_info.y;
        element_count = _expr14;
        uint _expr16 = element_count;
        elements_per_invocation = (_expr16 / 10u);
        uint _expr20 = slice_start;
        uint _expr21 = elements_per_invocation;
        working_slice_start = (_expr20 + (_expr21 * local_invocation_index));
        uint _expr25 = working_slice_start;
        uint _expr26 = elements_per_invocation;
        working_slice_end = (_expr25 + _expr26);
        uint _expr30 = chunk_element_info.z;
        uint _expr32 = chunk_element_info.w;
        chunk_offset = uint2(_expr30, _expr32);
        uint _expr35 = working_slice_start;
        current_vertex_index = _expr35;
        bool loop_init = true;
        while(true) {
            if (!loop_init) {
                uint _expr40 = current_vertex_index;
                current_vertex_index = (_expr40 + 1u);
            }
            loop_init = false;
            uint _expr37 = current_vertex_index;
            uint _expr38 = working_slice_end;
            if ((_expr37 < _expr38)) {
            } else {
                break;
            }
            int _e44; vbo_out_index.InterlockedAdd(0, 1, _e44);
            out_index = uint(_e44);
            uint _expr47 = current_vertex_index;
            Vertex _expr49 = ConstructVertex({asfloat(vbo_in.Load(_expr47*36+0+0)), asfloat(vbo_in.Load(_expr47*36+0+4)), asfloat(vbo_in.Load(_expr47*36+0+8)), asfloat(vbo_in.Load(_expr47*36+0+12)), asfloat(
vbo_in.Load(_expr47*36+0+16)), asfloat(vbo_in.Load(_expr47*36+0+20)), asfloat(vbo_in.Load(_expr47*36+0+24)), asfloat(vbo_in.Load(_expr47*36+0+28)), asfloat(vbo_in.Load(_expr47*36+0+32))});
            vert = _expr49;
            float _expr54 = vert.attributes[0];
            uint _expr56 = chunk_offset.x;
            vert.attributes[0] = (_expr54 + float(_expr56));
            float _expr62 = vert.attributes[2];
            uint _expr64 = chunk_offset.y;
            vert.attributes[2] = (_expr62 + float(_expr64));
            uint _expr67 = out_index;
            Vertex _expr69 = vert;
            {
                Vertex _value3 = _expr69;
                {
                    float _value4[9] = _value3.attributes;
                    vbo_out.Store(_expr67*36+0+0, asuint(_value4[0]));
                    vbo_out.Store(_expr67*36+0+4, asuint(_value4[1]));
                    vbo_out.Store(_expr67*36+0+8, asuint(_value4[2]));
                    vbo_out.Store(_expr67*36+0+12, asuint(_value4[3]));
                    vbo_out.Store(_expr67*36+0+16, asuint(_value4[4]));
                    vbo_out.Store(_expr67*36+0+20, asuint(_value4[5]));
                    vbo_out.Store(_expr67*36+0+24, asuint(_value4[6]));
                    vbo_out.Store(_expr67*36+0+28, asuint(_value4[7]));
                    vbo_out.Store(_expr67*36+0+32, asuint(_value4[8]));
                }
            }
        }
        return;
    }

[SparkyPotato]

MRE:

struct Struct {
    data: array<i32, 2>
}

@group(1) @binding(0)
var<storage> storage_buffer: array<Struct>;

@compute @workgroup_size(10)
fn assemble() {
    let x = storage_buffer[0];
}

Generates

struct Struct {
    int data[2];
};

ByteAddressBuffer storage_buffer : register(t0, space1);

Struct ConstructStruct(int arg0[2]) {
    Struct ret = (Struct)0;
    ret.data = arg0;
    return ret;
}

[numthreads(10, 1, 1)]
void assemble()
{
    Struct x = ConstructStruct(
        { asint(storage_buffer.Load(0+0+0)), asint(storage_buffer.Load(0+0+4)) }
    );
}

Whereas if we directly construct the array

struct Struct {
    data: array<i32, 2>
}

@group(1) @binding(0)
var<storage> storage_buffer: array<Struct>;

@compute @workgroup_size(10)
fn assemble() {
    let x = Struct(array(1, 2));
}

The backend generates this instead

struct Struct {
    int data[2];
};

ByteAddressBuffer storage_buffer : register(t0, space1);

typedef int ret_Constructarray2_int_[2];
ret_Constructarray2_int_ Constructarray2_int_(int arg0, int arg1) {
    int ret[2] = { arg0, arg1 };
    return ret;
}

Struct ConstructStruct(int arg0[2]) {
    Struct ret = (Struct)0;
    ret.data = arg0;
    return ret;
}

[numthreads(10, 1, 1)]
void assemble()
{
    Struct x = ConstructStruct(Constructarray2_int_(1, 2));
}

The issue is that

void x(int x[2]) {}

x({1, 2});

is invalid HLSL, but is generated on Expression::Load from a variable with storage address space (in write_storage_load with TypeInner::Array).

[birbe]

Thanks for the well formulated response; I forced wgpu to run on Vulkan and I got no such compile issues. I haven't been able to test if the compute shader works properly yet, but I assume that it was properly converted to SPIR-V.

[teoxoy]

Closing as duplicate since #2184 is more concise.