#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wincompatible-pointer-types-discards-qualifiers"

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

using namespace metal;

template<typename T>
struct spvDescriptor
{
    T value;
};

template<typename T>
struct spvDescriptorArray
{
    spvDescriptorArray(const device spvDescriptor<T>* ptr_) : ptr(&ptr_->value) {}
    spvDescriptorArray(const device void *ptr_) : spvDescriptorArray(static_cast<const device spvDescriptor<T>*>(ptr_)) {}
    const device T& operator [] (size_t i) const { return ptr[i]; }
    const device T* ptr;
};

template <typename ImageT>
void spvImageFence(ImageT img) { img.fence(); }

struct B10
{
    float v;
};

struct B11
{
    float v;
};

struct B20
{
    float v;
};

struct B21
{
    float v;
};

struct B30
{
    uint i;
};

struct B31
{
    float v;
};

struct B40
{
    float v;
};

struct B41
{
    float v;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);

struct spvDescriptorSetBuffer0
{
    array<texture2d<float>, 8> t00 [[id(0)]];
    // Overlapping binding: array<texture2d<uint>, 8> t01 [[id(0)]];
    // Overlapping binding: array<texture2d<int>, 8> t02 [[id(0)]];
    // Overlapping binding: array<texture_buffer<uint, access::read_write>, 8> u0 [[id(0)]];
    // Overlapping binding: array<sampler, 8> s00 [[id(0)]];
};

struct spvDescriptorSetBuffer1
{
    spvDescriptor<const device B30 *> b30 [[id(0)]][1] /* unsized array hack */;
    // Overlapping binding: spvDescriptor<constant B31 *> b31 [[id(0)]][1] /* unsized array hack */;
    // Overlapping binding: spvDescriptor<texture2d<float>> t30 [[id(0)]][1] /* unsized array hack */;
    // Overlapping binding: spvDescriptor<texture2d<uint>> t31 [[id(0)]][1] /* unsized array hack */;
    // Overlapping binding: spvDescriptor<texture2d<int>> t32 [[id(0)]][1] /* unsized array hack */;
    // Overlapping binding: spvDescriptor<texture_buffer<uint, access::read_write>> u3 [[id(0)]][1] /* unsized array hack */;
    // Overlapping binding: spvDescriptor<sampler> s30 [[id(0)]][1] /* unsized array hack */;
};

struct spvDescriptorSetBuffer2
{
    const device B20* b20 [[id(0)]][8];
    // Overlapping binding: constant B21* b21 [[id(0)]][8];
    // Overlapping binding: array<texture2d<float>, 8> t20 [[id(0)]];
    // Overlapping binding: array<texture2d<uint>, 8> t21 [[id(0)]];
    // Overlapping binding: array<texture2d<int>, 8> t22 [[id(0)]];
    // Overlapping binding: array<texture_buffer<uint, access::read_write>, 8> u2 [[id(0)]];
    // Overlapping binding: array<sampler, 8> s20 [[id(0)]];
};

struct spvDescriptorSetBuffer3
{
    const device B10* b10 [[id(0)]][8];
    // Overlapping binding: constant B11* b11 [[id(0)]][8];
    // Overlapping binding: array<texture_buffer<uint, access::read_write>, 8> u1 [[id(0)]];
};

struct spvDescriptorSetBuffer4
{
    const device B40* b40 [[id(0)]];
    // Overlapping binding: constant B41* b41 [[id(0)]];
    // Overlapping binding: depth2d<float> t40 [[id(0)]];
    // Overlapping binding: texture2d<uint> t41 [[id(0)]];
    // Overlapping binding: texture2d<int> t42 [[id(0)]];
    // Overlapping binding: texture_buffer<uint, access::read_write> u4 [[id(0)]];
    // Overlapping binding: sampler s40 [[id(0)]];
};

