Skip to content

Commit 8724e3f

Browse files
HansKristian-WorkHansKristian-Work
authored
Merge pull request #2619 from KhronosGroup/fix-2614
MSL: Fix some scenarios around writing array type to clip/cull.
2 parents e7f242f + 7a3b67f commit 8724e3f

7 files changed

Lines changed: 668 additions & 4 deletions

File tree

‎reference/opt/shaders-msl/vert/clip-copy.for-tess.vert‎

Lines changed: 76 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,76 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
#pragma clang diagnostic ignored "-Wmissing-braces"
3+
4+
#include <metal_stdlib>
5+
#include <simd/simd.h>
6+
7+
using namespace metal;
8+
9+
template<typename T, size_t Num>
10+
struct spvUnsafeArray
11+
{
12+
T elements[Num ? Num : 1];
13+
14+
thread T& operator [] (size_t pos) thread
15+
{
16+
return elements[pos];
17+
}
18+
constexpr const thread T& operator [] (size_t pos) const thread
19+
{
20+
return elements[pos];
21+
}
22+
23+
device T& operator [] (size_t pos) device
24+
{
25+
return elements[pos];
26+
}
27+
constexpr const device T& operator [] (size_t pos) const device
28+
{
29+
return elements[pos];
30+
}
31+
32+
constexpr const constant T& operator [] (size_t pos) const constant
33+
{
34+
return elements[pos];
35+
}
36+
37+
threadgroup T& operator [] (size_t pos) threadgroup
38+
{
39+
return elements[pos];
40+
}
41+
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
42+
{
43+
return elements[pos];
44+
}
45+
};
46+
47+
struct Block
48+
{
49+
spvUnsafeArray<float, 4> block0;
50+
};
51+
52+
constant spvUnsafeArray<float, 4> _30 = spvUnsafeArray<float, 4>({ 1.0, 2.0, -1.0, -2.0 });
53+
54+
struct main0_out
55+
{
56+
spvUnsafeArray<float, 4> F_array;
57+
spvUnsafeArray<float, 4> m_43_block0;
58+
float4 gl_Position;
59+
spvUnsafeArray<float, 4> gl_ClipDistance;
60+
};
61+
62+
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
63+
{
64+
Block _43 = {};
65+
device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
66+
if (any(gl_GlobalInvocationID >= spvStageInputSize))
67+
return;
68+
out.gl_Position = float4(1.0, 2.0, 3.0, 4.0);
69+
out.gl_ClipDistance = _30;
70+
spvUnsafeArray<float, 4> _51 = out.gl_ClipDistance;
71+
out.gl_ClipDistance = _51;
72+
out.F_array = _51;
73+
_43.block0 = _30;
74+
out.m_43_block0 = _43.block0;
75+
}
76+

‎reference/opt/shaders-msl/vert/clip-copy.vert‎

