// Generated by SharpEmu Gen5MslTranslator.
#include <metal_stdlib>

using namespace metal;

struct SharpEmuUniforms
{
    uint dispatch_limit_x;
    uint dispatch_limit_y;
    uint dispatch_limit_z;
    uint reserved;
    uint buffer_bytes[1];
};

static inline uint sharpemu_load_word(device uint* b, uint bytes, uint addr)
{
    if ((addr & 3u) == 0u)
    {
        return addr + 4u <= bytes ? b[addr >> 2] : 0u;
    }
    uint value = 0u;
    device const uchar* p = (device const uchar*)b;
    for (uint i = 0u; i < 4u; i++)
    {
        if (addr + i < bytes)
        {
            value |= (uint)p[addr + i] << (i * 8u);
        }
    }
    return value;
}

static inline uint sharpemu_load_bytes(device uint* b, uint bytes, uint addr, uint count, bool signExtend)
{
    uint value = 0u;
    device const uchar* p = (device const uchar*)b;
    for (uint i = 0u; i < count; i++)
    {
        if (addr + i < bytes)
        {
            value |= (uint)p[addr + i] << (i * 8u);
        }
    }
    if (signExtend && count < 4u)
    {
        uint shift = 32u - (count * 8u);
        value = (uint)(((int)(value << shift)) >> shift);
    }
    return value;
}

static inline void sharpemu_store_bytes(device uint* b, uint bytes, uint addr, uint value, uint count)
{
    device uchar* p = (device uchar*)b;
    for (uint i = 0u; i < count; i++)
    {
        if (addr + i < bytes)
        {
            p[addr + i] = (uchar)((value >> (i * 8u)) & 0xFFu);
        }
    }
}

static inline uint sharpemu_ballot(bool value)
{
    return (uint)(uint64_t)simd_ballot(value);
}

static constant float sharpemu_off_i4_table[16] =
{
    0.0f, 0.0625f, 0.1250f, 0.1875f, 0.2500f, 0.3125f, 0.3750f, 0.4375f,
    -0.5000f, -0.4375f, -0.3750f, -0.3125f, -0.2500f, -0.1875f, -0.1250f, -0.0625f,
};

kernel void gen5_cs(
    device uint* b0 [[buffer(0)]],
    constant SharpEmuUniforms& sharpemu_uniforms [[buffer(1)]],
    uint3 sharpemu_local_id [[thread_position_in_threadgroup]],
    uint3 sharpemu_group_id [[threadgroup_position_in_grid]],
    uint sharpemu_lane [[thread_index_in_simdgroup]])
{
    uint s[128] = {};
    uint v[256] = {};
    bool exec = true;
    bool vcc = false;
    bool scc = false;
    uint pc = 0u;
    bool active = true;
    uint steps = 0u;
    uint bias[1] = {};
    v[0] = sharpemu_local_id.x;
    v[1] = sharpemu_local_id.y;
    v[2] = sharpemu_local_id.z;
    active = (sharpemu_group_id.x * 32u + sharpemu_local_id.x) < sharpemu_uniforms.dispatch_limit_x
        && (sharpemu_group_id.y * 1u + sharpemu_local_id.y) < sharpemu_uniforms.dispatch_limit_y
        && (sharpemu_group_id.z * 1u + sharpemu_local_id.z) < sharpemu_uniforms.dispatch_limit_z;
    s[106] = 0u;
    s[107] = 0u;
    s[126] = sharpemu_ballot(true);
    s[127] = 0u;

    while (active)
    {
        switch (pc)
        {
        case 0u:
        {
            uint t0 = 0x3FC00000u;
            if (exec) { v[0] = t0; }
            uint t1 = 0x40100000u;
            if (exec) { v[1] = t1; }
            uint t2 = 0x41200000u;
            if (exec) { v[2] = t2; }
            uint t3 = as_type<uint>((fma(as_type<float>(v[0]), as_type<float>(v[1]), as_type<float>(v[2]))));
            if (exec) { v[2] = t3; }
            uint t4 = 0x7FFFFFFFu;
            if (exec) { v[3] = t4; }
            uint t5 = 0x10003u;
            if (exec) { v[4] = t5; }
            uint t6 = as_type<uint>(mulhi(as_type<int>(v[3]), as_type<int>(v[4])));
            if (exec) { v[5] = t6; }
            uint t7 = ((v[3]) * (v[4]));
            if (exec) { v[6] = t7; }
            uint t8 = ((0x0u + 0u + 0u + (0u * ((s[9] >> 16) & 0x3FFFu))) + bias[0]);
            if (exec)
            {
                sharpemu_store_bytes(b0, sharpemu_uniforms.buffer_bytes[0], t8 + 0u, v[2], 4u);
            }
            uint t9 = ((0x4u + 0u + 0u + (0u * ((s[9] >> 16) & 0x3FFFu))) + bias[0]);
            if (exec)
            {
                sharpemu_store_bytes(b0, sharpemu_uniforms.buffer_bytes[0], t9 + 0u, v[5], 4u);
            }
            uint t10 = ((0x8u + 0u + 0u + (0u * ((s[9] >> 16) & 0x3FFFu))) + bias[0]);
            if (exec)
            {
                sharpemu_store_bytes(b0, sharpemu_uniforms.buffer_bytes[0], t10 + 0u, v[6], 4u);
            }
            uint t11 = 0u;
            uint t12 = t11;
            s[126] = t12;
            exec = ((t12) >> sharpemu_lane & 1u) != 0u;
            uint t13 = ((0xCu + 0u + 0u + (0u * ((s[9] >> 16) & 0x3FFFu))) + bias[0]);
            if (exec)
            {
                sharpemu_store_bytes(b0, sharpemu_uniforms.buffer_bytes[0], t13 + 0u, v[2], 4u);
            }
            uint t14 = 0xFFFFFFFFu;
            uint t15 = t14;
            s[126] = t15;
            exec = ((t15) >> sharpemu_lane & 1u) != 0u;
            uint t16 = ((0x10u + 0u + 0u + (0u * ((s[9] >> 16) & 0x3FFFu))) + bias[0]);
            if (exec)
            {
                sharpemu_store_bytes(b0, sharpemu_uniforms.buffer_bytes[0], t16 + 0u, v[0], 4u);
            }
            active = false;
        }
        break;
        default:
            active = false;
            break;
        }
        if (++steps >= 100000u)
        {
            active = false;
        }
    }
}