95 lines
3.0 KiB
Metal
95 lines
3.0 KiB
Metal
|
// Automatically generated from files in pathfinder/shaders/. Do not edit!
|
||
|
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||
|
|
|
||
|
|
#include <metal_stdlib>
|
||
|
|
#include <simd/simd.h>
|
||
|
|
|
||
|
|
using namespace metal;
|
||
|
|
|
||
|
|
struct bFirstTileMap
|
||
|
|
{
|
||
|
|
int iFirstTileMap[1];
|
||
|
|
};
|
||
|
|
|
||
|
|
struct bTiles
|
||
|
|
{
|
||
|
|
uint iTiles[1];
|
||
|
|
};
|
||
|
|
|
||
|
|
struct bZBuffer
|
||
|
|
{
|
||
|
|
int iZBuffer[1];
|
||
|
|
};
|
||
|
|
|
||
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
|
||
|
|
|
||
|
|
static inline __attribute__((always_inline))
|
||
|
|
int getFirst(thread const uint& globalTileIndex, device bFirstTileMap& v_26)
|
||
|
|
{
|
||
|
|
return v_26.iFirstTileMap[globalTileIndex];
|
||
|
|
}
|
||
|
|
|
||
|
|
static inline __attribute__((always_inline))
|
||
|
|
int getNextTile(thread const int& tileIndex, device bTiles& v_37)
|
||
|
|
{
|
||
|
|
return int(v_37.iTiles[(tileIndex * 4) + 0]);
|
||
|
|
}
|
||
|
|
|
||
|
|
static inline __attribute__((always_inline))
|
||
|
|
void setNextTile(thread const int& tileIndex, thread const int& newNextTileIndex, device bTiles& v_37)
|
||
|
|
{
|
||
|
|
v_37.iTiles[(tileIndex * 4) + 0] = uint(newNextTileIndex);
|
||
|
|
}
|
||
|
|
|
||
|
|
kernel void main0(constant int& uTileCount [[buffer(2)]], device bFirstTileMap& v_26 [[buffer(0)]], device bTiles& v_37 [[buffer(1)]], const device bZBuffer& _76 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||
|
|
{
|
||
|
|
uint globalTileIndex = gl_GlobalInvocationID.x;
|
||
|
|
if (globalTileIndex >= uint(uTileCount))
|
||
|
|
{
|
||
|
|
return;
|
||
|
|
}
|
||
|
|
int zValue = _76.iZBuffer[globalTileIndex];
|
||
|
|
uint param = globalTileIndex;
|
||
|
|
int unsortedFirstTileIndex = getFirst(param, v_26);
|
||
|
|
int sortedFirstTileIndex = -1;
|
||
|
|
while (unsortedFirstTileIndex >= 0)
|
||
|
|
{
|
||
|
|
int currentTileIndex = unsortedFirstTileIndex;
|
||
|
|
int param_1 = currentTileIndex;
|
||
|
|
unsortedFirstTileIndex = getNextTile(param_1, v_37);
|
||
|
|
if (currentTileIndex >= zValue)
|
||
|
|
{
|
||
|
|
int prevTrialTileIndex = -1;
|
||
|
|
int trialTileIndex = sortedFirstTileIndex;
|
||
|
|
while (true)
|
||
|
|
{
|
||
|
|
if ((trialTileIndex < 0) || (currentTileIndex < trialTileIndex))
|
||
|
|
{
|
||
|
|
if (prevTrialTileIndex < 0)
|
||
|
|
{
|
||
|
|
int param_2 = currentTileIndex;
|
||
|
|
int param_3 = sortedFirstTileIndex;
|
||
|
|
setNextTile(param_2, param_3, v_37);
|
||
|
|
sortedFirstTileIndex = currentTileIndex;
|
||
|
|
}
|
||
|
|
else
|
||
|
|
{
|
||
|
|
int param_4 = currentTileIndex;
|
||
|
|
int param_5 = trialTileIndex;
|
||
|
|
setNextTile(param_4, param_5, v_37);
|
||
|
|
int param_6 = prevTrialTileIndex;
|
||
|
|
int param_7 = currentTileIndex;
|
||
|
|
setNextTile(param_6, param_7, v_37);
|
||
|
|
}
|
||
|
|
break;
|
||
|
|
}
|
||
|
|
prevTrialTileIndex = trialTileIndex;
|
||
|
|
int param_8 = trialTileIndex;
|
||
|
|
trialTileIndex = getNextTile(param_8, v_37);
|
||
|
|
}
|
||
|
|
}
|
||
|
|
}
|
||
|
|
v_26.iFirstTileMap[globalTileIndex] = sortedFirstTileIndex;
|
||
|
|
}
|
||
|
|
|