AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity
c084973c67
SPIRV-Cross/reference/shaders-msl/comp/shared-matrix-nested-struct.comp
Chip Davis fc4a12fd4f MSL: Use a wrapper type for matrices in workgroup storage. 
The standard `matrix` type in MSL lacked a constructor in the
`threadgroup` AS. This means that it was impossible to declare a
`threadgroup` variable that contains a matrix. This appears to have been
an oversight that was corrected in macOS 13/Xcode 14 beta 4. This
workaround continues to be required, however, for older systems.

To avoid changing interfaces unnecessarily (which shouldn't be a problem
regardless because the old and new types take up the same amount of
storage), only do this for structs if the struct is positively
identified as being used for workgroup storage.

I'm entirely aware this is inconsistent with the way packed matrices are
handled. One of them should be changed to match the other. Not sure
which one.

Fixes 23 CTS tests under `dEQP-VK.memory_model.shared`.
2022-08-07 17:31:41 -07:00

1474 lines
43 KiB
Plaintext
Raw Blame History

#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++;
}
}
Reference in New Issue View Git Blame Copy Permalink