#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wmissing-braces"

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

using namespace metal;

template<typename T, size_t Num>
struct spvUnsafeArray
{
    T elements[Num ? Num : 1];
    
    thread T& operator [] (size_t pos) thread
    {
        return elements[pos];
    }
    constexpr const thread T& operator [] (size_t pos) const thread
    {
        return elements[pos];
    }
    
    device T& operator [] (size_t pos) device
    {
        return elements[pos];
    }
    constexpr const device T& operator [] (size_t pos) const device
    {
        return elements[pos];
    }
    
    constexpr const constant T& operator [] (size_t pos) const constant
    {
        return elements[pos];
    }
    
    threadgroup T& operator [] (size_t pos) threadgroup
    {
        return elements[pos];
    }
    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
    {
        return elements[pos];
    }
};

template<typename T, int Cols, int Rows=Cols>
struct spvStorageMatrix
{
    vec<T, Rows> columns[Cols];
    
    spvStorageMatrix() thread = default;
    thread spvStorageMatrix& operator=(initializer_list<vec<T, Rows>> cols) thread
    {
        size_t i;
        thread vec<T, Rows>* col;
        for (i = 0, col = cols.begin(); i < Cols; ++i, ++col)
            columns[i] = *col;
        return *this;
    }
    
    spvStorageMatrix(const thread matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const thread spvStorageMatrix& m) thread = default;
    thread spvStorageMatrix& operator=(const thread matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    thread spvStorageMatrix& operator=(const thread spvStorageMatrix& m) thread = default;
    
    spvStorageMatrix(const constant matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const constant spvStorageMatrix& m) thread = default;
    thread spvStorageMatrix& operator=(const constant matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    thread spvStorageMatrix& operator=(const constant spvStorageMatrix& m) thread = default;
    
    spvStorageMatrix(const device matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const device spvStorageMatrix& m) thread = default;
    thread spvStorageMatrix& operator=(const device matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    thread spvStorageMatrix& operator=(const device spvStorageMatrix& m) thread = default;
    
    spvStorageMatrix(const threadgroup matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup spvStorageMatrix& m) thread = default;
    thread spvStorageMatrix& operator=(const threadgroup matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    thread spvStorageMatrix& operator=(const threadgroup spvStorageMatrix& m) thread = default;
    
    #ifdef __HAVE_IMAGEBLOCKS__
    spvStorageMatrix(const threadgroup_imageblock matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup_imageblock spvStorageMatrix& m) thread = default;
    thread spvStorageMatrix& operator=(const threadgroup_imageblock matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    thread spvStorageMatrix& operator=(const threadgroup_imageblock spvStorageMatrix& m) thread = default;
    #endif
    
    #ifdef __HAVE_RAYTRACING__
    spvStorageMatrix(const ray_data matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const ray_data spvStorageMatrix& m) thread = default;
    thread spvStorageMatrix& operator=(const ray_data matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    thread spvStorageMatrix& operator=(const ray_data spvStorageMatrix& m) thread = default;
    #endif
    
    #ifdef __HAVE_MESH__
    spvStorageMatrix(const object_data matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const object_data spvStorageMatrix& m) thread = default;
    thread spvStorageMatrix& operator=(const object_data matrix<T, Cols, Rows>& m) thread
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    thread spvStorageMatrix& operator=(const object_data spvStorageMatrix& m) thread = default;
    #endif
    
    operator matrix<T, Cols, Rows>() const thread
    {
        matrix<T, Cols, Rows> m;
        for (int i = 0; i < Cols; ++i)
            m.columns[i] = columns[i];
        return m;
    }
    
    vec<T, Rows> operator[](size_t idx) const thread
    {
        return columns[idx];
    }
    thread vec<T, Rows>& operator[](size_t idx) thread
    {
        return columns[idx];
    }
    
    spvStorageMatrix() constant = default;
    
    spvStorageMatrix(const thread matrix<T, Cols, Rows>& m) constant
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const thread spvStorageMatrix& m) constant = default;
    
    spvStorageMatrix(const constant matrix<T, Cols, Rows>& m) constant
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const constant spvStorageMatrix& m) constant = default;
    
    spvStorageMatrix(const device matrix<T, Cols, Rows>& m) constant
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const device spvStorageMatrix& m) constant = default;
    
    spvStorageMatrix(const threadgroup matrix<T, Cols, Rows>& m) constant
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup spvStorageMatrix& m) constant = default;
    
    #ifdef __HAVE_IMAGEBLOCKS__
    spvStorageMatrix(const threadgroup_imageblock matrix<T, Cols, Rows>& m) constant
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup_imageblock spvStorageMatrix& m) constant = default;
    #endif
    
    #ifdef __HAVE_RAYTRACING__
    spvStorageMatrix(const ray_data matrix<T, Cols, Rows>& m) constant
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const ray_data spvStorageMatrix& m) constant = default;
    #endif
    