Lines changed: 209 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,209 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
#pragma clang diagnostic ignored "-Wmissing-braces"
3+
4+
#include <metal_stdlib>
5+
#include <simd/simd.h>
6+
7+
using namespace metal;
8+
9+
template<typename T, size_t Num>
10+
struct spvUnsafeArray
11+
{
12+
T elements[Num ? Num : 1];
13+
14+
thread T& operator [] (size_t pos) thread
15+
{
16+
return elements[pos];
17+
}
18+
constexpr const thread T& operator [] (size_t pos) const thread
19+
{
20+
return elements[pos];
21+
}
22+
23+
device T& operator [] (size_t pos) device
24+
{
25+
return elements[pos];
26+
}
27+
constexpr const device T& operator [] (size_t pos) const device
28+
{
29+
return elements[pos];
30+
}
31+
32+
constexpr const constant T& operator [] (size_t pos) const constant
33+
{
34+
return elements[pos];
35+
}
36+
37+
threadgroup T& operator [] (size_t pos) threadgroup
38+
{
39+
return elements[pos];
40+
}
41+
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
42+
{
43+
return elements[pos];
44+
}
45+
};
46+
47+
template<typename T, uint N>
48+
inline void spvArrayCopyFromConstantToStack(thread T (&dst)[N], constant T (&src)[N])
49+
{
50+
for (uint i = 0; i < N; i++)
51+
{
52+
dst[i] = src[i];
53+
}
54+
}
55+
56+
template<typename T, uint N>
57+
inline void spvArrayCopyFromConstantToThreadGroup(threadgroup T (&dst)[N], constant T (&src)[N])
58+
{
59+
for (uint i = 0; i < N; i++)
60+
{
61+
dst[i] = src[i];
62+
}
63+
}
64+
65+
template<typename T, uint N>
66+
inline void spvArrayCopyFromStackToStack(thread T (&dst)[N], thread const T (&src)[N])
67+
{
68+
for (uint i = 0; i < N; i++)
69+
{
70+
dst[i] = src[i];
71+
}
72+
}
73+
74+
template<typename T, uint N>
75+
inline void spvArrayCopyFromStackToThreadGroup(threadgroup T (&dst)[N], thread const T (&src)[N])
76+
{
77+
for (uint i = 0; i < N; i++)
78+
{
79+
dst[i] = src[i];
80+
}
81+
}
82+
83+
template<typename T, uint N>
84+
inline void spvArrayCopyFromThreadGroupToStack(thread T (&dst)[N], threadgroup const T (&src)[N])
85+
{
86+
for (uint i = 0; i < N; i++)
87+
{
88+
dst[i] = src[i];
89+
}
90+
}
91+
92+
template<typename T, uint N>
93+
inline void spvArrayCopyFromThreadGroupToThreadGroup(threadgroup T (&dst)[N], threadgroup const T (&src)[N])
94+
{
95+
for (uint i = 0; i < N; i++)
96+
{
97+
dst[i] = src[i];
98+
}
99+
}
100+
101+
template<typename T, uint N>
102+
inline void spvArrayCopyFromDeviceToDevice(device T (&dst)[N], device const T (&src)[N])
103+
{
104+
for (uint i = 0; i < N; i++)
105+
{
106+
dst[i] = src[i];
107+
}
108+
}
109+
110+
template<typename T, uint N>
111+
inline void spvArrayCopyFromConstantToDevice(device T (&dst)[N], constant T (&src)[N])
112+
{
113+
for (uint i = 0; i < N; i++)
114+
{
115+
dst[i] = src[i];
116+
}
117+
}
118+
119+
template<typename T, uint N>
120+
inline void spvArrayCopyFromStackToDevice(device T (&dst)[N], thread const T (&src)[N])
121+
{
122+
for (uint i = 0; i < N; i++)
123+
{
124+
dst[i] = src[i];
125+
}
126+
}
127+
128+
template<typename T, uint N>
129+
inline void spvArrayCopyFromThreadGroupToDevice(device T (&dst)[N], threadgroup const T (&src)[N])
130+
{
131+
for (uint i = 0; i < N; i++)
132+
{
133+
dst[i] = src[i];
134+
}
135+
}
136+
137+
template<typename T, uint N>
138+
inline void spvArrayCopyFromDeviceToStack(thread T (&dst)[N], device const T (&src)[N])
139+
{
140+
for (uint i = 0; i < N; i++)
141+
{
142+
dst[i] = src[i];
143+
}
144+
}
145+
146+
template<typename T, uint N>
147+
inline void spvArrayCopyFromDeviceToThreadGroup(threadgroup T (&dst)[N], device const T (&src)[N])
148+
{
149+
for (uint i = 0; i < N; i++)
150+
{
151+
dst[i] = src[i];
152+
}
153+
}
154+
155+
struct Block
156+
{
157+
spvUnsafeArray<float, 4> block0;
158+
};
159+
160+
constant spvUnsafeArray<float, 4> _30 = spvUnsafeArray<float, 4>({ 1.0, 2.0, -1.0, -2.0 });
161+
162+
struct main0_out
163+
{
164+
float F_array_0 [[user(locn0)]];
165+
float F_array_1 [[user(locn1)]];
166+
float F_array_2 [[user(locn2)]];
167+
float F_array_3 [[user(locn3)]];
168+
float m_43_block0_0 [[user(locn4)]];
169+
float m_43_block0_1 [[user(locn5)]];
170+
float m_43_block0_2 [[user(locn6)]];
171+
float m_43_block0_3 [[user(locn7)]];
172+
float4 gl_Position [[position]];
173+
float gl_ClipDistance [[clip_distance]] [4];
174+
float gl_ClipDistance_0 [[user(clip0)]];
175+
float gl_ClipDistance_1 [[user(clip1)]];
176+
float gl_ClipDistance_2 [[user(clip2)]];
177+
float gl_ClipDistance_3 [[user(clip3)]];
178+
};
179+
180+
vertex main0_out main0()
181+
{
182+
main0_out out = {};
183+
spvUnsafeArray<float, 4> F_array = {};
184+
Block _43 = {};
185+
out.gl_Position = float4(1.0, 2.0, 3.0, 4.0);
186+
spvArrayCopyFromConstantToStack(out.gl_ClipDistance, _30.elements);
187+
spvUnsafeArray<float, 4> _51;
188+
_51[0] = out.gl_ClipDistance[0];
189+
_51[1] = out.gl_ClipDistance[1];
190+
_51[2] = out.gl_ClipDistance[2];
191+
_51[3] = out.gl_ClipDistance[3];
192+
spvArrayCopyFromStackToStack(out.gl_ClipDistance, _51.elements);
193+
F_array = _51;
194+
_43.block0 = _30;
195+
out.gl_ClipDistance_0 = out.gl_ClipDistance[0];
196+
out.gl_ClipDistance_1 = out.gl_ClipDistance[1];
197+
out.gl_ClipDistance_2 = out.gl_ClipDistance[2];
198+
out.gl_ClipDistance_3 = out.gl_ClipDistance[3];
199+
out.F_array_0 = F_array[0];
200+
out.F_array_1 = F_array[1];
201+
out.F_array_2 = F_array[2];
202+
out.F_array_3 = F_array[3];
203+
out.m_43_block0_0 = _43.block0[0];
204+
out.m_43_block0_1 = _43.block0[1];
205+
out.m_43_block0_2 = _43.block0[2];
206+
out.m_43_block0_3 = _43.block0[3];
207+
return out;
208+
}
209+

