#pragma clang diagnostic ignored "-Wmissing-prototypes"

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

using namespace metal;

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
{
    float4 a;
    spvStorage_float3x2 b;
    short4 c;
};

struct block
{
    uint passed;
};

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

static inline __attribute__((always_inline))
bool compare_float(thread const float& a, thread const float& b)
{
    return abs(a - b) < 0.0500000007450580596923828125;
}

static inline __attribute__((always_inline))
bool compare_vec4(thread const float4& a, thread const float4& b)
{
    float param = a.x;
    float param_1 = b.x;
    bool _78 = compare_float(param, param_1);
    bool _88;
    if (_78)
    {
        float param_2 = a.y;
        float param_3 = b.y;
        _88 = compare_float(param_2, param_3);
    }
    else
    {
        _88 = _78;
    }
    bool _99;
    if (_88)
    {
        float param_4 = a.z;
        float param_5 = b.z;
        _99 = compare_float(param_4, param_5);
    }
    else
    {
        _99 = _88;
    }
    bool _110;
    if (_99)
    {
        float param_6 = a.w;
        float param_7 = b.w;
        _110 = compare_float(param_6, param_7);
    }
    else
    {
        _110 = _99;
    }
    return _110;
}

static inline __attribute__((always_inline))
bool compare_vec2(thread const float2& a, thread const float2& b)
{
    float param = a.x;
    float param_1 = b.x;
    bool _58 = compare_float(param, param_1);
    bool _69;
    if (_58)
    {
        float param_2 = a.y;
        float param_3 = b.y;
        _69 = compare_float(param_2, param_3);
    }
    else
    {
        _69 = _58;
    }
    return _69;
}

static inline __attribute__((always_inline))
bool compare_mat3x2(thread const float3x2& a, thread const float3x2& b)
{
    float2 param = a[0];
    float2 param_1 = b[0];
    bool _121 = compare_vec2(param, param_1);
    bool _132;
    if (_121)
    {
        float2 param_2 = a[1];
        float2 param_3 = b[1];
        _132 = compare_vec2(param_2, param_3);
    }
    else
    {
        _132 = _121;
    }
    bool _143;
    if (_132)
    {
        float2 param_4 = a[2];
        float2 param_5 = b[2];
        _143 = compare_vec2(param_4, param_5);
    }
    else
    {
        _143 = _132;
    }
    return _143;
}

static inline __attribute__((always_inline))
bool compare_bvec4(thread const bool4& a, thread const bool4& b)
{
    return all(a == b);
}

kernel void main0(device block& _212 [[buffer(0)]])
{
    threadgroup S1 s1;
    s1.a = float4(1.0, -5.0, -9.0, -5.0);
    s1.b = spvStorage_float3x2(float3x2(float2(1.0, -7.0), float2(1.0, 2.0), float2(8.0, 7.0)));
    s1.c = short4(bool4(false, true, false, false));
    threadgroup_barrier(mem_flags::mem_threadgroup);
    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
    bool allOk = true;
    bool _188;
    if (allOk)
    {
        float4 param = float4(1.0, -5.0, -9.0, -5.0);
        float4 param_1 = s1.a;
        _188 = compare_vec4(param, param_1);
    }
    else
    {
        _188 = allOk;
    }
    allOk = _188;
    bool _197;
    if (allOk)
    {
        float3x2 param_2 = float3x2(float2(1.0, -7.0), float2(1.0, 2.0), float2(8.0, 7.0));
        float3x2 param_3 = float3x2(s1.b);
        _197 = compare_mat3x2(param_2, param_3);
    }
    else
    {
        _197 = allOk;
    }
    allOk = _197;
    bool _206;
    if (allOk)
    {
        bool4 param_4 = bool4(false, true, false, false);
        bool4 param_5 = bool4(s1.c);
        _206 = compare_bvec4(param_4, param_5);
    }
    else
    {
        _206 = allOk;
    }
    allOk = _206;
    if (allOk)
    {
        _212.passed++;
    }
}