    #ifdef __HAVE_MESH__
    spvStorageMatrix(const object_data matrix<T, Cols, Rows>& m) constant
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const object_data spvStorageMatrix& m) constant = default;
    #endif
    
    operator matrix<T, Cols, Rows>() const constant
    {
        matrix<T, Cols, Rows> m;
        for (int i = 0; i < Cols; ++i)
            m.columns[i] = columns[i];
        return m;
    }
    
    vec<T, Rows> operator[](size_t idx) const constant
    {
        return columns[idx];
    }
    
    spvStorageMatrix() device = default;
    device spvStorageMatrix& operator=(initializer_list<vec<T, Rows>> cols) device
    {
        size_t i;
        thread vec<T, Rows>* col;
        for (i = 0, col = cols.begin(); i < Cols; ++i, ++col)
            columns[i] = *col;
        return *this;
    }
    
    spvStorageMatrix(const thread matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const thread spvStorageMatrix& m) device = default;
    device spvStorageMatrix& operator=(const thread matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    device spvStorageMatrix& operator=(const thread spvStorageMatrix& m) device = default;
    
    spvStorageMatrix(const constant matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const constant spvStorageMatrix& m) device = default;
    device spvStorageMatrix& operator=(const constant matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    device spvStorageMatrix& operator=(const constant spvStorageMatrix& m) device = default;
    
    spvStorageMatrix(const device matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const device spvStorageMatrix& m) device = default;
    device spvStorageMatrix& operator=(const device matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    device spvStorageMatrix& operator=(const device spvStorageMatrix& m) device = default;
    
    spvStorageMatrix(const threadgroup matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup spvStorageMatrix& m) device = default;
    device spvStorageMatrix& operator=(const threadgroup matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    device spvStorageMatrix& operator=(const threadgroup spvStorageMatrix& m) device = default;
    
    #ifdef __HAVE_IMAGEBLOCKS__
    spvStorageMatrix(const threadgroup_imageblock matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup_imageblock spvStorageMatrix& m) device = default;
    device spvStorageMatrix& operator=(const threadgroup_imageblock matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    device spvStorageMatrix& operator=(const threadgroup_imageblock spvStorageMatrix& m) device = default;
    #endif
    
    #ifdef __HAVE_RAYTRACING__
    spvStorageMatrix(const ray_data matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const ray_data spvStorageMatrix& m) device = default;
    device spvStorageMatrix& operator=(const ray_data matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    device spvStorageMatrix& operator=(const ray_data spvStorageMatrix& m) device = default;
    #endif
    
    #ifdef __HAVE_MESH__
    spvStorageMatrix(const object_data matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const object_data spvStorageMatrix& m) device = default;
    device spvStorageMatrix& operator=(const object_data matrix<T, Cols, Rows>& m) device
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    device spvStorageMatrix& operator=(const object_data spvStorageMatrix& m) device = default;
    #endif
    
    operator matrix<T, Cols, Rows>() const device
    {
        matrix<T, Cols, Rows> m;
        for (int i = 0; i < Cols; ++i)
            m.columns[i] = columns[i];
        return m;
    }
    
    vec<T, Rows> operator[](size_t idx) const device
    {
        return columns[idx];
    }
    device vec<T, Rows>& operator[](size_t idx) device
    {
        return columns[idx];
    }
    
    spvStorageMatrix() threadgroup = default;
    threadgroup spvStorageMatrix& operator=(initializer_list<vec<T, Rows>> cols) threadgroup
    {
        size_t i;
        thread vec<T, Rows>* col;
        for (i = 0, col = cols.begin(); i < Cols; ++i, ++col)
            columns[i] = *col;
        return *this;
    }
    
    spvStorageMatrix(const thread matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const thread spvStorageMatrix& m) threadgroup = default;
    threadgroup spvStorageMatrix& operator=(const thread matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup spvStorageMatrix& operator=(const thread spvStorageMatrix& m) threadgroup = default;
    
    spvStorageMatrix(const constant matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const constant spvStorageMatrix& m) threadgroup = default;
    threadgroup spvStorageMatrix& operator=(const constant matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup spvStorageMatrix& operator=(const constant spvStorageMatrix& m) threadgroup = default;
    
    spvStorageMatrix(const device matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const device spvStorageMatrix& m) threadgroup = default;
    threadgroup spvStorageMatrix& operator=(const device matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup spvStorageMatrix& operator=(const device spvStorageMatrix& m) threadgroup = default;
    
    spvStorageMatrix(const threadgroup matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup spvStorageMatrix& m) threadgroup = default;
    threadgroup spvStorageMatrix& operator=(const threadgroup matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup spvStorageMatrix& operator=(const threadgroup spvStorageMatrix& m) threadgroup = default;
    
    #ifdef __HAVE_IMAGEBLOCKS__
    spvStorageMatrix(const threadgroup_imageblock matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup_imageblock spvStorageMatrix& m) threadgroup = default;
    threadgroup spvStorageMatrix& operator=(const threadgroup_imageblock matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup spvStorageMatrix& operator=(const threadgroup_imageblock spvStorageMatrix& m) threadgroup = default;
    #endif
    
    #ifdef __HAVE_RAYTRACING__
    spvStorageMatrix(const ray_data matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const ray_data spvStorageMatrix& m) threadgroup = default;
    threadgroup spvStorageMatrix& operator=(const ray_data matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup spvStorageMatrix& operator=(const ray_data spvStorageMatrix& m) threadgroup = default;
    #endif
    
    #ifdef __HAVE_MESH__
    spvStorageMatrix(const object_data matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const object_data spvStorageMatrix& m) threadgroup = default;
    threadgroup spvStorageMatrix& operator=(const object_data matrix<T, Cols, Rows>& m) threadgroup
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup spvStorageMatrix& operator=(const object_data spvStorageMatrix& m) threadgroup = default;
    #endif
    
    operator matrix<T, Cols, Rows>() const threadgroup
    {
        matrix<T, Cols, Rows> m;
        for (int i = 0; i < Cols; ++i)
            m.columns[i] = columns[i];
        return m;
    }
    
    vec<T, Rows> operator[](size_t idx) const threadgroup
    {
        return columns[idx];
    }
    threadgroup vec<T, Rows>& operator[](size_t idx) threadgroup
    {
        return columns[idx];
    }
    
    #ifdef __HAVE_IMAGEBLOCKS__
    spvStorageMatrix() threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(initializer_list<vec<T, Rows>> cols) threadgroup_imageblock
    {
        size_t i;
        thread vec<T, Rows>* col;
        for (i = 0, col = cols.begin(); i < Cols; ++i, ++col)
            columns[i] = *col;
        return *this;
    }
    
    spvStorageMatrix(const thread matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const thread spvStorageMatrix& m) threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(const thread matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup_imageblock spvStorageMatrix& operator=(const thread spvStorageMatrix& m) threadgroup_imageblock = default;
    
    spvStorageMatrix(const constant matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const constant spvStorageMatrix& m) threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(const constant matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup_imageblock spvStorageMatrix& operator=(const constant spvStorageMatrix& m) threadgroup_imageblock = default;
    
    spvStorageMatrix(const device matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const device spvStorageMatrix& m) threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(const device matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup_imageblock spvStorageMatrix& operator=(const device spvStorageMatrix& m) threadgroup_imageblock = default;
    
    spvStorageMatrix(const threadgroup matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup spvStorageMatrix& m) threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(const threadgroup matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup_imageblock spvStorageMatrix& operator=(const threadgroup spvStorageMatrix& m) threadgroup_imageblock = default;
    
    spvStorageMatrix(const threadgroup_imageblock matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup_imageblock spvStorageMatrix& m) threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(const threadgroup_imageblock matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup_imageblock spvStorageMatrix& operator=(const threadgroup_imageblock spvStorageMatrix& m) threadgroup_imageblock = default;
    
    #ifdef __HAVE_RAYTRACING__
    spvStorageMatrix(const ray_data matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const ray_data spvStorageMatrix& m) threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(const ray_data matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup_imageblock spvStorageMatrix& operator=(const ray_data spvStorageMatrix& m) threadgroup_imageblock = default;
    #endif
    
    #ifdef __HAVE_MESH__
    spvStorageMatrix(const object_data matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const object_data spvStorageMatrix& m) threadgroup_imageblock = default;
    threadgroup_imageblock spvStorageMatrix& operator=(const object_data matrix<T, Cols, Rows>& m) threadgroup_imageblock
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    threadgroup_imageblock spvStorageMatrix& operator=(const object_data spvStorageMatrix& m) threadgroup_imageblock = default;
    #endif
    
    operator matrix<T, Cols, Rows>() const threadgroup_imageblock
    {
        matrix<T, Cols, Rows> m;
        for (int i = 0; i < Cols; ++i)
            m.columns[i] = columns[i];
        return m;
    }
    
    vec<T, Rows> operator[](size_t idx) const threadgroup_imageblock
    {
        return columns[idx];
    }
    threadgroup_imageblock vec<T, Rows>& operator[](size_t idx) threadgroup_imageblock
    {
        return columns[idx];
    }
    #endif
    
    #ifdef __HAVE_RAYTRACING__
    spvStorageMatrix() ray_data = default;
    ray_data spvStorageMatrix& operator=(initializer_list<vec<T, Rows>> cols) ray_data
    {
        size_t i;
        thread vec<T, Rows>* col;
        for (i = 0, col = cols.begin(); i < Cols; ++i, ++col)
            columns[i] = *col;
        return *this;
    }
    
    spvStorageMatrix(const thread matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const thread spvStorageMatrix& m) ray_data = default;
    ray_data spvStorageMatrix& operator=(const thread matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    ray_data spvStorageMatrix& operator=(const thread spvStorageMatrix& m) ray_data = default;
    
    spvStorageMatrix(const constant matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const constant spvStorageMatrix& m) ray_data = default;
    ray_data spvStorageMatrix& operator=(const constant matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    ray_data spvStorageMatrix& operator=(const constant spvStorageMatrix& m) ray_data = default;
    
    spvStorageMatrix(const device matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const device spvStorageMatrix& m) ray_data = default;
    ray_data spvStorageMatrix& operator=(const device matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    ray_data spvStorageMatrix& operator=(const device spvStorageMatrix& m) ray_data = default;
    
    spvStorageMatrix(const threadgroup matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup spvStorageMatrix& m) ray_data = default;
    ray_data spvStorageMatrix& operator=(const threadgroup matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    ray_data spvStorageMatrix& operator=(const threadgroup spvStorageMatrix& m) ray_data = default;
    
    #ifdef __HAVE_IMAGEBLOCKS__
    spvStorageMatrix(const threadgroup_imageblock matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup_imageblock spvStorageMatrix& m) ray_data = default;
    ray_data spvStorageMatrix& operator=(const threadgroup_imageblock matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    ray_data spvStorageMatrix& operator=(const threadgroup_imageblock spvStorageMatrix& m) ray_data = default;
    #endif
    
    spvStorageMatrix(const ray_data matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const ray_data spvStorageMatrix& m) ray_data = default;
    ray_data spvStorageMatrix& operator=(const ray_data matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    ray_data spvStorageMatrix& operator=(const ray_data spvStorageMatrix& m) ray_data = default;
    
    #ifdef __HAVE_MESH__
    spvStorageMatrix(const object_data matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const object_data spvStorageMatrix& m) ray_data = default;
    ray_data spvStorageMatrix& operator=(const object_data matrix<T, Cols, Rows>& m) ray_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    ray_data spvStorageMatrix& operator=(const object_data spvStorageMatrix& m) ray_data = default;
    #endif
    
    operator matrix<T, Cols, Rows>() const ray_data
    {
        matrix<T, Cols, Rows> m;
        for (int i = 0; i < Cols; ++i)
            m.columns[i] = columns[i];
        return m;
    }
    
    vec<T, Rows> operator[](size_t idx) const ray_data
    {
        return columns[idx];
    }
    ray_data vec<T, Rows>& operator[](size_t idx) ray_data
    {
        return columns[idx];
    }
    #endif
    
    #ifdef __HAVE_MESH__
    spvStorageMatrix() object_data = default;
    object_data spvStorageMatrix& operator=(initializer_list<vec<T, Rows>> cols) object_data
    {
        size_t i;
        thread vec<T, Rows>* col;
        for (i = 0, col = cols.begin(); i < Cols; ++i, ++col)
            columns[i] = *col;
        return *this;
    }
    
    spvStorageMatrix(const thread matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const thread spvStorageMatrix& m) object_data = default;
    object_data spvStorageMatrix& operator=(const thread matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    object_data spvStorageMatrix& operator=(const thread spvStorageMatrix& m) object_data = default;
    
    spvStorageMatrix(const constant matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const constant spvStorageMatrix& m) object_data = default;
    object_data spvStorageMatrix& operator=(const constant matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    object_data spvStorageMatrix& operator=(const constant spvStorageMatrix& m) object_data = default;
    
    spvStorageMatrix(const device matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const device spvStorageMatrix& m) object_data = default;
    object_data spvStorageMatrix& operator=(const device matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    object_data spvStorageMatrix& operator=(const device spvStorageMatrix& m) object_data = default;
    
    spvStorageMatrix(const threadgroup matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup spvStorageMatrix& m) object_data = default;
    object_data spvStorageMatrix& operator=(const threadgroup matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    object_data spvStorageMatrix& operator=(const threadgroup spvStorageMatrix& m) object_data = default;
    
    #ifdef __HAVE_IMAGEBLOCKS__
    spvStorageMatrix(const threadgroup_imageblock matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const threadgroup_imageblock spvStorageMatrix& m) object_data = default;
    object_data spvStorageMatrix& operator=(const threadgroup_imageblock matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    object_data spvStorageMatrix& operator=(const threadgroup_imageblock spvStorageMatrix& m) object_data = default;
    #endif
    
    #ifdef __HAVE_RAYTRACING__
    spvStorageMatrix(const ray_data matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const ray_data spvStorageMatrix& m) object_data = default;
    object_data spvStorageMatrix& operator=(const ray_data matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    object_data spvStorageMatrix& operator=(const ray_data spvStorageMatrix& m) object_data = default;
    #endif
    
    spvStorageMatrix(const object_data matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
    }
    spvStorageMatrix(const object_data spvStorageMatrix& m) object_data = default;
    object_data spvStorageMatrix& operator=(const object_data matrix<T, Cols, Rows>& m) object_data
    {
        for (size_t i = 0; i < Cols; ++i)
            columns[i] = m.columns[i];
        return *this;
    }
    object_data spvStorageMatrix& operator=(const object_data spvStorageMatrix& m) object_data = default;
    
    operator matrix<T, Cols, Rows>() const object_data
    {
        matrix<T, Cols, Rows> m;
        for (int i = 0; i < Cols; ++i)
            m.columns[i] = columns[i];
        return m;
    }
    
    vec<T, Rows> operator[](size_t idx) const object_data
    {
        return columns[idx];
    }
    object_data vec<T, Rows>& operator[](size_t idx) object_data
    {
        return columns[idx];
    }
    #endif
    
};

template<typename T, int Cols, int Rows>
matrix<T, Rows, Cols> transpose(spvStorageMatrix<T, Cols, Rows> m)
{
    return transpose(matrix<T, Cols, Rows>(m));
}

typedef spvStorageMatrix<half, 2, 2> spvStorage_half2x2;
typedef spvStorageMatrix<half, 2, 3> spvStorage_half2x3;
typedef spvStorageMatrix<half, 2, 4> spvStorage_half2x4;
typedef spvStorageMatrix<half, 3, 2> spvStorage_half3x2;
typedef spvStorageMatrix<half, 3, 3> spvStorage_half3x3;
typedef spvStorageMatrix<half, 3, 4> spvStorage_half3x4;
typedef spvStorageMatrix<half, 4, 2> spvStorage_half4x2;
typedef spvStorageMatrix<half, 4, 3> spvStorage_half4x3;
typedef spvStorageMatrix<half, 4, 4> spvStorage_half4x4;
typedef spvStorageMatrix<float, 2, 2> spvStorage_float2x2;
typedef spvStorageMatrix<float, 2, 3> spvStorage_float2x3;
typedef spvStorageMatrix<float, 2, 4> spvStorage_float2x4;
typedef spvStorageMatrix<float, 3, 2> spvStorage_float3x2;
typedef spvStorageMatrix<float, 3, 3> spvStorage_float3x3;
typedef spvStorageMatrix<float, 3, 4> spvStorage_float3x4;
typedef spvStorageMatrix<float, 4, 2> spvStorage_float4x2;
typedef spvStorageMatrix<float, 4, 3> spvStorage_float4x3;
typedef spvStorageMatrix<float, 4, 4> spvStorage_float4x4;

struct S1
{
    spvStorage_float4x3 a[2];
    float b;
    spvUnsafeArray<float2, 3> c;
};

struct S2
{
    int4 a;
    spvUnsafeArray<spvUnsafeArray<spvUnsafeArray<short, 3>, 1>, 3> b;
};

struct block
{
    uint passed;
};

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

kernel void main0(device block& _383 [[buffer(0)]])
{
    threadgroup S1 s1;
    threadgroup S2 s2;
    s1.a[0] = spvStorage_float4x3(float4x3(float3(0.0, 2.0, -8.0), float3(6.0, 7.0, 5.0), float3(-6.0, 1.0, 9.0), float3(-4.0, -3.0, 4.0)));
    s1.a[1] = spvStorage_float4x3(float4x3(float3(4.0, 9.0, -9.0), float3(-8.0, -9.0, 8.0), float3(0.0, 4.0, -4.0), float3(7.0, 2.0, -1.0)));
    s1.b = 7.0;
    s1.c[0] = float2(-5.0, -4.0);
    s1.c[1] = float2(3.0, -5.0);
    s1.c[2] = float2(-3.0, -1.0);
    s2.a = int4(1, 0, -3, 1);
    s2.b[0][0][0] = short(true);
    s2.b[0][0][1] = short(false);
    s2.b[0][0][2] = short(false);
    s2.b[1][0][0] = short(true);
    s2.b[1][0][1] = short(false);
    s2.b[1][0][2] = short(true);
    s2.b[2][0][0] = short(false);
    s2.b[2][0][1] = short(true);
    s2.b[2][0][2] = short(true);
    threadgroup_barrier(mem_flags::mem_threadgroup);
    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
    bool _464 = abs(-float4x3(s1.a[0])[0].x) < 0.0500000007450580596923828125;
    bool _449;
    if (_464)
    {
        _449 = abs(2.0 - float4x3(s1.a[0])[0].y) < 0.0500000007450580596923828125;
    }
    else
    {
        _449 = _464;
    }
    bool _457;
    if (_449)
    {
        _457 = abs((-8.0) - float4x3(s1.a[0])[0].z) < 0.0500000007450580596923828125;
    }
    else
    {
        _457 = _449;
    }
    bool _412;
    if (_457)
    {
        bool _514 = abs(6.0 - float4x3(s1.a[0])[1].x) < 0.0500000007450580596923828125;
        bool _499;
        if (_514)
        {
            _499 = abs(7.0 - float4x3(s1.a[0])[1].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _499 = _514;
        }
        bool _507;
        if (_499)
        {
            _507 = abs(5.0 - float4x3(s1.a[0])[1].z) < 0.0500000007450580596923828125;
        }
        else
        {
            _507 = _499;
        }
        _412 = _507;
    }
    else
    {
        _412 = _457;
    }
    bool _420;
    if (_412)
    {
        bool _564 = abs((-6.0) - float4x3(s1.a[0])[2].x) < 0.0500000007450580596923828125;
        bool _549;
        if (_564)
        {
            _549 = abs(1.0 - float4x3(s1.a[0])[2].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _549 = _564;
        }
        bool _557;
        if (_549)
        {
            _557 = abs(9.0 - float4x3(s1.a[0])[2].z) < 0.0500000007450580596923828125;
        }
        else
        {
            _557 = _549;
        }
        _420 = _557;
    }
    else
    {
        _420 = _412;
    }
    bool _428;
    if (_420)
    {
        bool _614 = abs((-4.0) - float4x3(s1.a[0])[3].x) < 0.0500000007450580596923828125;
        bool _599;
        if (_614)
        {
            _599 = abs((-3.0) - float4x3(s1.a[0])[3].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _599 = _614;
        }
        bool _607;
        if (_599)
        {
            _607 = abs(4.0 - float4x3(s1.a[0])[3].z) < 0.0500000007450580596923828125;
        }
        else
        {
            _607 = _599;
        }
        _428 = _607;
    }
    else
    {
        _428 = _420;
    }
    bool _251;
    if (_428)
    {
        bool _703 = abs(4.0 - float4x3(s1.a[1])[0].x) < 0.0500000007450580596923828125;
        bool _688;
        if (_703)
        {
            _688 = abs(9.0 - float4x3(s1.a[1])[0].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _688 = _703;
        }
        bool _696;
        if (_688)
        {
            _696 = abs((-9.0) - float4x3(s1.a[1])[0].z) < 0.0500000007450580596923828125;
        }
        else
        {
            _696 = _688;
        }
        bool _651;
        if (_696)
        {
            bool _753 = abs((-8.0) - float4x3(s1.a[1])[1].x) < 0.0500000007450580596923828125;
            bool _738;
            if (_753)
            {
                _738 = abs((-9.0) - float4x3(s1.a[1])[1].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _738 = _753;
            }
            bool _746;
            if (_738)
            {
                _746 = abs(8.0 - float4x3(s1.a[1])[1].z) < 0.0500000007450580596923828125;
            }
            else
            {
                _746 = _738;
            }
            _651 = _746;
        }
        else
        {
            _651 = _696;
        }
        bool _659;
        if (_651)
        {
            bool _803 = abs(-float4x3(s1.a[1])[2].x) < 0.0500000007450580596923828125;
            bool _788;
            if (_803)
            {
                _788 = abs(4.0 - float4x3(s1.a[1])[2].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _788 = _803;
            }
            bool _796;
            if (_788)
            {
                _796 = abs((-4.0) - float4x3(s1.a[1])[2].z) < 0.0500000007450580596923828125;
            }
            else
            {
                _796 = _788;
            }
            _659 = _796;
        }
        else
        {
            _659 = _651;
        }
        bool _667;
        if (_659)
        {
            bool _853 = abs(7.0 - float4x3(s1.a[1])[3].x) < 0.0500000007450580596923828125;
            bool _838;
            if (_853)
            {
                _838 = abs(2.0 - float4x3(s1.a[1])[3].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _838 = _853;
            }
            bool _846;
            if (_838)
            {
                _846 = abs((-1.0) - float4x3(s1.a[1])[3].z) < 0.0500000007450580596923828125;
            }
            else
            {
                _846 = _838;
            }
            _667 = _846;
        }
        else
        {
            _667 = _659;
        }
        _251 = _667;
    }
    else
    {
        _251 = _428;
    }
    bool _260;
    if (_251)
    {
        _260 = abs(7.0 - s1.b) < 0.0500000007450580596923828125;
    }
    else
    {
        _260 = _251;
    }
    bool _269;
    if (_260)
    {
        bool _900 = abs((-5.0) - s1.c[0].x) < 0.0500000007450580596923828125;
        bool _893;
        if (_900)
        {
            _893 = abs((-4.0) - s1.c[0].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _893 = _900;
        }
        _269 = _893;
    }
    else
    {
        _269 = _260;
    }
    bool _278;
    if (_269)
    {
        bool _933 = abs(3.0 - s1.c[1].x) < 0.0500000007450580596923828125;
        bool _926;
        if (_933)
        {
            _926 = abs((-5.0) - s1.c[1].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _926 = _933;
        }
        _278 = _926;
    }
    else
    {
        _278 = _269;
    }
    bool _287;
    if (_278)
    {
        bool _966 = abs((-3.0) - s1.c[2].x) < 0.0500000007450580596923828125;
        bool _959;
        if (_966)
        {
            _959 = abs((-1.0) - s1.c[2].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _959 = _966;
        }
        _287 = _959;
    }
    else
    {
        _287 = _278;
    }
    bool _296;
    if (_287)
    {
        _296 = all(int4(1, 0, -3, 1) == s2.a);
    }
    else
    {
        _296 = _287;
    }
    bool _305;
    if (_296)
    {
        _305 = true == bool(s2.b[0][0][0]);
    }
    else
    {
        _305 = _296;
    }
    bool _314;
    if (_305)
    {
        _314 = false == bool(s2.b[0][0][1]);
    }
    else
    {
        _314 = _305;
    }
    bool _323;
    if (_314)
    {
        _323 = false == bool(s2.b[0][0][2]);
    }
    else
    {
        _323 = _314;
    }
    bool _332;
    if (_323)
    {
        _332 = true == bool(s2.b[1][0][0]);
    }
    else
    {
        _332 = _323;
    }
    bool _341;
    if (_332)
    {
        _341 = false == bool(s2.b[1][0][1]);
    }
    else
    {
        _341 = _332;
    }
    bool _350;
    if (_341)
    {
        _350 = true == bool(s2.b[1][0][2]);
    }
    else
    {
        _350 = _341;
    }
    bool _359;
    if (_350)
    {
        _359 = false == bool(s2.b[2][0][0]);
    }
    else
    {
        _359 = _350;
    }
    bool _368;
    if (_359)
    {
        _368 = true == bool(s2.b[2][0][1]);
    }
    else
    {
        _368 = _359;
    }
    bool _377;
    if (_368)
    {
        _377 = true == bool(s2.b[2][0][2]);
    }
    else
    {
        _377 = _368;
    }
    if (_377)
    {
        _383.passed++;
    }
}