‎reference/shaders-msl/vert/clip-copy.for-tess.vert‎

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
1+
#pragma clang diagnostic ignored "-Wmissing-prototypes"
2+
#pragma clang diagnostic ignored "-Wmissing-braces"
3+
4+
#include <metal_stdlib>
5+
#include <simd/simd.h>
6+
7+
using namespace metal;
8+
9+
template<typename T, size_t Num>
10+
struct spvUnsafeArray
11+
{
12+
T elements[Num ? Num : 1];
13+
14+
thread T& operator [] (size_t pos) thread
15+
{
16+
return elements[pos];
17+
}
18+
constexpr const thread T& operator [] (size_t pos) const thread
19+
{
20+
return elements[pos];
21+
}
22+
23+
device T& operator [] (size_t pos) device
24+
{
25+
return elements[pos];
26+
}
27+
constexpr const device T& operator [] (size_t pos) const device
28+
{
29+
return elements[pos];
30+
}
31+
32+
constexpr const constant T& operator [] (size_t pos) const constant
33+
{
34+
return elements[pos];
35+
}
36+
37+
threadgroup T& operator [] (size_t pos) threadgroup
38+
{
39+
return elements[pos];
40+
}
41+
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
42+
{
43+
return elements[pos];
44+
}
45+
};
46+
47+
struct Block
48+
{
49+
spvUnsafeArray<float, 4> block0;
50+
};
51+
52+
constant spvUnsafeArray<float, 4> _30 = spvUnsafeArray<float, 4>({ 1.0, 2.0, -1.0, -2.0 });
53+
54+
struct main0_out
55+
{
56+
spvUnsafeArray<float, 4> F_array;
57+
spvUnsafeArray<float, 4> m_43_block0;
58+
float4 gl_Position;
59+
spvUnsafeArray<float, 4> gl_ClipDistance;
60+
};
61+
62+
static inline __attribute__((always_inline))
63+
void in_func(device float4& gl_Position, device spvUnsafeArray<float, 4>& gl_ClipDistance, device spvUnsafeArray<float, 4>& F_array, thread Block& _43)
64+
{
65+
gl_Position = float4(1.0, 2.0, 3.0, 4.0);
66+
gl_ClipDistance = _30;
67+
spvUnsafeArray<float, 4> non_const_clips = gl_ClipDistance;
68+
gl_ClipDistance = non_const_clips;
69+
F_array = non_const_clips;
70+
_43.block0 = _30;
71+
}
72+
73+
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
74+
{
75+
Block _43 = {};
76+
device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
77+
if (any(gl_GlobalInvocationID >= spvStageInputSize))
78+
return;
79+
in_func(out.gl_Position, out.gl_ClipDistance, out.F_array, _43);
80+
out.m_43_block0 = _43.block0;
81+
}
82+

0 commit comments

Comments
 (0)