MSL: Invalid code generated using argument buffers and atomic image operations

[stuartcarnie]

[!NOTE]

This issue was originally discovered whilst adding Metal support to Godot during the compilation of FSR 2.

Issue

The following GLSL code:

#version 450

layout (set = 1, binding = 0, r32ui) coherent uniform uimage2D   rw_spd_global_atomic;

void SPD_IncreaseAtomicCounter(inout uint spdCounter)
{
    spdCounter = imageAtomicAdd(rw_spd_global_atomic, ivec2(0,0), 1);
}

void ComputeAutoExposure() {
    uint v = 0;
    SPD_IncreaseAtomicCounter(v);
}

layout (local_size_x = 256, local_size_y = 1, local_size_z = 1) in;

void main()
{
    ComputeAutoExposure();
}

when transpiled to MSL (version < 3.1), and using argument buffers, drops the volatile attribute for the argument to the ComputeAutoExposure function, which fails to compile:

#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wunused-variable"

#include <metal_stdlib>
#include <simd/simd.h>
#include <metal_atomic>

using namespace metal;

// The required alignment of a linear texture of R32Uint format.
constant uint spvLinearTextureAlignmentOverride [[function_constant(65535)]];
constant uint spvLinearTextureAlignment = is_function_constant_defined(spvLinearTextureAlignmentOverride) ? spvLinearTextureAlignmentOverride : 4;
// Returns buffer coords corresponding to 2D texture coords for emulating 2D texture atomics
#define spvImage2DAtomicCoord(tc, tex) (((((tex).get_width() +  spvLinearTextureAlignment / 4 - 1) & ~( spvLinearTextureAlignment / 4 - 1)) * (tc).y) + (tc).x)

struct spvDescriptorSetBuffer1
{
    texture2d<uint> rw_spd_global_atomic [[id(0)]];
    volatile device atomic_uint* rw_spd_global_atomic_atomic [[id(1)]];
};

static inline __attribute__((always_inline))
void SPD_IncreaseAtomicCounter(thread uint& spdCounter, texture2d<uint> rw_spd_global_atomic, device atomic_uint* rw_spd_global_atomic_atomic)
{
    uint _25 = atomic_fetch_add_explicit((device atomic_uint*)&rw_spd_global_atomic_atomic[spvImage2DAtomicCoord(int2(0), rw_spd_global_atomic)], 1u, memory_order_relaxed);
    spdCounter = _25;
}

static inline __attribute__((always_inline))
void ComputeAutoExposure(texture2d<uint> rw_spd_global_atomic, device atomic_uint* rw_spd_global_atomic_atomic)
{
    uint v = 0u;
    uint param = v;
    SPD_IncreaseAtomicCounter(param, rw_spd_global_atomic, rw_spd_global_atomic_atomic);
    v = param;
}

kernel void main0(constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]])
{
    ComputeAutoExposure(spvDescriptorSet1.rw_spd_global_atomic, spvDescriptorSet1.rw_spd_global_atomic_atomic);
}

Error:

<stdin>:40:5: error: no matching function for call to 'ComputeAutoExposure'
    ComputeAutoExposure(spvDescriptorSet1.rw_spd_global_atomic, spvDescriptorSet1.rw_spd_global_atomic_atomic);
    ^~~~~~~~~~~~~~~~~~~
<stdin>:30:6: note: candidate function not viable: 2nd argument ('volatile device metal::atomic_uint *const constant' (aka 'volatile device _atomic<unsigned int> *const constant')) would lose volatile qualifier
void ComputeAutoExposure(texture2d<uint> rw_spd_global_atomic, device atomic_uint* rw_spd_global_atomic_atomic)
     ^
1 error generated.

---

The following commands were used.

Compiled to SPIR-V:

external/glslang-build/output/bin/glslang --target-env vulkan1.3 -V shaders-msl-no-opt/comp/volatile.argument.comp

Convert to MSL:

cmake-build-debug/spirv-cross --msl --msl-argument-buffers --msl-version 20200 comp.spv

[HansKristian-Work]

Do be aware that the FidelityFX SPD shader (assuming that's what this code is) cannot work correctly on Metal since Metal does not support the concept of "coherent" storage images or buffers.

I'll still have a look to see what is going wrong here. You can also try enabling Metal 3.1, which uses proper image atomics and not this broken workaround.

[stuartcarnie]

Do be aware that the FidelityFX SPD shader (assuming that's what this code is) cannot work correctly on Metal since Metal does not support the concept of "coherent" storage images or buffers.

You are correct. Interestingly, visually things look ok, as I'm implementing a Metal renderer for Godot, which can enable FSR 2.

I'll still have a look to see what is going wrong here. You can also try enabling Metal 3.1, which uses proper image atomics and not this broken workaround.

I did try that, but ran into a separate issue. I can try again and report as an issue.