Use recursive template for spvArrayCopy
Comment claims we can't, but I tested a number of older Metal compilers (Xcode 8 targeting macOS 10.11, macOS 10.13 online compiler, Xcode 14 targeting iOS 8) and none of them had any issues with it
This commit is contained in:
Evan Tang
parent
7d92d7d879
commit
cc542f8147
@ -5,109 +5,109 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -141,7 +141,7 @@ kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadg
|
||||
|
||||
Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
|
||||
Data data2[2];
|
||||
spvArrayCopyFromStackToStack1(data2, _31);
|
||||
spvArrayCopyFromStackToStack(data2, _31);
|
||||
_53.outdata[gl_WorkGroupID.x].a = _25[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a;
|
||||
_53.outdata[gl_WorkGroupID.x].b = _25[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b;
|
||||
}
|
||||
|
||||
@ -44,109 +44,109 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -167,8 +167,8 @@ kernel void main0(device SSBO& ssbo [[buffer(0)]])
|
||||
{
|
||||
threadgroup spvUnsafeArray<uint2, 2> _18;
|
||||
spvUnsafeArray<uint2, 2> _27;
|
||||
spvArrayCopyFromDeviceToStack1(_27.elements, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _27.elements);
|
||||
spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _27.elements);
|
||||
spvArrayCopyFromDeviceToStack(_27.elements, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromStackToDevice(ssbo._m0[0u]._m0, _27.elements);
|
||||
spvArrayCopyFromStackToDevice(ssbo._m0[0u]._m0, _27.elements);
|
||||
}
|
||||
|
||||
|
||||
@ -5,109 +5,109 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -128,8 +128,8 @@ kernel void main0(device SSBO& ssbo [[buffer(0)]])
|
||||
{
|
||||
threadgroup uint2 _18[2];
|
||||
uint2 _27[2];
|
||||
spvArrayCopyFromDeviceToStack1(_27, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _27);
|
||||
spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _27);
|
||||
spvArrayCopyFromDeviceToStack(_27, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromStackToDevice(ssbo._m0[0u]._m0, _27);
|
||||
spvArrayCopyFromStackToDevice(ssbo._m0[0u]._m0, _27);
|
||||
}
|
||||
|
||||
|
||||
@ -44,109 +44,109 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -166,14 +166,14 @@ struct SSBO
|
||||
kernel void main0(device SSBO& ssbo [[buffer(0)]], constant SSBO& ubo [[buffer(1)]])
|
||||
{
|
||||
threadgroup spvUnsafeArray<uint2, 2> _18;
|
||||
spvArrayCopyFromDeviceToDevice1(ssbo._m0[0u]._m0, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToDevice1(ssbo._m0[0u]._m0, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromDeviceToDevice(ssbo._m0[0u]._m0, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToDevice(ssbo._m0[0u]._m0, ubo._m0[0u]._m1);
|
||||
spvUnsafeArray<uint2, 2> _23;
|
||||
spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _23.elements);
|
||||
spvArrayCopyFromThreadGroupToDevice1(ssbo._m0[0u]._m0, _18.elements);
|
||||
spvArrayCopyFromDeviceToThreadGroup1(_18.elements, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromDeviceToStack1(_23.elements, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToThreadGroup1(_18.elements, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToStack1(_23.elements, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromStackToDevice(ssbo._m0[0u]._m0, _23.elements);
|
||||
spvArrayCopyFromThreadGroupToDevice(ssbo._m0[0u]._m0, _18.elements);
|
||||
spvArrayCopyFromDeviceToThreadGroup(_18.elements, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromDeviceToStack(_23.elements, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToThreadGroup(_18.elements, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToStack(_23.elements, ubo._m0[0u]._m1);
|
||||
}
|
||||
|
||||
|
||||
@ -5,109 +5,109 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -127,14 +127,14 @@ struct SSBO
|
||||
kernel void main0(device SSBO& ssbo [[buffer(0)]], constant SSBO& ubo [[buffer(1)]])
|
||||
{
|
||||
threadgroup uint2 _18[2];
|
||||
spvArrayCopyFromDeviceToDevice1(ssbo._m0[0u]._m0, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToDevice1(ssbo._m0[0u]._m0, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromDeviceToDevice(ssbo._m0[0u]._m0, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToDevice(ssbo._m0[0u]._m0, ubo._m0[0u]._m1);
|
||||
uint2 _23[2];
|
||||
spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _23);
|
||||
spvArrayCopyFromThreadGroupToDevice1(ssbo._m0[0u]._m0, _18);
|
||||
spvArrayCopyFromDeviceToThreadGroup1(_18, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromDeviceToStack1(_23, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToThreadGroup1(_18, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToStack1(_23, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromStackToDevice(ssbo._m0[0u]._m0, _23);
|
||||
spvArrayCopyFromThreadGroupToDevice(ssbo._m0[0u]._m0, _18);
|
||||
spvArrayCopyFromDeviceToThreadGroup(_18, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromDeviceToStack(_23, ssbo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToThreadGroup(_18, ubo._m0[0u]._m1);
|
||||
spvArrayCopyFromConstantToStack(_23, ubo._m0[0u]._m1);
|
||||
}
|
||||
|
||||
|
||||
@ -44,109 +44,109 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -165,6 +165,6 @@ fragment void main0()
|
||||
_34[2u] = 0.0;
|
||||
_34[3u] = 0.0;
|
||||
_3 _33;
|
||||
spvArrayCopyFromStackToStack1(_33._m0, _34.elements);
|
||||
spvArrayCopyFromStackToStack(_33._m0, _34.elements);
|
||||
}
|
||||
|
||||
|
||||
@ -5,109 +5,109 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -130,9 +130,9 @@ static inline __attribute__((always_inline))
|
||||
float4 consume_constant_arrays2(thread const float4 (&positions)[4], thread const float4 (&positions2)[4], thread int& Index1, thread int& Index2)
|
||||
{
|
||||
float4 indexable[4];
|
||||
spvArrayCopyFromStackToStack1(indexable, positions);
|
||||
spvArrayCopyFromStackToStack(indexable, positions);
|
||||
float4 indexable_1[4];
|
||||
spvArrayCopyFromStackToStack1(indexable_1, positions2);
|
||||
spvArrayCopyFromStackToStack(indexable_1, positions2);
|
||||
return indexable[Index1] + indexable_1[Index2];
|
||||
}
|
||||
|
||||
|
||||
@ -5,109 +5,109 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -148,7 +148,7 @@ kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadg
|
||||
Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
|
||||
Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
|
||||
Data data2[2];
|
||||
spvArrayCopyFromStackToStack1(data2, _31);
|
||||
spvArrayCopyFromStackToStack(data2, _31);
|
||||
Data param = data[gl_LocalInvocationID.x];
|
||||
Data param_1 = data2[gl_LocalInvocationID.x];
|
||||
Data _73 = combine(param, param_1);
|
||||
|
||||
@ -5,327 +5,219 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromConstantToStack2(thread T (&dst)[A][B], constant T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N][M], constant T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromConstantToStack1(dst[i], src[i]);
|
||||
spvArrayCopyFromConstantToStack(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup2(threadgroup T (&dst)[A][B], constant T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N][M], constant T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromConstantToThreadGroup1(dst[i], src[i]);
|
||||
spvArrayCopyFromConstantToThreadGroup(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromStackToStack2(thread T (&dst)[A][B], thread const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N][M], thread const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromStackToStack1(dst[i], src[i]);
|
||||
spvArrayCopyFromStackToStack(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromStackToThreadGroup2(threadgroup T (&dst)[A][B], thread const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N][M], thread const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromStackToThreadGroup1(dst[i], src[i]);
|
||||
spvArrayCopyFromStackToThreadGroup(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromThreadGroupToStack2(thread T (&dst)[A][B], threadgroup const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N][M], threadgroup const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromThreadGroupToStack1(dst[i], src[i]);
|
||||
spvArrayCopyFromThreadGroupToStack(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup2(threadgroup T (&dst)[A][B], threadgroup const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N][M], threadgroup const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromThreadGroupToThreadGroup1(dst[i], src[i]);
|
||||
spvArrayCopyFromThreadGroupToThreadGroup(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromDeviceToDevice2(device T (&dst)[A][B], device const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N][M], device const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromDeviceToDevice1(dst[i], src[i]);
|
||||
spvArrayCopyFromDeviceToDevice(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromConstantToDevice2(device T (&dst)[A][B], constant T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N][M], constant T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromConstantToDevice1(dst[i], src[i]);
|
||||
spvArrayCopyFromConstantToDevice(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromStackToDevice2(device T (&dst)[A][B], thread const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N][M], thread const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromStackToDevice1(dst[i], src[i]);
|
||||
spvArrayCopyFromStackToDevice(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice2(device T (&dst)[A][B], threadgroup const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N][M], threadgroup const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromThreadGroupToDevice1(dst[i], src[i]);
|
||||
spvArrayCopyFromThreadGroupToDevice(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromDeviceToStack2(thread T (&dst)[A][B], device const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N][M], device const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromDeviceToStack1(dst[i], src[i]);
|
||||
spvArrayCopyFromDeviceToStack(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup2(threadgroup T (&dst)[A][B], device const T (&src)[A][B])
|
||||
template<typename T, uint N, uint M>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N][M], device const T (&src)[N][M])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
spvArrayCopyFromDeviceToThreadGroup1(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromConstantToStack3(thread T (&dst)[A][B][C], constant T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromConstantToStack2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup3(threadgroup T (&dst)[A][B][C], constant T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromConstantToThreadGroup2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromStackToStack3(thread T (&dst)[A][B][C], thread const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromStackToStack2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromStackToThreadGroup3(threadgroup T (&dst)[A][B][C], thread const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromStackToThreadGroup2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromThreadGroupToStack3(thread T (&dst)[A][B][C], threadgroup const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromThreadGroupToStack2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup3(threadgroup T (&dst)[A][B][C], threadgroup const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromThreadGroupToThreadGroup2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromDeviceToDevice3(device T (&dst)[A][B][C], device const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromDeviceToDevice2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromConstantToDevice3(device T (&dst)[A][B][C], constant T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromConstantToDevice2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromStackToDevice3(device T (&dst)[A][B][C], thread const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromStackToDevice2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice3(device T (&dst)[A][B][C], threadgroup const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromThreadGroupToDevice2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromDeviceToStack3(thread T (&dst)[A][B][C], device const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromDeviceToStack2(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A, uint B, uint C>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup3(threadgroup T (&dst)[A][B][C], device const T (&src)[A][B][C])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
{
|
||||
spvArrayCopyFromDeviceToThreadGroup2(dst[i], src[i]);
|
||||
spvArrayCopyFromDeviceToThreadGroup(dst[i], src[i]);
|
||||
}
|
||||
}
|
||||
|
||||
@ -346,7 +238,7 @@ constant float _21[2][2][2] = { { { 1.0, 2.0 }, { 3.0, 4.0 } }, { { 1.0, 2.0 },
|
||||
kernel void main0(device BUF& o [[buffer(0)]])
|
||||
{
|
||||
float c[2][2][2];
|
||||
spvArrayCopyFromConstantToStack3(c, _21);
|
||||
spvArrayCopyFromConstantToStack(c, _21);
|
||||
o.a = int(c[1][1][1]);
|
||||
float _43[2] = { o.b, o.c };
|
||||
float _48[2] = { o.b, o.b };
|
||||
@ -356,9 +248,9 @@ kernel void main0(device BUF& o [[buffer(0)]])
|
||||
float _60[2][2] = { { _54[0], _54[1] }, { _59[0], _59[1] } };
|
||||
float _61[2][2][2] = { { { _49[0][0], _49[0][1] }, { _49[1][0], _49[1][1] } }, { { _60[0][0], _60[0][1] }, { _60[1][0], _60[1][1] } } };
|
||||
float d[2][2][2];
|
||||
spvArrayCopyFromStackToStack3(d, _61);
|
||||
spvArrayCopyFromStackToStack(d, _61);
|
||||
float e[2][2][2];
|
||||
spvArrayCopyFromStackToStack3(e, d);
|
||||
spvArrayCopyFromStackToStack(e, d);
|
||||
o.b = e[1][0][1];
|
||||
}
|
||||
|
||||
|
||||
@ -5,109 +5,109 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
}
|
||||
|
||||
template<typename T, uint A>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
|
||||
template<typename T, uint N>
|
||||
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
|
||||
{
|
||||
for (uint i = 0; i < A; i++)
|
||||
for (uint i = 0; i < N; i++)
|
||||
{
|
||||
dst[i] = src[i];
|
||||
}
|
||||
@ -129,7 +129,7 @@ struct main0_in
|
||||
static inline __attribute__((always_inline))
|
||||
void test(thread float4 (&spvReturnValue)[2])
|
||||
{
|
||||
spvArrayCopyFromConstantToStack1(spvReturnValue, _20);
|
||||
spvArrayCopyFromConstantToStack(spvReturnValue, _20);
|
||||
}
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
@ -138,7 +138,7 @@ void test2(thread float4 (&spvReturnValue)[2], thread float4& vInput0, thread fl
|
||||
float4 foobar[2];
|
||||
foobar[0] = vInput0;
|
||||
foobar[1] = vInput1;
|
||||
spvArrayCopyFromStackToStack1(spvReturnValue, foobar);
|
||||
spvArrayCopyFromStackToStack(spvReturnValue, foobar);
|
||||
}
|
||||
|
||||
vertex main0_out main0(main0_in in [[stage_in]])
|
||||
|
||||
@ -5561,9 +5561,8 @@ void CompilerMSL::emit_custom_templates()
|
||||
// otherwise they will cause problems when linked together in a single Metallib.
|
||||
void CompilerMSL::emit_custom_functions()
|
||||
{
|
||||
for (uint32_t i = kArrayCopyMultidimMax; i >= 2; i--)
|
||||
if (spv_function_implementations.count(static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + i)))
|
||||
spv_function_implementations.insert(static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + i - 1));
|
||||
if (spv_function_implementations.count(SPVFuncImplArrayCopyMultidim))
|
||||
spv_function_implementations.insert(SPVFuncImplArrayCopy);
|
||||
|
||||
if (spv_function_implementations.count(SPVFuncImplDynamicImageSampler))
|
||||
{
|
||||
@ -5674,11 +5673,7 @@ void CompilerMSL::emit_custom_functions()
|
||||
break;
|
||||
|
||||
case SPVFuncImplArrayCopy:
|
||||
case SPVFuncImplArrayOfArrayCopy2Dim:
|
||||
case SPVFuncImplArrayOfArrayCopy3Dim:
|
||||
case SPVFuncImplArrayOfArrayCopy4Dim:
|
||||
case SPVFuncImplArrayOfArrayCopy5Dim:
|
||||
case SPVFuncImplArrayOfArrayCopy6Dim:
|
||||
case SPVFuncImplArrayCopyMultidim:
|
||||
{
|
||||
// Unfortunately we cannot template on the address space, so combinatorial explosion it is.
|
||||
static const char *function_name_tags[] = {
|
||||
@ -5701,36 +5696,19 @@ void CompilerMSL::emit_custom_functions()
|
||||
|
||||
for (uint32_t variant = 0; variant < 12; variant++)
|
||||
{
|
||||
uint8_t dimensions = spv_func - SPVFuncImplArrayCopyMultidimBase;
|
||||
string tmp = "template<typename T";
|
||||
for (uint8_t i = 0; i < dimensions; i++)
|
||||
{
|
||||
tmp += ", uint ";
|
||||
tmp += 'A' + i;
|
||||
}
|
||||
tmp += ">";
|
||||
statement(tmp);
|
||||
|
||||
string array_arg;
|
||||
for (uint8_t i = 0; i < dimensions; i++)
|
||||
{
|
||||
array_arg += "[";
|
||||
array_arg += 'A' + i;
|
||||
array_arg += "]";
|
||||
}
|
||||
|
||||
statement("inline void spvArrayCopy", function_name_tags[variant], dimensions, "(",
|
||||
dst_address_space[variant], " T (&dst)", array_arg, ", ", src_address_space[variant],
|
||||
" T (&src)", array_arg, ")");
|
||||
|
||||
bool is_multidim = spv_func == SPVFuncImplArrayCopyMultidim;
|
||||
const char* dim = is_multidim ? "[N][M]" : "[N]";
|
||||
statement("template<typename T, uint N", is_multidim ? ", uint M>" : ">");
|
||||
statement("inline void spvArrayCopy", function_name_tags[variant], "(",
|
||||
dst_address_space[variant], " T (&dst)", dim, ", ",
|
||||
src_address_space[variant], " T (&src)", dim, ")");
|
||||
begin_scope();
|
||||
statement("for (uint i = 0; i < A; i++)");
|
||||
statement("for (uint i = 0; i < N; i++)");
|
||||
begin_scope();
|
||||
|
||||
if (dimensions == 1)
|
||||
statement("dst[i] = src[i];");
|
||||
if (is_multidim)
|
||||
statement("spvArrayCopy", function_name_tags[variant], "(dst[i], src[i]);");
|
||||
else
|
||||
statement("spvArrayCopy", function_name_tags[variant], dimensions - 1, "(dst[i], src[i]);");
|
||||
statement("dst[i] = src[i];");
|
||||
end_scope();
|
||||
end_scope();
|
||||
statement("");
|
||||
@ -10006,15 +9984,7 @@ bool CompilerMSL::emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rh
|
||||
// we cannot easily detect this case ahead of time since it's
|
||||
// context dependent. We might have to force a recompile here
|
||||
// if this is the only use of array copies in our shader.
|
||||
if (type.array.size() > 1)
|
||||
{
|
||||
if (type.array.size() > kArrayCopyMultidimMax)
|
||||
SPIRV_CROSS_THROW("Cannot support this many dimensions for arrays of arrays.");
|
||||
auto func = static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + type.array.size());
|
||||
add_spv_func_and_recompile(func);
|
||||
}
|
||||
else
|
||||
add_spv_func_and_recompile(SPVFuncImplArrayCopy);
|
||||
add_spv_func_and_recompile(type.array.size() > 1 ? SPVFuncImplArrayCopyMultidim : SPVFuncImplArrayCopy);
|
||||
|
||||
const char *tag = nullptr;
|
||||
if (lhs_is_thread_storage && is_constant)
|
||||
@ -10046,13 +10016,13 @@ bool CompilerMSL::emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rh
|
||||
|
||||
// Pass internal array of spvUnsafeArray<> into wrapper functions
|
||||
if (lhs_is_array_template && rhs_is_array_template && !msl_options.force_native_arrays)
|
||||
statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ".elements, ", to_expression(rhs_id), ".elements);");
|
||||
statement("spvArrayCopy", tag, "(", lhs, ".elements, ", to_expression(rhs_id), ".elements);");
|
||||
if (lhs_is_array_template && !msl_options.force_native_arrays)
|
||||
statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ".elements, ", to_expression(rhs_id), ");");
|
||||
statement("spvArrayCopy", tag, "(", lhs, ".elements, ", to_expression(rhs_id), ");");
|
||||
else if (rhs_is_array_template && !msl_options.force_native_arrays)
|
||||
statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ", ", to_expression(rhs_id), ".elements);");
|
||||
statement("spvArrayCopy", tag, "(", lhs, ", ", to_expression(rhs_id), ".elements);");
|
||||
else
|
||||
statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ", ", to_expression(rhs_id), ");");
|
||||
statement("spvArrayCopy", tag, "(", lhs, ", ", to_expression(rhs_id), ");");
|
||||
}
|
||||
|
||||
return true;
|
||||
|
||||
@ -287,9 +287,6 @@ static const uint32_t kArgumentBufferBinding = ~(3u);
|
||||
|
||||
static const uint32_t kMaxArgumentBuffers = 8;
|
||||
|
||||
// The arbitrary maximum for the nesting of array of array copies.
|
||||
static const uint32_t kArrayCopyMultidimMax = 6;
|
||||
|
||||
// Decompiles SPIR-V to Metal Shading Language
|
||||
class CompilerMSL : public CompilerGLSL
|
||||
{
|
||||
@ -752,15 +749,8 @@ protected:
|
||||
SPVFuncImplFindSMsb,
|
||||
SPVFuncImplFindUMsb,
|
||||
SPVFuncImplSSign,
|
||||
SPVFuncImplArrayCopyMultidimBase,
|
||||
// Unfortunately, we cannot use recursive templates in the MSL compiler properly,
|
||||
// so stamp out variants up to some arbitrary maximum.
|
||||
SPVFuncImplArrayCopy = SPVFuncImplArrayCopyMultidimBase + 1,
|
||||
SPVFuncImplArrayOfArrayCopy2Dim = SPVFuncImplArrayCopyMultidimBase + 2,
|
||||
SPVFuncImplArrayOfArrayCopy3Dim = SPVFuncImplArrayCopyMultidimBase + 3,
|
||||
SPVFuncImplArrayOfArrayCopy4Dim = SPVFuncImplArrayCopyMultidimBase + 4,
|
||||
SPVFuncImplArrayOfArrayCopy5Dim = SPVFuncImplArrayCopyMultidimBase + 5,
|
||||
SPVFuncImplArrayOfArrayCopy6Dim = SPVFuncImplArrayCopyMultidimBase + 6,
|
||||
SPVFuncImplArrayCopy,
|
||||
SPVFuncImplArrayCopyMultidim,
|
||||
SPVFuncImplTexelBufferCoords,
|
||||
SPVFuncImplImage2DAtomicCoords, // Emulate texture2D atomic operations
|
||||
SPVFuncImplGradientCube,
|
||||
|
||||
Loading…
Reference in New Issue
Block a user