• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1#include <metal_stdlib>
2
3using namespace metal;
4
5template<typename T, int N, int M>
6inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
7  return lhs * vec<T, N>(rhs);
8}
9
10template<typename T, int N, int M>
11inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
12  return vec<T, M>(lhs) * rhs;
13}
14
15struct LightData {
16  /* 0x0000 */ float4 position;
17  /* 0x0010 */ packed_float3 color;
18  /* 0x001c */ float radius;
19};
20struct LightsBuffer {
21  /* 0x0000 */ LightData lights[1];
22};
23struct tint_array_wrapper {
24  /* 0x0000 */ uint arr[64];
25};
26struct TileLightIdData {
27  /* 0x0000 */ atomic_uint count;
28  /* 0x0004 */ tint_array_wrapper lightId;
29};
30struct tint_array_wrapper_1 {
31  /* 0x0000 */ TileLightIdData arr[4];
32};
33struct Tiles {
34  /* 0x0000 */ tint_array_wrapper_1 data;
35};
36struct Config {
37  /* 0x0000 */ uint numLights;
38  /* 0x0004 */ uint numTiles;
39  /* 0x0008 */ uint tileCountX;
40  /* 0x000c */ uint tileCountY;
41  /* 0x0010 */ uint numTileLightSlot;
42  /* 0x0014 */ uint tileSize;
43};
44struct Uniforms {
45  /* 0x0000 */ float4 min;
46  /* 0x0010 */ float4 max;
47  /* 0x0020 */ float4x4 viewMatrix;
48  /* 0x0060 */ float4x4 projectionMatrix;
49  /* 0x00a0 */ float4 fullScreenSize;
50};
51struct tint_array_wrapper_2 {
52  float4 arr[6];
53};
54
55void tint_symbol_inner(uint3 GlobalInvocationID, const constant Config* const tint_symbol_1, device LightsBuffer* const tint_symbol_2, const constant Uniforms* const tint_symbol_3, device Tiles* const tint_symbol_4) {
56  uint index = GlobalInvocationID[0];
57  if ((index >= (*(tint_symbol_1)).numLights)) {
58    return;
59  }
60  (*(tint_symbol_2)).lights[index].position[1] = (((*(tint_symbol_2)).lights[index].position[1] - 0.100000001f) + (0.001f * (float(index) - (64.0f * floor((float(index) / 64.0f))))));
61  if (((*(tint_symbol_2)).lights[index].position[1] < (*(tint_symbol_3)).min[1])) {
62    (*(tint_symbol_2)).lights[index].position[1] = (*(tint_symbol_3)).max[1];
63  }
64  float4x4 M = (*(tint_symbol_3)).projectionMatrix;
65  float viewNear = (-(M[3][2]) / (-1.0f + M[2][2]));
66  float viewFar = (-(M[3][2]) / (1.0f + M[2][2]));
67  float4 lightPos = (*(tint_symbol_2)).lights[index].position;
68  lightPos = ((*(tint_symbol_3)).viewMatrix * lightPos);
69  lightPos = (lightPos / lightPos[3]);
70  float lightRadius = (*(tint_symbol_2)).lights[index].radius;
71  float4 boxMin = (lightPos - float4(float3(lightRadius), 0.0f));
72  float4 boxMax = (lightPos + float4(float3(lightRadius), 0.0f));
73  tint_array_wrapper_2 frustumPlanes = {};
74  frustumPlanes.arr[4] = float4(0.0f, 0.0f, -1.0f, viewNear);
75  frustumPlanes.arr[5] = float4(0.0f, 0.0f, 1.0f, -(viewFar));
76  int const TILE_SIZE = 16;
77  int const TILE_COUNT_X = 2;
78  int const TILE_COUNT_Y = 2;
79  for(int y_1 = 0; (y_1 < TILE_COUNT_Y); y_1 = as_type<int>((as_type<uint>(y_1) + as_type<uint>(1)))) {
80    for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
81      int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_SIZE))));
82      float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
83      float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(TILE_SIZE)))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
84      float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
85      float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
86      frustumPlanes.arr[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0]) / viewNear), 0.0f);
87      frustumPlanes.arr[1] = float4(-1.0f, 0.0f, (viewCeilCoord[0] / viewNear), 0.0f);
88      frustumPlanes.arr[2] = float4(0.0f, 1.0f, (-(viewFloorCoord[1]) / viewNear), 0.0f);
89      frustumPlanes.arr[3] = float4(0.0f, -1.0f, (viewCeilCoord[1] / viewNear), 0.0f);
90      float dp = 0.0f;
91      for(uint i = 0u; (i < 6u); i = (i + 1u)) {
92        float4 p = 0.0f;
93        if ((frustumPlanes.arr[i][0] > 0.0f)) {
94          p[0] = boxMax[0];
95        } else {
96          p[0] = boxMin[0];
97        }
98        if ((frustumPlanes.arr[i][1] > 0.0f)) {
99          p[1] = boxMax[1];
100        } else {
101          p[1] = boxMin[1];
102        }
103        if ((frustumPlanes.arr[i][2] > 0.0f)) {
104          p[2] = boxMax[2];
105        } else {
106          p[2] = boxMin[2];
107        }
108        p[3] = 1.0f;
109        dp = (dp + fmin(0.0f, dot(p, frustumPlanes.arr[i])));
110      }
111      if ((dp >= 0.0f)) {
112        uint tileId = uint(as_type<int>((as_type<uint>(x_1) + as_type<uint>(as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_COUNT_X)))))));
113        if (((tileId < 0u) || (tileId >= (*(tint_symbol_1)).numTiles))) {
114          continue;
115        }
116        uint offset = atomic_fetch_add_explicit(&((*(tint_symbol_4)).data.arr[tileId].count), 1u, memory_order_relaxed);
117        if ((offset >= (*(tint_symbol_1)).numTileLightSlot)) {
118          continue;
119        }
120        (*(tint_symbol_4)).data.arr[tileId].lightId.arr[offset] = GlobalInvocationID[0];
121      }
122    }
123  }
124}
125
126kernel void tint_symbol(const constant Config* tint_symbol_5 [[buffer(0)]], device LightsBuffer* tint_symbol_6 [[buffer(2)]], const constant Uniforms* tint_symbol_7 [[buffer(1)]], device Tiles* tint_symbol_8 [[buffer(3)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
127  tint_symbol_inner(GlobalInvocationID, tint_symbol_5, tint_symbol_6, tint_symbol_7, tint_symbol_8);
128  return;
129}
130
131