#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);

static inline __attribute__((always_inline))
bool compare_uint(thread const uint& a, thread const uint& b)
{
    return a == b;
}

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 _147 = compare_float(param, param_1);
    bool _157;
    if (_147)
    {
        float param_2 = a.y;
        float param_3 = b.y;
        _157 = compare_float(param_2, param_3);
    }
    else
    {
        _157 = _147;
    }
    bool _168;
    if (_157)
    {
        float param_4 = a.z;
        float param_5 = b.z;
        _168 = compare_float(param_4, param_5);
    }
    else
    {
        _168 = _157;
    }
    bool _179;
    if (_168)
    {
        float param_6 = a.w;
        float param_7 = b.w;
        _179 = compare_float(param_6, param_7);
    }
    else
    {
        _179 = _168;
    }
    return _179;
}

static inline __attribute__((always_inline))
bool compare_mat4(thread const float4x4& a, thread const float4x4& b)
{
    float4 param = a[0];
    float4 param_1 = b[0];
    bool _239 = compare_vec4(param, param_1);
    bool _249;
    if (_239)
    {
        float4 param_2 = a[1];
        float4 param_3 = b[1];
        _249 = compare_vec4(param_2, param_3);
    }
    else
    {
        _249 = _239;
    }
    bool _259;
    if (_249)
    {
        float4 param_4 = a[2];
        float4 param_5 = b[2];
        _259 = compare_vec4(param_4, param_5);
    }
    else
    {
        _259 = _249;
    }
    bool _270;
    if (_259)
    {
        float4 param_6 = a[3];
        float4 param_7 = b[3];
        _270 = compare_vec4(param_6, param_7);
    }
    else
    {
        _270 = _259;
    }
    return _270;
}

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

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

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

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

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

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

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

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 _127 = compare_float(param, param_1);
    bool _138;
    if (_127)
    {
        float param_2 = a.y;
        float param_3 = b.y;
        _138 = compare_float(param_2, param_3);
    }
    else
    {
        _138 = _127;
    }
    return _138;
}

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 _209 = compare_vec2(param, param_1);
    bool _219;
    if (_209)
    {
        float2 param_2 = a[1];
        float2 param_3 = b[1];
        _219 = compare_vec2(param_2, param_3);
    }
    else
    {
        _219 = _209;
    }
    bool _230;
    if (_219)
    {
        float2 param_4 = a[2];
        float2 param_5 = b[2];
        _230 = compare_vec2(param_4, param_5);
    }
    else
    {
        _230 = _219;
    }
    return _230;
}

