twinaphex 420b7913d4 Squashed 'deps/SPIRV-Cross/' changes from 8aa6731925..2820ab0b51
2820ab0b51 Merge pull request #1076 from KhronosGroup/bitcast-pre-330-glsl
63bcbd511e GLSL: Need extension to use bitcast on GLSL < 330.
9f3bebe3d0 Merge pull request #1075 from lifpan/master
b11c20fc1d Remove unreasonable assertion for OpTypeImage Sampled parameter.
1a592b7c0f Merge pull request #1067 from cdavis5e/msl-scalar-block-layout
28454facbb MSL: Handle packed matrices.
ea5c0ed82f MSL: Fix alignment of packed types.
44f688bf0b Merge pull request #1070 from KhronosGroup/fix-1066
25c74b324e Forget loop variable enables after emitting block chain.
6b010e0cbc Merge pull request #1069 from KhronosGroup/fix-1053
f6f849397e MSL: Re-roll array expressions in initializers.
e5fa7edfd6 MSL: Support scalar block layout.

git-subtree-dir: deps/SPIRV-Cross
git-subtree-split: 2820ab0b51bf5e4187435d904b34e762b988f48b
2019-07-11 20:37:45 +02:00

40 lines
1.0 KiB
Plaintext

#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO0
{
float4 as[1];
};
struct SSBO1
{
float4 bs[1];
};
// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
template<typename T, uint N>
void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
{
for (uint i = 0; i < N; dst[i] = src[i], i++);
}
template<typename T, uint N>
void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
{
for (uint i = 0; i < N; dst[i] = src[i], i++);
}
kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO1& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
{
float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.bs[gl_GlobalInvocationID.x] };
float4 values[2];
spvArrayCopyFromStack1(values, _37);
_16.as[0] = values[gl_LocalInvocationIndex];
_32.bs[1] = float4(40.0);
}