Slang-RHI/WGPU: Too small buffer is bound

[aleino-nv]

Affected tests:

• tests/language-feature/shader-params/interface-shader-param-ordinary.slang

Example error:

WGPU error: The shader uses more bytes of the buffer (64) than the layout's minBindingSize (48).
 - While validating that the entry-point's declaration for @group(0) @binding(0) matches [BindGroupLayout (unlabeled)]
 - While validating the entry-point's compatibility for group 0 with [BindGroupLayout (unlabeled)]
 - While validating compute stage ([ShaderModule (unlabeled)], entryPoint: computeMain).
 - While calling [Device].CreateComputePipeline([ComputePipelineDescriptor]).

[aleino-nv]

Note: The test tests/language-feature/shader-params/interface-shader-param-ordinary.slang also triggers this issue: #5610

[aleino-nv]

I've dumped the shader that slang-test winds up passing to Slang-RHI/WGPU:

@binding(1) @group(0) var<storage, read_write> gOutputBuffer_0 : array<i32>;

struct _std140_0
{
    @align(16) field0_0 : u32,
    @align(4) field1_0 : u32,
    @align(8) field2_0 : u32,
    @align(4) field3_0 : u32,
};

struct Tuple_std140_0
{
    @align(16) value0_0 : vec2<u32>,
    @align(8) value1_0 : vec2<u32>,
    @align(16) value3_0 : _std140_0,
};

struct GlobalParams_std140_0
{
    @align(16) delta_0 : i32,
    @align(16) gModifier_0 : Tuple_std140_0,
};

struct MyModifier_std140_0
{
    @align(16) extra_0 : i32,
};

struct _std140_1
{
    @align(16) _S1 : GlobalParams_std140_0,
    @align(16) _S2 : MyModifier_std140_0,
};

@binding(0) @group(0) var<uniform> globalParams_0 : _std140_1;
@binding(2) @group(0) var<storage, read_write> _S3 : array<i32>;

struct MyModifier_0
{
    extra_0 : i32,
};

fn unpackStorage_0( _S4 : MyModifier_std140_0) -> MyModifier_0
{
    var _S5 : MyModifier_0 = MyModifier_0( _S4.extra_0 );
    return _S5;
}


fn MyModifier_modify_0( _S6 : MyModifier_0,  _S7 : i32) -> i32
{
    return _S7 * i32(65536) + _S3[_S7] * i32(256) + _S7 * _S6.extra_0;
}

fn test_0( val_0 : i32) -> i32
{
    var _S8 : i32 = MyModifier_modify_0(unpackStorage_0(globalParams_0._S2), val_0);
    return _S8 + globalParams_0._S1.delta_0;
}

@compute
@workgroup_size(4, 1, 1)
fn computeMain(@builtin(global_invocation_id) dispatchThreadID_0 : vec3<u32>)
{
    var dispatchThreadID_1 : vec3<i32> = vec3<i32>(dispatchThreadID_0);
    var _S9 : i32 = dispatchThreadID_1.x;
    var _S10 : i32 = test_0(_S9);
    gOutputBuffer_0[_S9] = _S10;
    return;
}

[aleino-nv]

A uniform buffer size of 64 seems to make sense to me:

struct _std140_0
{
    @align(16) field0_0 : u32, // offs: 0
    @align(4) field1_0 : u32,  // offs: 4
    @align(8) field2_0 : u32,  // offs: 8
    @align(4) field3_0 : u32,  // offs: 12
}; // Size: 16

struct Tuple_std140_0
{
    @align(16) value0_0 : vec2<u32>, // offs: 0
    @align(8) value1_0 : vec2<u32>,  // offs: 8
    @align(16) value3_0 : _std140_0, // offs: 16
}; // Size: 32

struct GlobalParams_std140_0
{
    @align(16) delta_0 : i32,                // offs: 0
    @align(16) gModifier_0 : Tuple_std140_0, // offs: 16         // NOTE: Padded offset due to alignment
}; // Size: 48

struct MyModifier_std140_0
{
    @align(16) extra_0 : i32, // offs: 0
}; // Size: 16                // NOTE: Size is rounded up to alignment (16)

struct _std140_1
{
    @align(16) _S1 : GlobalParams_std140_0, // offs: 0
    @align(16) _S2 : MyModifier_std140_0,   // offs: 48
}; // Size: 64

[aleino-nv]

We do indeed get 48 from the reflection data, here:

https://github.com/shader-slang/slang-rhi/blob/539322d79c1af13bbc5624a333b65a9e856a13a4/src/wgpu/wgpu-shader-object-layout.cpp#L319

There is probably an issue with the layout. Probably I should check carefully that Slang does indeed take the things I marked with NOTE into account.

[aleino-nv]

...
There is probably an issue with the layout. Probably I should check carefully that Slang does indeed take the things I marked with NOTE into account.

Scratch that...
After some more debugging, it seems as though Slang-RHI/WGPU is creating bind group layout entries for two different uniform buffers: one corresponding to GlobalParams_std140_0 which should indeed have a size of 48, and another one corresponding to MyModifier_std140_0 of size 16.

The error comes from the fact that the shader has a uniform with a _std140_1 struct in it. This struct contains the above mentioned structs and would be 64 in size, but Slang-RHI/WGPU doesn't create any bind group layout entry of such a size and instead apparently tries to use the bind group layout entry corresponding to GlobalParams_std140_0, and that's too small.

[aleino-nv]

I need to first verify that the shader is getting specialized as appropriate.
That extra stuff in MyModifier that's not being accounted for depends gModifier being specialized as appropriate, of course!

https://github.com/shader-slang/slang/blob/master/tests/language-feature/shader-params/interface-shader-param-ordinary.slang#L27

[aleino-nv]

As discussed with @csyonghe, not working on this yet because there are some things to resolve about uniform parameters taking interface values.

[aleino-nv]

@bmillsNV FYI I'm moving this to the next sprint.

[aleino-nv]

@csyonghe You mentioned a while back that 'global interface typed uniforms' needs some cleanup that should happen before working on this. Do we have an issue filed for that?

In any case, I'm clearing the sprint field on this task, since I don't know if something else is required before working on this task.