1#pragma clang diagnostic ignored "-Wmissing-prototypes" 2 3#include <metal_stdlib> 4#include <simd/simd.h> 5 6using namespace metal; 7 8struct SSBO 9{ 10 float res; 11 int ires; 12 uint ures; 13 float4 f32; 14 int4 s32; 15 uint4 u32; 16 float2x2 m2; 17 float3x3 m3; 18 float4x4 m4; 19}; 20 21struct ResType 22{ 23 float _m0; 24 int _m1; 25}; 26 27constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); 28 29// Implementation of the GLSL radians() function 30template<typename T> 31inline T radians(T d) 32{ 33 return d * T(0.01745329251); 34} 35 36// Implementation of the GLSL degrees() function 37template<typename T> 38inline T degrees(T r) 39{ 40 return r * T(57.2957795131); 41} 42 43// Implementation of the GLSL findLSB() function 44template<typename T> 45inline T spvFindLSB(T x) 46{ 47 return select(ctz(x), T(-1), x == T(0)); 48} 49 50// Implementation of the signed GLSL findMSB() function 51template<typename T> 52inline T spvFindSMSB(T x) 53{ 54 T v = select(x, T(-1) - x, x < T(0)); 55 return select(clz(T(0)) - (clz(v) + T(1)), T(-1), v == T(0)); 56} 57 58// Implementation of the unsigned GLSL findMSB() function 59template<typename T> 60inline T spvFindUMSB(T x) 61{ 62 return select(clz(T(0)) - (clz(x) + T(1)), T(-1), x == T(0)); 63} 64 65// Implementation of the GLSL sign() function for integer types 66template<typename T, typename E = typename enable_if<is_integral<T>::value>::type> 67inline T sign(T x) 68{ 69 return select(select(select(x, T(0), x == T(0)), T(1), x > T(0)), T(-1), x < T(0)); 70} 71 72// Returns the determinant of a 2x2 matrix. 73static inline __attribute__((always_inline)) 74float spvDet2x2(float a1, float a2, float b1, float b2) 75{ 76 return a1 * b2 - b1 * a2; 77} 78 79// Returns the determinant of a 3x3 matrix. 80static inline __attribute__((always_inline)) 81float spvDet3x3(float a1, float a2, float a3, float b1, float b2, float b3, float c1, float c2, float c3) 82{ 83 return a1 * spvDet2x2(b2, b3, c2, c3) - b1 * spvDet2x2(a2, a3, c2, c3) + c1 * spvDet2x2(a2, a3, b2, b3); 84} 85 86// Returns the inverse of a matrix, by using the algorithm of calculating the classical 87// adjoint and dividing by the determinant. The contents of the matrix are changed. 88static inline __attribute__((always_inline)) 89float4x4 spvInverse4x4(float4x4 m) 90{ 91 float4x4 adj; // The adjoint matrix (inverse after dividing by determinant) 92 93 // Create the transpose of the cofactors, as the classical adjoint of the matrix. 94 adj[0][0] = spvDet3x3(m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]); 95 adj[0][1] = -spvDet3x3(m[0][1], m[0][2], m[0][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]); 96 adj[0][2] = spvDet3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[3][1], m[3][2], m[3][3]); 97 adj[0][3] = -spvDet3x3(m[0][1], m[0][2], m[0][3], m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3]); 98 99 adj[1][0] = -spvDet3x3(m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]); 100 adj[1][1] = spvDet3x3(m[0][0], m[0][2], m[0][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]); 101 adj[1][2] = -spvDet3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[3][0], m[3][2], m[3][3]); 102 adj[1][3] = spvDet3x3(m[0][0], m[0][2], m[0][3], m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3]); 103 104 adj[2][0] = spvDet3x3(m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]); 105 adj[2][1] = -spvDet3x3(m[0][0], m[0][1], m[0][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]); 106 adj[2][2] = spvDet3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[3][0], m[3][1], m[3][3]); 107 adj[2][3] = -spvDet3x3(m[0][0], m[0][1], m[0][3], m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3]); 108 109 adj[3][0] = -spvDet3x3(m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]); 110 adj[3][1] = spvDet3x3(m[0][0], m[0][1], m[0][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]); 111 adj[3][2] = -spvDet3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[3][0], m[3][1], m[3][2]); 112 adj[3][3] = spvDet3x3(m[0][0], m[0][1], m[0][2], m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2]); 113 114 // Calculate the determinant as a combination of the cofactors of the first row. 115 float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]) + (adj[0][2] * m[2][0]) + (adj[0][3] * m[3][0]); 116 117 // Divide the classical adjoint matrix by the determinant. 118 // If determinant is zero, matrix is not invertable, so leave it unchanged. 119 return (det != 0.0f) ? (adj * (1.0f / det)) : m; 120} 121 122// Returns the inverse of a matrix, by using the algorithm of calculating the classical 123// adjoint and dividing by the determinant. The contents of the matrix are changed. 124static inline __attribute__((always_inline)) 125float3x3 spvInverse3x3(float3x3 m) 126{ 127 float3x3 adj; // The adjoint matrix (inverse after dividing by determinant) 128 129 // Create the transpose of the cofactors, as the classical adjoint of the matrix. 130 adj[0][0] = spvDet2x2(m[1][1], m[1][2], m[2][1], m[2][2]); 131 adj[0][1] = -spvDet2x2(m[0][1], m[0][2], m[2][1], m[2][2]); 132 adj[0][2] = spvDet2x2(m[0][1], m[0][2], m[1][1], m[1][2]); 133 134 adj[1][0] = -spvDet2x2(m[1][0], m[1][2], m[2][0], m[2][2]); 135 adj[1][1] = spvDet2x2(m[0][0], m[0][2], m[2][0], m[2][2]); 136 adj[1][2] = -spvDet2x2(m[0][0], m[0][2], m[1][0], m[1][2]); 137 138 adj[2][0] = spvDet2x2(m[1][0], m[1][1], m[2][0], m[2][1]); 139 adj[2][1] = -spvDet2x2(m[0][0], m[0][1], m[2][0], m[2][1]); 140 adj[2][2] = spvDet2x2(m[0][0], m[0][1], m[1][0], m[1][1]); 141 142 // Calculate the determinant as a combination of the cofactors of the first row. 143 float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]) + (adj[0][2] * m[2][0]); 144 145 // Divide the classical adjoint matrix by the determinant. 146 // If determinant is zero, matrix is not invertable, so leave it unchanged. 147 return (det != 0.0f) ? (adj * (1.0f / det)) : m; 148} 149 150// Returns the inverse of a matrix, by using the algorithm of calculating the classical 151// adjoint and dividing by the determinant. The contents of the matrix are changed. 152static inline __attribute__((always_inline)) 153float2x2 spvInverse2x2(float2x2 m) 154{ 155 float2x2 adj; // The adjoint matrix (inverse after dividing by determinant) 156 157 // Create the transpose of the cofactors, as the classical adjoint of the matrix. 158 adj[0][0] = m[1][1]; 159 adj[0][1] = -m[0][1]; 160 161 adj[1][0] = -m[1][0]; 162 adj[1][1] = m[0][0]; 163 164 // Calculate the determinant as a combination of the cofactors of the first row. 165 float det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]); 166 167 // Divide the classical adjoint matrix by the determinant. 168 // If determinant is zero, matrix is not invertable, so leave it unchanged. 169 return (det != 0.0f) ? (adj * (1.0f / det)) : m; 170} 171 172template<typename T> 173inline T spvReflect(T i, T n) 174{ 175 return i - T(2) * i * n * n; 176} 177 178template<typename T> 179inline T spvRefract(T i, T n, T eta) 180{ 181 T NoI = n * i; 182 T NoI2 = NoI * NoI; 183 T k = T(1) - eta * eta * (T(1) - NoI2); 184 if (k < T(0)) 185 { 186 return T(0); 187 } 188 else 189 { 190 return eta * i - (eta * NoI + sqrt(k)) * n; 191 } 192} 193 194template<typename T> 195inline T spvFaceForward(T n, T i, T nref) 196{ 197 return i * nref < T(0) ? n : -n; 198} 199 200kernel void main0(device SSBO& _19 [[buffer(0)]]) 201{ 202 _19.res = round(((device float*)&_19.f32)[0u]); 203 _19.res = rint(((device float*)&_19.f32)[0u]); 204 _19.res = trunc(((device float*)&_19.f32)[0u]); 205 _19.res = abs(((device float*)&_19.f32)[0u]); 206 _19.ires = abs(((device int*)&_19.s32)[0u]); 207 _19.res = sign(((device float*)&_19.f32)[0u]); 208 _19.ires = sign(((device int*)&_19.s32)[0u]); 209 _19.res = floor(((device float*)&_19.f32)[0u]); 210 _19.res = ceil(((device float*)&_19.f32)[0u]); 211 _19.res = fract(((device float*)&_19.f32)[0u]); 212 _19.res = radians(((device float*)&_19.f32)[0u]); 213 _19.res = degrees(((device float*)&_19.f32)[0u]); 214 _19.res = sin(((device float*)&_19.f32)[0u]); 215 _19.res = cos(((device float*)&_19.f32)[0u]); 216 _19.res = tan(((device float*)&_19.f32)[0u]); 217 _19.res = asin(((device float*)&_19.f32)[0u]); 218 _19.res = acos(((device float*)&_19.f32)[0u]); 219 _19.res = atan(((device float*)&_19.f32)[0u]); 220 _19.res = sinh(((device float*)&_19.f32)[0u]); 221 _19.res = cosh(((device float*)&_19.f32)[0u]); 222 _19.res = tanh(((device float*)&_19.f32)[0u]); 223 _19.res = asinh(((device float*)&_19.f32)[0u]); 224 _19.res = acosh(((device float*)&_19.f32)[0u]); 225 _19.res = atanh(((device float*)&_19.f32)[0u]); 226 _19.res = atan2(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); 227 _19.res = pow(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); 228 _19.res = exp(((device float*)&_19.f32)[0u]); 229 _19.res = log(((device float*)&_19.f32)[0u]); 230 _19.res = exp2(((device float*)&_19.f32)[0u]); 231 _19.res = log2(((device float*)&_19.f32)[0u]); 232 _19.res = sqrt(((device float*)&_19.f32)[0u]); 233 _19.res = rsqrt(((device float*)&_19.f32)[0u]); 234 _19.res = abs(((device float*)&_19.f32)[0u]); 235 _19.res = abs(((device float*)&_19.f32)[0u] - ((device float*)&_19.f32)[1u]); 236 _19.res = sign(((device float*)&_19.f32)[0u]); 237 _19.res = spvFaceForward(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u], ((device float*)&_19.f32)[2u]); 238 _19.res = spvReflect(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); 239 _19.res = spvRefract(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u], ((device float*)&_19.f32)[2u]); 240 _19.res = length(_19.f32.xy); 241 _19.res = distance(_19.f32.xy, _19.f32.zw); 242 float2 v2 = normalize(_19.f32.xy); 243 v2 = faceforward(_19.f32.xy, _19.f32.yz, _19.f32.zw); 244 v2 = reflect(_19.f32.xy, _19.f32.zw); 245 v2 = refract(_19.f32.xy, _19.f32.yz, ((device float*)&_19.f32)[3u]); 246 float3 v3 = cross(_19.f32.xyz, _19.f32.yzw); 247 _19.res = determinant(_19.m2); 248 _19.res = determinant(_19.m3); 249 _19.res = determinant(_19.m4); 250 _19.m2 = spvInverse2x2(_19.m2); 251 _19.m3 = spvInverse3x3(_19.m3); 252 _19.m4 = spvInverse4x4(_19.m4); 253 float tmp; 254 float _287 = modf(((device float*)&_19.f32)[0u], tmp); 255 _19.res = _287; 256 _19.res = fast::min(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); 257 _19.ures = min(((device uint*)&_19.u32)[0u], ((device uint*)&_19.u32)[1u]); 258 _19.ires = min(((device int*)&_19.s32)[0u], ((device int*)&_19.s32)[1u]); 259 _19.res = fast::max(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); 260 _19.ures = max(((device uint*)&_19.u32)[0u], ((device uint*)&_19.u32)[1u]); 261 _19.ires = max(((device int*)&_19.s32)[0u], ((device int*)&_19.s32)[1u]); 262 _19.res = fast::clamp(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u], ((device float*)&_19.f32)[2u]); 263 _19.ures = clamp(((device uint*)&_19.u32)[0u], ((device uint*)&_19.u32)[1u], ((device uint*)&_19.u32)[2u]); 264 _19.ires = clamp(((device int*)&_19.s32)[0u], ((device int*)&_19.s32)[1u], ((device int*)&_19.s32)[2u]); 265 _19.res = mix(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u], ((device float*)&_19.f32)[2u]); 266 _19.res = step(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); 267 _19.res = smoothstep(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u], ((device float*)&_19.f32)[2u]); 268 _19.res = fma(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u], ((device float*)&_19.f32)[2u]); 269 ResType _387; 270 _387._m0 = frexp(((device float*)&_19.f32)[0u], _387._m1); 271 int itmp = _387._m1; 272 _19.res = _387._m0; 273 _19.res = ldexp(((device float*)&_19.f32)[0u], itmp); 274 _19.ures = pack_float_to_snorm4x8(_19.f32); 275 _19.ures = pack_float_to_unorm4x8(_19.f32); 276 _19.ures = pack_float_to_snorm2x16(_19.f32.xy); 277 _19.ures = pack_float_to_unorm2x16(_19.f32.xy); 278 _19.ures = as_type<uint>(half2(_19.f32.xy)); 279 v2 = unpack_snorm2x16_to_float(((device uint*)&_19.u32)[0u]); 280 v2 = unpack_unorm2x16_to_float(((device uint*)&_19.u32)[0u]); 281 v2 = float2(as_type<half2>(((device uint*)&_19.u32)[0u])); 282 float4 v4 = unpack_snorm4x8_to_float(((device uint*)&_19.u32)[0u]); 283 v4 = unpack_unorm4x8_to_float(((device uint*)&_19.u32)[0u]); 284 _19.s32 = spvFindLSB(_19.s32); 285 _19.s32 = int4(spvFindLSB(_19.u32)); 286 _19.s32 = spvFindSMSB(_19.s32); 287 _19.s32 = int4(spvFindUMSB(_19.u32)); 288} 289 290