hlsl-out miscompilation of atomic workgroup variable

[simonask]

Some atomic operations on workgroup-shared variables produce invalid HLSL. Similar to #2137, but without arrays.

This is using naga v0.11.0.

With this compute shader (WGSL):

var<workgroup> foo: atomic<u32>;

@compute
@workgroup_size(2)
fn test_miscompilation(@builtin(local_invocation_id) id: vec3<u32>) {
    if id.x == 0u {
        atomicStore(&foo, 0u);
    }

    workgroupBarrier();

    atomicAdd(&foo, 1u);
}

Naga generates the following HLSL:

groupshared uint foo;

[numthreads(2, 1, 1)]
void test_miscompilation(uint3 id : SV_GroupThreadID, uint3 __global_invocation_id : SV_DispatchThreadID)
{
    if (all(__global_invocation_id == uint3(0u, 0u, 0u))) {
        foo = (uint)0;
    }
    GroupMemoryBarrierWithGroupSync();
    if ((id.x == 0u)) {
        foo = 0u;
    }
    GroupMemoryBarrierWithGroupSync();
    uint _e8; foo.InterlockedAdd(0, 1u, _e8);
    return;
}

... which, when compiled with dxc -T cs_6_7, produces the following error:

test.hlsl:15:19: error: invalid format for vector swizzle 'InterlockedAdd'
    uint _e8; foo.InterlockedAdd(0, 1u, _e8);

[simonask]

I don't know much about HLSL, but it looks like foo.InterlockedAdd(...) (a "method" call style) is actually nonsense?

[teoxoy]

The underlying issue seems to be the same as #2137.
The atomic functions are only methods for RWByteAddressBuffer and we should instead use the intrinsic functions (ex InterlockedAdd) for variables in the workgroup address space.

[ErichDonGubler]

I'll take this one!