static inline __attribute__((always_inline))
void in_function(thread float4& r0, const device array<texture2d<float>, 8>& t00, const device array<sampler, 8>& s00, const device array<texture2d<uint>, 8>& t01, const device array<texture2d<int>, 8>& t02, const device array<texture_buffer<uint, access::read_write>, 8>& u0, thread float4& r1, const device B10* constant (&b10)[8], constant B11* constant (&b11)[8], constant array<texture_buffer<uint, access::read_write>, 8>& u1, thread float4& r2, constant array<texture2d<float>, 8>& t20, constant array<sampler, 8>& s20, constant array<texture2d<uint>, 8>& t21, constant array<texture2d<int>, 8>& t22, const device B20* constant (&b20)[8], constant array<texture_buffer<uint, access::read_write>, 8>& u2, constant B21* constant (&b21)[8], const spvDescriptorArray<const device B30*> b30, thread uint3& gl_WorkGroupID, thread float4& r3, const spvDescriptorArray<texture2d<float>> t30, const spvDescriptorArray<sampler> s30, const spvDescriptorArray<texture2d<uint>> t31, const spvDescriptorArray<texture2d<int>> t32, const spvDescriptorArray<constant B31*> b31, const spvDescriptorArray<texture_buffer<uint, access::read_write>> u3, thread float4& r4, depth2d<float> t40, sampler s40, texture2d<uint> t41, texture2d<int> t42, const device B40& b40, constant B41& b41, texture_buffer<uint, access::read_write> u4)
{
    r0 = t00[0].sample(s00[3], float2(0.0), level(0.0));
    r0.x = as_type<float>(t01[1].read(uint2(int2(0)), 0).x);
    r0.y = as_type<float>(t02[2].read(uint2(int2(0)), 0).x);
    spvImageFence(u0[2]);
    r0.z = as_type<float>(u0[2].read(uint(0)).x);
    r1.x = b10[3]->v;
    r1.y = b11[4]->v;
    spvImageFence(u1[2]);
    r1.z = as_type<float>(u1[2].read(uint(0)).x);
    r2 = t20[0].sample(s20[3], float2(0.0), level(0.0));
    r2.x = as_type<float>(t21[1].read(uint2(int2(0)), 0).x);
    r2.y = as_type<float>(t22[2].read(uint2(int2(0)), 0).x);
    spvImageFence(u2[2]);
    r2.z = b20[3]->v + as_type<float>(u2[2].read(uint(0)).x);
    r2.w = b21[4]->v;
    uint i = b30[gl_WorkGroupID.x]->i;
    r3 = t30[i].sample(s30[i + 1u], float2(0.0), level(0.0));
    r3.x = as_type<float>(t31[i + 2u].read(uint2(int2(0)), 0).x);
    r3.y = as_type<float>(t32[i + 3u].read(uint2(int2(0)), 0).x);
    r3.z = b31[i + 5u]->v;
    uint _218 = i + 6u;
    spvImageFence(u3[_218]);
    r3.w = as_type<float>(u3[_218].read(uint(0)).x);
    r4 = float4(t40.sample(s40, float2(0.0), level(0.0)));
    r4.x = as_type<float>(t41.read(uint2(int2(0)), 0).x);
    r4.y = as_type<float>(t42.read(uint2(int2(0)), 0).x);
    r4.z = b40.v + b41.v;
    spvImageFence(u4);
    r4.w = as_type<float>(u4.read(uint(0)).x);
    u0[0].write(as_type<uint4>(r0), uint(0));
    u1[0].write(as_type<uint4>(r1), uint(0));
    u2[0].write(as_type<uint4>(r2), uint(0));
    u3[0].write(as_type<uint4>(r3), uint(0));
    u4.write(as_type<uint4>(r4), uint(0));
}

kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]], constant spvDescriptorSetBuffer3& spvDescriptorSet3 [[buffer(3)]], constant spvDescriptorSetBuffer4& spvDescriptorSet4 [[buffer(4)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
{
    spvDescriptorArray<const device B30*> b30 {spvDescriptorSet1.b30};
    spvDescriptorArray<texture2d<float>> t30 {reinterpret_cast<const device spvDescriptor<texture2d<float>>*>(&spvDescriptorSet1.b30)};
    spvDescriptorArray<sampler> s30 {reinterpret_cast<const device spvDescriptor<sampler>*>(&spvDescriptorSet1.b30)};
    spvDescriptorArray<texture2d<uint>> t31 {reinterpret_cast<const device spvDescriptor<texture2d<uint>>*>(&spvDescriptorSet1.b30)};
    spvDescriptorArray<texture2d<int>> t32 {reinterpret_cast<const device spvDescriptor<texture2d<int>>*>(&spvDescriptorSet1.b30)};
    spvDescriptorArray<constant B31*> b31 {reinterpret_cast<spvDescriptor<constant B31 *> const device *>(&spvDescriptorSet1.b30)};
    spvDescriptorArray<texture_buffer<uint, access::read_write>> u3 {reinterpret_cast<const device spvDescriptor<texture_buffer<uint, access::read_write>>*>(&spvDescriptorSet1.b30)};

    const device auto &t01 = reinterpret_cast<const device array<texture2d<uint>, 8> &>(spvDescriptorSet0.t00);
    const device auto &t02 = reinterpret_cast<const device array<texture2d<int>, 8> &>(spvDescriptorSet0.t00);
    const device auto &u0 = reinterpret_cast<const device array<texture_buffer<uint, access::read_write>, 8> &>(spvDescriptorSet0.t00);
    const device auto &s00 = reinterpret_cast<const device array<sampler, 8> &>(spvDescriptorSet0.t00);
    constant auto &b21 = reinterpret_cast<constant B21* constant (&)[8]>(spvDescriptorSet2.b20);
    constant auto &t20 = reinterpret_cast<constant array<texture2d<float>, 8> &>(spvDescriptorSet2.b20);
    constant auto &t21 = reinterpret_cast<constant array<texture2d<uint>, 8> &>(spvDescriptorSet2.b20);
    constant auto &t22 = reinterpret_cast<constant array<texture2d<int>, 8> &>(spvDescriptorSet2.b20);
    constant auto &u2 = reinterpret_cast<constant array<texture_buffer<uint, access::read_write>, 8> &>(spvDescriptorSet2.b20);
    constant auto &s20 = reinterpret_cast<constant array<sampler, 8> &>(spvDescriptorSet2.b20);
    constant auto &b11 = reinterpret_cast<constant B11* constant (&)[8]>(spvDescriptorSet3.b10);
    constant auto &u1 = reinterpret_cast<constant array<texture_buffer<uint, access::read_write>, 8> &>(spvDescriptorSet3.b10);
    constant auto &b41 = *reinterpret_cast<constant B41* constant &>(spvDescriptorSet4.b40);
    constant auto &t40 = reinterpret_cast<constant depth2d<float> &>(spvDescriptorSet4.b40);
    constant auto &t41 = reinterpret_cast<constant texture2d<uint> &>(spvDescriptorSet4.b40);
    constant auto &t42 = reinterpret_cast<constant texture2d<int> &>(spvDescriptorSet4.b40);
    constant auto &u4 = reinterpret_cast<constant texture_buffer<uint, access::read_write> &>(spvDescriptorSet4.b40);
    constant auto &s40 = reinterpret_cast<constant sampler &>(spvDescriptorSet4.b40);
    float4 r0;
    float4 r1;
    float4 r2;
    float4 r3;
    float4 r4;
    in_function(r0, spvDescriptorSet0.t00, s00, t01, t02, u0, r1, spvDescriptorSet3.b10, b11, u1, r2, t20, s20, t21, t22, spvDescriptorSet2.b20, u2, b21, b30, gl_WorkGroupID, r3, t30, s30, t31, t32, b31, u3, r4, t40, s40, t41, t42, (*spvDescriptorSet4.b40), b41, u4);
}

