285 lines
9.9 KiB
Metal
285 lines
9.9 KiB
Metal
// Automatically generated from files in pathfinder/shaders/. Do not edit!
|
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
#pragma clang diagnostic ignored "-Wunused-variable"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
#include <metal_atomic>
|
|
|
|
using namespace metal;
|
|
|
|
struct bIndirectDrawParams
|
|
{
|
|
uint iIndirectDrawParams[1];
|
|
};
|
|
|
|
struct bTiles
|
|
{
|
|
uint iTiles[1];
|
|
};
|
|
|
|
struct bFills
|
|
{
|
|
uint iFills[1];
|
|
};
|
|
|
|
struct bBackdrops
|
|
{
|
|
uint iBackdrops[1];
|
|
};
|
|
|
|
struct bMicrolines
|
|
{
|
|
uint4 iMicrolines[1];
|
|
};
|
|
|
|
struct bMetadata
|
|
{
|
|
int4 iMetadata[1];
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
|
|
|
|
static inline __attribute__((always_inline))
|
|
float4 unpackMicroline(thread const uint4& packedMicroline, thread uint& outPathIndex)
|
|
{
|
|
outPathIndex = packedMicroline.w;
|
|
int4 signedMicroline = int4(packedMicroline);
|
|
return float4(float((signedMicroline.x << 16) >> 16), float(signedMicroline.x >> 16), float((signedMicroline.y << 16) >> 16), float(signedMicroline.y >> 16)) + (float4(float(signedMicroline.z & 255), float((signedMicroline.z >> 8) & 255), float((signedMicroline.z >> 16) & 255), float((signedMicroline.z >> 24) & 255)) / float4(256.0));
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
uint computeTileIndexNoCheck(thread const int2& tileCoords, thread const int4& pathTileRect, thread const uint& pathTileOffset)
|
|
{
|
|
int2 offsetCoords = tileCoords - pathTileRect.xy;
|
|
return (pathTileOffset + uint(offsetCoords.x)) + uint(offsetCoords.y * (pathTileRect.z - pathTileRect.x));
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
bool4 computeTileOutcodes(thread const int2& tileCoords, thread const int4& pathTileRect)
|
|
{
|
|
return bool4(tileCoords < pathTileRect.xy, tileCoords >= pathTileRect.zw);
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
bool computeTileIndex(thread const int2& tileCoords, thread const int4& pathTileRect, thread const uint& pathTileOffset, thread uint& outTileIndex)
|
|
{
|
|
int2 param = tileCoords;
|
|
int4 param_1 = pathTileRect;
|
|
uint param_2 = pathTileOffset;
|
|
outTileIndex = computeTileIndexNoCheck(param, param_1, param_2);
|
|
int2 param_3 = tileCoords;
|
|
int4 param_4 = pathTileRect;
|
|
return !any(computeTileOutcodes(param_3, param_4));
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void addFill(thread const float4& lineSegment, thread const int2& tileCoords, thread const int4& pathTileRect, thread const uint& pathTileOffset, device bIndirectDrawParams& v_155, device bTiles& v_165, thread int uMaxFillCount, device bFills& v_186)
|
|
{
|
|
int2 param = tileCoords;
|
|
int4 param_1 = pathTileRect;
|
|
uint param_2 = pathTileOffset;
|
|
uint param_3;
|
|
bool _124 = computeTileIndex(param, param_1, param_2, param_3);
|
|
uint tileIndex = param_3;
|
|
if (!_124)
|
|
{
|
|
return;
|
|
}
|
|
uint4 scaledLocalLine = uint4((lineSegment - float4(tileCoords.xyxy * int4(16))) * float4(256.0));
|
|
if (scaledLocalLine.x == scaledLocalLine.z)
|
|
{
|
|
return;
|
|
}
|
|
uint _160 = atomic_fetch_add_explicit((device atomic_uint*)&v_155.iIndirectDrawParams[1], 1u, memory_order_relaxed);
|
|
uint fillIndex = _160;
|
|
uint _174 = atomic_exchange_explicit((device atomic_uint*)&v_165.iTiles[(tileIndex * 4u) + 1u], uint(int(fillIndex)), memory_order_relaxed);
|
|
uint fillLink = _174;
|
|
if (fillIndex < uint(uMaxFillCount))
|
|
{
|
|
v_186.iFills[(fillIndex * 3u) + 0u] = scaledLocalLine.x | (scaledLocalLine.y << uint(16));
|
|
v_186.iFills[(fillIndex * 3u) + 1u] = scaledLocalLine.z | (scaledLocalLine.w << uint(16));
|
|
v_186.iFills[(fillIndex * 3u) + 2u] = fillLink;
|
|
}
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void adjustBackdrop(thread const int& backdropDelta, thread const int2& tileCoords, thread const int4& pathTileRect, thread const uint& pathTileOffset, thread const uint& pathBackdropOffset, device bTiles& v_165, device bBackdrops& v_251)
|
|
{
|
|
int2 param = tileCoords;
|
|
int4 param_1 = pathTileRect;
|
|
bool4 outcodes = computeTileOutcodes(param, param_1);
|
|
if (any(outcodes))
|
|
{
|
|
bool _230 = (!outcodes.x) && outcodes.y;
|
|
bool _236;
|
|
if (_230)
|
|
{
|
|
_236 = !outcodes.z;
|
|
}
|
|
else
|
|
{
|
|
_236 = _230;
|
|
}
|
|
if (_236)
|
|
{
|
|
uint backdropIndex = pathBackdropOffset + uint(tileCoords.x - pathTileRect.x);
|
|
uint _257 = atomic_fetch_add_explicit((device atomic_uint*)&v_251.iBackdrops[backdropIndex * 3u], uint(backdropDelta), memory_order_relaxed);
|
|
}
|
|
}
|
|
else
|
|
{
|
|
int2 param_2 = tileCoords;
|
|
int4 param_3 = pathTileRect;
|
|
uint param_4 = pathTileOffset;
|
|
uint tileIndex = computeTileIndexNoCheck(param_2, param_3, param_4);
|
|
uint _275 = atomic_fetch_add_explicit((device atomic_uint*)&v_165.iTiles[(tileIndex * 4u) + 2u], uint(backdropDelta) << uint(24), memory_order_relaxed);
|
|
}
|
|
}
|
|
|
|
kernel void main0(constant int& uMaxFillCount [[buffer(2)]], constant int& uMicrolineCount [[buffer(5)]], device bIndirectDrawParams& v_155 [[buffer(0)]], device bTiles& v_165 [[buffer(1)]], device bFills& v_186 [[buffer(3)]], device bBackdrops& v_251 [[buffer(4)]], const device bMicrolines& _346 [[buffer(6)]], const device bMetadata& _360 [[buffer(7)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
{
|
|
uint segmentIndex = gl_GlobalInvocationID.x;
|
|
if (segmentIndex >= uint(uMicrolineCount))
|
|
{
|
|
return;
|
|
}
|
|
uint4 param = _346.iMicrolines[segmentIndex];
|
|
uint param_1;
|
|
float4 _354 = unpackMicroline(param, param_1);
|
|
uint pathIndex = param_1;
|
|
float4 lineSegment = _354;
|
|
int4 pathTileRect = _360.iMetadata[(pathIndex * 3u) + 0u];
|
|
uint pathTileOffset = uint(_360.iMetadata[(pathIndex * 3u) + 1u].x);
|
|
uint pathBackdropOffset = uint(_360.iMetadata[(pathIndex * 3u) + 2u].x);
|
|
int2 tileSize = int2(16);
|
|
int4 tileLineSegment = int4(floor(lineSegment / float4(tileSize.xyxy)));
|
|
int2 fromTileCoords = tileLineSegment.xy;
|
|
int2 toTileCoords = tileLineSegment.zw;
|
|
float2 vector = lineSegment.zw - lineSegment.xy;
|
|
float2 vectorIsNegative = float2((vector.x < 0.0) ? (-1.0) : 0.0, (vector.y < 0.0) ? (-1.0) : 0.0);
|
|
int2 tileStep = int2((vector.x < 0.0) ? (-1) : 1, (vector.y < 0.0) ? (-1) : 1);
|
|
float2 firstTileCrossing = float2((fromTileCoords + int2(int(vector.x >= 0.0), int(vector.y >= 0.0))) * tileSize);
|
|
float2 tMax = (firstTileCrossing - lineSegment.xy) / vector;
|
|
float2 tDelta = abs(float2(tileSize) / vector);
|
|
float2 currentPosition = lineSegment.xy;
|
|
int2 tileCoords = fromTileCoords;
|
|
int lastStepDirection = 0;
|
|
uint iteration = 0u;
|
|
int nextStepDirection;
|
|
float _501;
|
|
float4 auxiliarySegment;
|
|
while (iteration < 1024u)
|
|
{
|
|
if (tMax.x < tMax.y)
|
|
{
|
|
nextStepDirection = 1;
|
|
}
|
|
else
|
|
{
|
|
if (tMax.x > tMax.y)
|
|
{
|
|
nextStepDirection = 2;
|
|
}
|
|
else
|
|
{
|
|
if (float(tileStep.x) > 0.0)
|
|
{
|
|
nextStepDirection = 1;
|
|
}
|
|
else
|
|
{
|
|
nextStepDirection = 2;
|
|
}
|
|
}
|
|
}
|
|
if (nextStepDirection == 1)
|
|
{
|
|
_501 = tMax.x;
|
|
}
|
|
else
|
|
{
|
|
_501 = tMax.y;
|
|
}
|
|
float nextT = fast::min(_501, 1.0);
|
|
if (all(tileCoords == toTileCoords))
|
|
{
|
|
nextStepDirection = 0;
|
|
}
|
|
float2 nextPosition = mix(lineSegment.xy, lineSegment.zw, float2(nextT));
|
|
float4 clippedLineSegment = float4(currentPosition, nextPosition);
|
|
float4 param_2 = clippedLineSegment;
|
|
int2 param_3 = tileCoords;
|
|
int4 param_4 = pathTileRect;
|
|
uint param_5 = pathTileOffset;
|
|
addFill(param_2, param_3, param_4, param_5, v_155, v_165, uMaxFillCount, v_186);
|
|
bool haveAuxiliarySegment = false;
|
|
if ((tileStep.y < 0) && (nextStepDirection == 2))
|
|
{
|
|
auxiliarySegment = float4(clippedLineSegment.zw, float2(tileCoords * tileSize));
|
|
haveAuxiliarySegment = true;
|
|
}
|
|
else
|
|
{
|
|
if ((tileStep.y > 0) && (lastStepDirection == 2))
|
|
{
|
|
auxiliarySegment = float4(float2(tileCoords * tileSize), clippedLineSegment.xy);
|
|
haveAuxiliarySegment = true;
|
|
}
|
|
}
|
|
if (haveAuxiliarySegment)
|
|
{
|
|
float4 param_6 = auxiliarySegment;
|
|
int2 param_7 = tileCoords;
|
|
int4 param_8 = pathTileRect;
|
|
uint param_9 = pathTileOffset;
|
|
addFill(param_6, param_7, param_8, param_9, v_155, v_165, uMaxFillCount, v_186);
|
|
}
|
|
if ((tileStep.x < 0) && (lastStepDirection == 1))
|
|
{
|
|
int param_10 = 1;
|
|
int2 param_11 = tileCoords;
|
|
int4 param_12 = pathTileRect;
|
|
uint param_13 = pathTileOffset;
|
|
uint param_14 = pathBackdropOffset;
|
|
adjustBackdrop(param_10, param_11, param_12, param_13, param_14, v_165, v_251);
|
|
}
|
|
else
|
|
{
|
|
if ((tileStep.x > 0) && (nextStepDirection == 1))
|
|
{
|
|
int param_15 = -1;
|
|
int2 param_16 = tileCoords;
|
|
int4 param_17 = pathTileRect;
|
|
uint param_18 = pathTileOffset;
|
|
uint param_19 = pathBackdropOffset;
|
|
adjustBackdrop(param_15, param_16, param_17, param_18, param_19, v_165, v_251);
|
|
}
|
|
}
|
|
if (nextStepDirection == 1)
|
|
{
|
|
tMax.x += tDelta.x;
|
|
tileCoords.x += tileStep.x;
|
|
}
|
|
else
|
|
{
|
|
if (nextStepDirection == 2)
|
|
{
|
|
tMax.y += tDelta.y;
|
|
tileCoords.y += tileStep.y;
|
|
}
|
|
else
|
|
{
|
|
if (nextStepDirection == 0)
|
|
{
|
|
break;
|
|
}
|
|
}
|
|
}
|
|
currentPosition = nextPosition;
|
|
lastStepDirection = nextStepDirection;
|
|
iteration++;
|
|
}
|
|
}
|
|
|