1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO_A
{
float data[1];
};
struct UBO_C
{
float4 data[1024];
};
struct Registers
{
float reg;
};
struct SSBO_B
{
uint2 data[1];
};
struct UBO_D
{
uint4 data[1024];
};
struct SSBO_BRO
{
uint2 data[1];
};
struct SSBO_As
{
float data[1];
};
struct UBO_Cs
{
float4 data[1024];
};
struct SSBO_Bs
{
uint2 data[1024];
};
struct UBO_Ds
{
uint4 data[1024];
};
struct SSBO_BsRO
{
uint2 data[1024];
};
struct SSBO_E
{
float data[1];
};
struct UBO_G
{
float4 data[1024];
};
struct SSBO_F
{
uint2 data[1];
};
struct UBO_H
{
uint4 data[1024];
};
struct SSBO_I
{
uint2 data[1];
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
struct spvDescriptorSetBuffer0
{
device SSBO_A* ssbo_a [[id(0)]];
constant UBO_C* ubo_c [[id(1)]];
device SSBO_As* ssbo_as [[id(2)]][4];
constant UBO_Cs* ubo_cs [[id(6)]][4];
};
static inline __attribute__((always_inline))
void func0(device SSBO_A& ssbo_a, thread uint3& gl_GlobalInvocationID, constant UBO_C& ubo_c, thread uint3& gl_WorkGroupID, constant Registers& v_42, device SSBO_B& ssbo_b, constant UBO_D& ubo_d, const device SSBO_BRO& ssbo_b_readonly)
{
ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x].x + v_42.reg;
ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
}
static inline __attribute__((always_inline))
void func1(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_As* constant (&ssbo_as)[4], constant UBO_Cs* constant (&ubo_cs)[4])
{
ssbo_as[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_cs[gl_WorkGroupID.x]->data[0].x;
}
static inline __attribute__((always_inline))
void func2(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_Bs* constant (&ssbo_bs)[4], constant UBO_Ds* constant (&ubo_ds)[4], const device SSBO_BsRO* constant (&ssbo_bs_readonly)[4])
{
ssbo_bs[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_ds[gl_WorkGroupID.x]->data[0].xy + ssbo_bs_readonly[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x];
}
static inline __attribute__((always_inline))
void func3(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_E& ssbo_e, constant UBO_G& ubo_g, device SSBO_F& ssbo_f, constant UBO_H& ubo_h, const device SSBO_I& ssbo_i)
{
ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x].x;
ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y].xy + ssbo_i.data[gl_GlobalInvocationID.x];
}
kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant Registers& v_42 [[buffer(1)]], device void* spvBufferAliasSet2Binding0 [[buffer(2)]], constant void* spvBufferAliasSet2Binding1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
{
device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding0;
constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding1;
device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding0;
constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding1;
const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding0;
device auto& ssbo_b = (device SSBO_B&)(*spvDescriptorSet0.ssbo_a);
constant auto& ubo_d = (constant UBO_D&)(*spvDescriptorSet0.ubo_c);
const device auto& ssbo_b_readonly = (const device SSBO_BRO&)(*spvDescriptorSet0.ssbo_a);
constant auto& ssbo_bs = (device SSBO_Bs* constant (&)[4])spvDescriptorSet0.ssbo_as;
constant auto& ubo_ds = (constant UBO_Ds* constant (&)[4])spvDescriptorSet0.ubo_cs;
constant auto& ssbo_bs_readonly = (const device SSBO_BsRO* constant (&)[4])spvDescriptorSet0.ssbo_as;
func0((*spvDescriptorSet0.ssbo_a), gl_GlobalInvocationID, (*spvDescriptorSet0.ubo_c), gl_WorkGroupID, v_42, ssbo_b, ubo_d, ssbo_b_readonly);
func1(gl_GlobalInvocationID, gl_WorkGroupID, spvDescriptorSet0.ssbo_as, spvDescriptorSet0.ubo_cs);
func2(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_bs, ubo_ds, ssbo_bs_readonly);
func3(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_e, ubo_g, ssbo_f, ubo_h, ssbo_i);
}
|