static inline __attribute__((always_inline))
bool compare_mat2(thread const float2x2& a, thread const float2x2& b)
{
    float2 param = a[0];
    float2 param_1 = b[0];
    bool _189 = compare_vec2(param, param_1);
    bool _200;
    if (_189)
    {
        float2 param_2 = a[1];
        float2 param_3 = b[1];
        _200 = compare_vec2(param_2, param_3);
    }
    else
    {
        _200 = _189;
    }
    return _200;
}

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

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 allOk = true;
    bool _435;
    if (allOk)
    {
        uint param = 0u;
        uint param_1 = s1.a;
        _435 = compare_uint(param, param_1);
    }
    else
    {
        _435 = allOk;
    }
    allOk = _435;
    bool _444;
    if (allOk)
    {
        float4 param_2 = float4(8.0, 8.0, 0.0, -4.0);
        float4 param_3 = s1.b;
        _444 = compare_vec4(param_2, param_3);
    }
    else
    {
        _444 = allOk;
    }
    allOk = _444;
    bool _453;
    if (allOk)
    {
        float4x4 param_4 = 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));
        float4x4 param_5 = float4x4(s2.a.mA.mA.mA);
        _453 = compare_mat4(param_4, param_5);
    }
    else
    {
        _453 = allOk;
    }
    allOk = _453;
    bool _462;
    if (allOk)
    {
        bool3 param_6 = bool3(true, false, false);
        bool3 param_7 = bool3(s2.a.mA.mA.mB);
        _462 = compare_bvec3(param_6, param_7);
    }
    else
    {
        _462 = allOk;
    }
    allOk = _462;
    bool _471;
    if (allOk)
    {
        bool4 param_8 = bool4(true, true, true, false);
        bool4 param_9 = bool4(s2.a.mA.mA.mC);
        _471 = compare_bvec4(param_8, param_9);
    }
    else
    {
        _471 = allOk;
    }
    allOk = _471;
    bool _480;
    if (allOk)
    {
        bool2 param_10 = bool2(true);
        bool2 param_11 = bool2(s2.a.mA.mB.mA);
        _480 = compare_bvec2(param_10, param_11);
    }
    else
    {
        _480 = allOk;
    }
    allOk = _480;
    bool _489;
    if (allOk)
    {
        float param_12 = 7.0;
        float param_13 = s2.a.mA.mC.mA;
        _489 = compare_float(param_12, param_13);
    }
    else
    {
        _489 = allOk;
    }
    allOk = _489;
    bool _498;
    if (allOk)
    {
        uint4 param_14 = uint4(8u, 6u, 2u, 0u);
        uint4 param_15 = s2.a.mA.mC.mB;
        _498 = compare_uvec4(param_14, param_15);
    }
    else
    {
        _498 = allOk;
    }
    allOk = _498;
    bool _507;
    if (allOk)
    {
        float param_16 = -9.0;
        float param_17 = s2.a.mA.mC.mC;
        _507 = compare_float(param_16, param_17);
    }
    else
    {
        _507 = allOk;
    }
    allOk = _507;
    bool _516;
    if (allOk)
    {
        int3 param_18 = int3(1, -4, 0);
        int3 param_19 = s2.b;
        _516 = compare_ivec3(param_18, param_19);
    }
    else
    {
        _516 = allOk;
    }
    allOk = _516;
    bool _525;
    if (allOk)
    {
        uint3 param_20 = uint3(4u, 9u, 1u);
        uint3 param_21 = s2.c.mA.mA.mA;
        _525 = compare_uvec3(param_20, param_21);
    }
    else
    {
        _525 = allOk;
    }
    allOk = _525;
    bool _534;
    if (allOk)
    {
        bool param_22 = false;
        bool param_23 = bool(s2.c.mA.mA.mB);
        _534 = compare_bool(param_22, param_23);
    }
    else
    {
        _534 = allOk;
    }
    allOk = _534;
    bool _543;
    if (allOk)
    {
        float3x2 param_24 = float3x2(float2(3.0, -5.0), float2(-1.0, -5.0), float2(-1.0, -9.0));
        float3x2 param_25 = float3x2(s2.c.mA.mB);
        _543 = compare_mat3x2(param_24, param_25);
    }
    else
    {
        _543 = allOk;
    }
    allOk = _543;
    bool _552;
    if (allOk)
    {
        float2 param_26 = float2(-6.0, -9.0);
        float2 param_27 = s2.c.mB;
        _552 = compare_vec2(param_26, param_27);
    }
    else
    {
        _552 = allOk;
    }
    allOk = _552;
    bool _561;
    if (allOk)
    {
        bool2 param_28 = bool2(true, false);
        bool2 param_29 = bool2(s2.d.mA);
        _561 = compare_bvec2(param_28, param_29);
    }
    else
    {
        _561 = allOk;
    }
    allOk = _561;
    bool _570;
    if (allOk)
    {
        float2x2 param_30 = float2x2(float2(-2.0, 3.0), float2(7.0, 2.0));
        float2x2 param_31 = float2x2(s2.d.mB.mA.mA);
        _570 = compare_mat2(param_30, param_31);
    }
    else
    {
        _570 = allOk;
    }
    allOk = _570;
    bool _579;
    if (allOk)
    {
        bool3 param_32 = bool3(false);
        bool3 param_33 = bool3(s2.d.mB.mA.mB);
        _579 = compare_bvec3(param_32, param_33);
    }
    else
    {
        _579 = allOk;
    }
    allOk = _579;
    bool _588;
    if (allOk)
    {
        bool4 param_34 = bool4(false, false, false, true);
        bool4 param_35 = bool4(s2.d.mB.mA.mC);
        _588 = compare_bvec4(param_34, param_35);
    }
    else
    {
        _588 = allOk;
    }
    allOk = _588;
    bool _597;
    if (allOk)
    {
        bool3 param_36 = bool3(true, false, false);
        bool3 param_37 = bool3(s2.d.mB.mB);
        _597 = compare_bvec3(param_36, param_37);
    }
    else
    {
        _597 = allOk;
    }
    allOk = _597;
    bool _606;
    if (allOk)
    {
        int2 param_38 = int2(-9, 0);
        int2 param_39 = s2.d.mC;
        _606 = compare_ivec2(param_38, param_39);
    }
    else
    {
        _606 = allOk;
    }
    allOk = _606;
    if (allOk)
    {
        _612.passed++;
    }
}