dawn-cmake/test/tint/bug/tint/1121.wgsl.expected.msl

129 lines
5.6 KiB
Plaintext

#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct LightData {
/* 0x0000 */ float4 position;
/* 0x0010 */ packed_float3 color;
/* 0x001c */ float radius;
};
struct LightsBuffer {
/* 0x0000 */ tint_array<LightData, 1> lights;
};
struct TileLightIdData {
/* 0x0000 */ atomic_uint count;
/* 0x0004 */ tint_array<uint, 64> lightId;
};
struct Tiles {
/* 0x0000 */ tint_array<TileLightIdData, 4> data;
};
struct Config {
/* 0x0000 */ uint numLights;
/* 0x0004 */ uint numTiles;
/* 0x0008 */ uint tileCountX;
/* 0x000c */ uint tileCountY;
/* 0x0010 */ uint numTileLightSlot;
/* 0x0014 */ uint tileSize;
};
struct Uniforms {
/* 0x0000 */ float4 min;
/* 0x0010 */ float4 max;
/* 0x0020 */ float4x4 viewMatrix;
/* 0x0060 */ float4x4 projectionMatrix;
/* 0x00a0 */ float4 fullScreenSize;
};
void 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) {
uint index = GlobalInvocationID[0];
if ((index >= (*(tint_symbol_1)).numLights)) {
return;
}
(*(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))))));
if (((*(tint_symbol_2)).lights[index].position[1] < (*(tint_symbol_3)).min[1])) {
(*(tint_symbol_2)).lights[index].position[1] = (*(tint_symbol_3)).max[1];
}
float4x4 M = (*(tint_symbol_3)).projectionMatrix;
float viewNear = (-(M[3][2]) / (-1.0f + M[2][2]));
float viewFar = (-(M[3][2]) / (1.0f + M[2][2]));
float4 lightPos = (*(tint_symbol_2)).lights[index].position;
lightPos = ((*(tint_symbol_3)).viewMatrix * lightPos);
lightPos = (lightPos / lightPos[3]);
float lightRadius = (*(tint_symbol_2)).lights[index].radius;
float4 boxMin = (lightPos - float4(float3(lightRadius), 0.0f));
float4 boxMax = (lightPos + float4(float3(lightRadius), 0.0f));
tint_array<float4, 6> frustumPlanes = {};
frustumPlanes[4] = float4(0.0f, 0.0f, -1.0f, viewNear);
frustumPlanes[5] = float4(0.0f, 0.0f, 1.0f, -(viewFar));
int const TILE_SIZE = 16;
int const TILE_COUNT_X = 2;
int const TILE_COUNT_Y = 2;
for(int y = 0; (y < TILE_COUNT_Y); y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)))) {
for(int x = 0; (x < TILE_COUNT_X); x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)))) {
int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y) * as_type<uint>(TILE_SIZE))));
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
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));
float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0]) / viewNear), 0.0f);
frustumPlanes[1] = float4(-1.0f, 0.0f, (viewCeilCoord[0] / viewNear), 0.0f);
frustumPlanes[2] = float4(0.0f, 1.0f, (-(viewFloorCoord[1]) / viewNear), 0.0f);
frustumPlanes[3] = float4(0.0f, -1.0f, (viewCeilCoord[1] / viewNear), 0.0f);
float dp = 0.0f;
for(uint i = 0u; (i < 6u); i = (i + 1u)) {
float4 p = 0.0f;
if ((frustumPlanes[i][0] > 0.0f)) {
p[0] = boxMax[0];
} else {
p[0] = boxMin[0];
}
if ((frustumPlanes[i][1] > 0.0f)) {
p[1] = boxMax[1];
} else {
p[1] = boxMin[1];
}
if ((frustumPlanes[i][2] > 0.0f)) {
p[2] = boxMax[2];
} else {
p[2] = boxMin[2];
}
p[3] = 1.0f;
dp = (dp + fmin(0.0f, dot(p, frustumPlanes[i])));
}
if ((dp >= 0.0f)) {
uint tileId = uint(as_type<int>((as_type<uint>(x) + as_type<uint>(as_type<int>((as_type<uint>(y) * as_type<uint>(TILE_COUNT_X)))))));
if (((tileId < 0u) || (tileId >= (*(tint_symbol_1)).numTiles))) {
continue;
}
uint offset = atomic_fetch_add_explicit(&((*(tint_symbol_4)).data[tileId].count), 1u, memory_order_relaxed);
if ((offset >= (*(tint_symbol_1)).numTileLightSlot)) {
continue;
}
(*(tint_symbol_4)).data[tileId].lightId[offset] = GlobalInvocationID[0];
}
}
}
}
kernel 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]]) {
tint_symbol_inner(GlobalInvocationID, tint_symbol_5, tint_symbol_6, tint_symbol_7, tint_symbol_8);
return;
}