SPIRV-Cross/reference/shaders-msl-no-opt/asm/comp/bitscan.asm.comp

#pragma clang diagnostic ignored "-Wmissing-prototypes"

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

using namespace metal;

// Implementation of the GLSL findLSB() function
template<typename T>
inline T spvFindLSB(T x)
{
    return select(ctz(x), T(-1), x == T(0));
}

// Implementation of the signed GLSL findMSB() function
template<typename T>
inline T spvFindSMSB(T x)
{
    T v = select(x, T(-1) - x, x < T(0));
    return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0));
}

// Implementation of the unsigned GLSL findMSB() function
template<typename T>
inline T spvFindUMSB(T x)
{
    return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0));
}

struct SSBO
{
    uint4 u;
    int4 i;
};

kernel void main0(device SSBO& _4 [[buffer(0)]])
{
    uint4 _19 = _4.u;
    int4 _20 = _4.i;
    _4.u = spvFindLSB(_19);
    _4.i = int4(spvFindLSB(_19));
    _4.u = uint4(spvFindLSB(_20));
    _4.i = spvFindLSB(_20);
    _4.u = spvFindUMSB(_19);
    _4.i = int4(spvFindUMSB(_19));
    _4.u = spvFindUMSB(uint4(_20));
    _4.i = int4(spvFindUMSB(uint4(_20)));
    _4.u = uint4(spvFindSMSB(int4(_19)));
    _4.i = spvFindSMSB(int4(_19));
    _4.u = uint4(spvFindSMSB(_20));
    _4.i = spvFindSMSB(_20);
}