1444 lines
43 KiB
Plaintext
1444 lines
43 KiB
Plaintext
|
#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++;
|
||
|
}
|
||
|
}
|
||
|
|