#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
{
    uint a;
    float4 b;
};

struct sA
{
    spvStorage_float4x4 mA;
    short3 mB;
    short4 mC;
};

struct sB
{
    short2 mA;
};

struct sC
{
    float mA;
    uint4 mB;
    float mC;
};

struct sD
{
    sA mA;
    sB mB;
    sC mC;
};

struct sE
{
    sD mA;
};

struct sF
{
    uint3 mA;
    short mB;
};

struct sG
{
    sF mA;
    spvStorage_float3x2 mB;
};

struct sH
{
    sG mA;
    float2 mB;
};

struct sI
{
    spvStorage_float2x2 mA;
    short3 mB;
    short4 mC;
};

struct sJ
{
    sI mA;
    short3 mB;
};

struct sK
{
    short2 mA;
    sJ mB;
    int2 mC;
};

struct S2
{
    sE a;
    int3 b;
    sH c;
    sK d;
};

struct block
{
    uint passed;
};

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

kernel void main0(device block& _612 [[buffer(0)]])
{
    threadgroup S1 s1;
    threadgroup S2 s2;
    s1.a = 0u;
    s1.b = float4(8.0, 8.0, 0.0, -4.0);
    s2.a.mA.mA.mA = spvStorage_float4x4(float4x4(float4(-5.0, 9.0, -4.0, -6.0), float4(-1.0, -1.0, -2.0, 1.0), float4(6.0, 5.0, 7.0, -2.0), float4(-4.0, -9.0, 8.0, 3.0)));
    s2.a.mA.mA.mB = short3(bool3(true, false, false));
    s2.a.mA.mA.mC = short4(bool4(true, true, true, false));
    s2.a.mA.mB.mA = short2(bool2(true));
    s2.a.mA.mC.mA = 7.0;
    s2.a.mA.mC.mB = uint4(8u, 6u, 2u, 0u);
    s2.a.mA.mC.mC = -9.0;
    s2.b = int3(1, -4, 0);
    s2.c.mA.mA.mA = uint3(4u, 9u, 1u);
    s2.c.mA.mA.mB = short(false);
    s2.c.mA.mB = spvStorage_float3x2(float3x2(float2(3.0, -5.0), float2(-1.0, -5.0), float2(-1.0, -9.0)));
    s2.c.mB = float2(-6.0, -9.0);
    s2.d.mA = short2(bool2(true, false));
    s2.d.mB.mA.mA = spvStorage_float2x2(float2x2(float2(-2.0, 3.0), float2(7.0, 2.0)));
    s2.d.mB.mA.mB = short3(bool3(false));
    s2.d.mB.mA.mC = short4(bool4(false, false, false, true));
    s2.d.mB.mB = short3(bool3(true, false, false));
    s2.d.mC = int2(-9, 0);
    threadgroup_barrier(mem_flags::mem_threadgroup);
    threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
    bool _622 = 0u == s1.a;
    bool _444;
    if (_622)
    {
        bool _668 = abs(8.0 - s1.b.x) < 0.0500000007450580596923828125;
        bool _645;
        if (_668)
        {
            _645 = abs(8.0 - s1.b.y) < 0.0500000007450580596923828125;
        }
        else
        {
            _645 = _668;
        }
        bool _653;
        if (_645)
        {
            _653 = abs(-s1.b.z) < 0.0500000007450580596923828125;
        }
        else
        {
            _653 = _645;
        }
        bool _661;
        if (_653)
        {
            _661 = abs((-4.0) - s1.b.w) < 0.0500000007450580596923828125;
        }
        else
        {
            _661 = _653;
        }
        _444 = _661;
    }
    else
    {
        _444 = _622;
    }
    bool _453;
    if (_444)
    {
        bool _774 = abs((-5.0) - float4x4(s2.a.mA.mA.mA)[0].x) < 0.0500000007450580596923828125;
        bool _751;
        if (_774)
        {
            _751 = abs(9.0 - float4x4(s2.a.mA.mA.mA)[0].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _751 = _774;
        }
        bool _759;
        if (_751)
        {
            _759 = abs((-4.0) - float4x4(s2.a.mA.mA.mA)[0].z) < 0.0500000007450580596923828125;
        }
        else
        {
            _759 = _751;
        }
        bool _767;
        if (_759)
        {
            _767 = abs((-6.0) - float4x4(s2.a.mA.mA.mA)[0].w) < 0.0500000007450580596923828125;
        }
        else
        {
            _767 = _759;
        }
        bool _712;
        if (_767)
        {
            bool _841 = abs((-1.0) - float4x4(s2.a.mA.mA.mA)[1].x) < 0.0500000007450580596923828125;
            bool _818;
            if (_841)
            {
                _818 = abs((-1.0) - float4x4(s2.a.mA.mA.mA)[1].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _818 = _841;
            }
            bool _826;
            if (_818)
            {
                _826 = abs((-2.0) - float4x4(s2.a.mA.mA.mA)[1].z) < 0.0500000007450580596923828125;
            }
            else
            {
                _826 = _818;
            }
            bool _834;
            if (_826)
            {
                _834 = abs(1.0 - float4x4(s2.a.mA.mA.mA)[1].w) < 0.0500000007450580596923828125;
            }
            else
            {
                _834 = _826;
            }
            _712 = _834;
        }
        else
        {
            _712 = _767;
        }
        bool _720;
        if (_712)
        {
            bool _908 = abs(6.0 - float4x4(s2.a.mA.mA.mA)[2].x) < 0.0500000007450580596923828125;
            bool _885;
            if (_908)
            {
                _885 = abs(5.0 - float4x4(s2.a.mA.mA.mA)[2].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _885 = _908;
            }
            bool _893;
            if (_885)
            {
                _893 = abs(7.0 - float4x4(s2.a.mA.mA.mA)[2].z) < 0.0500000007450580596923828125;
            }
            else
            {
                _893 = _885;
            }
            bool _901;
            if (_893)
            {
                _901 = abs((-2.0) - float4x4(s2.a.mA.mA.mA)[2].w) < 0.0500000007450580596923828125;
            }
            else
            {
                _901 = _893;
            }
            _720 = _901;
        }
        else
        {
            _720 = _712;
        }
        bool _728;
        if (_720)
        {
            bool _975 = abs((-4.0) - float4x4(s2.a.mA.mA.mA)[3].x) < 0.0500000007450580596923828125;
            bool _952;
            if (_975)
            {
                _952 = abs((-9.0) - float4x4(s2.a.mA.mA.mA)[3].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _952 = _975;
            }
            bool _960;
            if (_952)
            {
                _960 = abs(8.0 - float4x4(s2.a.mA.mA.mA)[3].z) < 0.0500000007450580596923828125;
            }
            else
            {
                _960 = _952;
            }
            bool _968;
            if (_960)
            {
                _968 = abs(3.0 - float4x4(s2.a.mA.mA.mA)[3].w) < 0.0500000007450580596923828125;
            }
            else
            {
                _968 = _960;
            }
            _728 = _968;
        }
        else
        {
            _728 = _720;
        }
        _453 = _728;
    }
    else
    {
        _453 = _444;
    }
    bool _462;
    if (_453)
    {
        _462 = all(bool3(true, false, false) == bool3(s2.a.mA.mA.mB));
    }
    else
    {
        _462 = _453;
    }
    bool _471;
    if (_462)
    {
        _471 = all(bool4(true, true, true, false) == bool4(s2.a.mA.mA.mC));
    }
    else
    {
        _471 = _462;
    }
    bool _480;
    if (_471)
    {
        _480 = all(bool2(true) == bool2(s2.a.mA.mB.mA));
    }
    else
    {
        _480 = _471;
    }
    bool _489;
    if (_480)
    {
        _489 = abs(7.0 - s2.a.mA.mC.mA) < 0.0500000007450580596923828125;
    }
    else
    {
        _489 = _480;
    }
    bool _498;
    if (_489)
    {
        _498 = all(uint4(8u, 6u, 2u, 0u) == s2.a.mA.mC.mB);
    }
    else
    {
        _498 = _489;
    }
    bool _507;
    if (_498)
    {
        _507 = abs((-9.0) - s2.a.mA.mC.mC) < 0.0500000007450580596923828125;
    }
    else
    {
        _507 = _498;
    }
    bool _516;
    if (_507)
    {
        _516 = all(int3(1, -4, 0) == s2.b);
    }
    else
    {
        _516 = _507;
    }
    bool _525;
    if (_516)
    {
        _525 = all(uint3(4u, 9u, 1u) == s2.c.mA.mA.mA);
    }
    else
    {
        _525 = _516;
    }
    bool _534;
    if (_525)
    {
        _534 = false == bool(s2.c.mA.mA.mB);
    }
    else
    {
        _534 = _525;
    }
    bool _543;
    if (_534)
    {
        bool _1106 = abs(3.0 - float3x2(s2.c.mA.mB)[0].x) < 0.0500000007450580596923828125;
        bool _1099;
        if (_1106)
        {
            _1099 = abs((-5.0) - float3x2(s2.c.mA.mB)[0].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _1099 = _1106;
        }
        bool _1072;
        if (_1099)
        {
            bool _1139 = abs((-1.0) - float3x2(s2.c.mA.mB)[1].x) < 0.0500000007450580596923828125;
            bool _1132;
            if (_1139)
            {
                _1132 = abs((-5.0) - float3x2(s2.c.mA.mB)[1].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _1132 = _1139;
            }
            _1072 = _1132;
        }
        else
        {
            _1072 = _1099;
        }
        bool _1080;
        if (_1072)
        {
            bool _1172 = abs((-1.0) - float3x2(s2.c.mA.mB)[2].x) < 0.0500000007450580596923828125;
            bool _1165;
            if (_1172)
            {
                _1165 = abs((-9.0) - float3x2(s2.c.mA.mB)[2].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _1165 = _1172;
            }
            _1080 = _1165;
        }
        else
        {
            _1080 = _1072;
        }
        _543 = _1080;
    }
    else
    {
        _543 = _534;
    }
    bool _552;
    if (_543)
    {
        bool _1205 = abs((-6.0) - s2.c.mB.x) < 0.0500000007450580596923828125;
        bool _1198;
        if (_1205)
        {
            _1198 = abs((-9.0) - s2.c.mB.y) < 0.0500000007450580596923828125;
        }
        else
        {
            _1198 = _1205;
        }
        _552 = _1198;
    }
    else
    {
        _552 = _543;
    }
    bool _561;
    if (_552)
    {
        _561 = all(bool2(true, false) == bool2(s2.d.mA));
    }
    else
    {
        _561 = _552;
    }
    bool _570;
    if (_561)
    {
        bool _1263 = abs((-2.0) - float2x2(s2.d.mB.mA.mA)[0].x) < 0.0500000007450580596923828125;
        bool _1256;
        if (_1263)
        {
            _1256 = abs(3.0 - float2x2(s2.d.mB.mA.mA)[0].y) < 0.0500000007450580596923828125;
        }
        else
        {
            _1256 = _1263;
        }
        bool _1237;
        if (_1256)
        {
            bool _1296 = abs(7.0 - float2x2(s2.d.mB.mA.mA)[1].x) < 0.0500000007450580596923828125;
            bool _1289;
            if (_1296)
            {
                _1289 = abs(2.0 - float2x2(s2.d.mB.mA.mA)[1].y) < 0.0500000007450580596923828125;
            }
            else
            {
                _1289 = _1296;
            }
            _1237 = _1289;
        }
        else
        {
            _1237 = _1256;
        }
        _570 = _1237;
    }
    else
    {
        _570 = _561;
    }
    bool _579;
    if (_570)
    {
        _579 = all(bool3(false) == bool3(s2.d.mB.mA.mB));
    }
    else
    {
        _579 = _570;
    }
    bool _588;
    if (_579)
    {
        _588 = all(bool4(false, false, false, true) == bool4(s2.d.mB.mA.mC));
    }
    else
    {
        _588 = _579;
    }
    bool _597;
    if (_588)
    {
        _597 = all(bool3(true, false, false) == bool3(s2.d.mB.mB));
    }
    else
    {
        _597 = _588;
    }
    bool _606;
    if (_597)
    {
        _606 = all(int2(-9, 0) == s2.d.mC);
    }
    else
    {
        _606 = _597;
    }
    if (_606)
    {
        _612.passed++;
    }
}