From 5c2531da748443bf10f1bbed90081e6c59a1fd6e Mon Sep 17 00:00:00 2001 From: y-jan137 Date: Mon, 27 Apr 2026 18:47:16 +0300 Subject: Optimize span sort dispatch --- Sources/Shaders.metal | 126 ++++++++++++++++++-------------------------------- 1 file changed, 46 insertions(+), 80 deletions(-) (limited to 'Sources/Shaders.metal') diff --git a/Sources/Shaders.metal b/Sources/Shaders.metal index 31608a6..c8989e9 100644 --- a/Sources/Shaders.metal +++ b/Sources/Shaders.metal @@ -1,21 +1,6 @@ #include using namespace metal; -// -------------------------------------------------------------------- -// Span-based pixel sorting (Unity Pixel-Sorting port) -// -// Pipeline: -// 1) createMask: mark pixels whose luminance is within thresholds -// 2) clearSpanBuffer -// 3) identifySpans: for each row, write span length at span start pixel -// 4) rgbToSortValue: compute per-pixel sort value (R/G/B/L/S/H) -// 5) pixelSortSpan: for each span start, sort pixels within the span -// 6) composite: apply sorted pixels only where mask==1 -// -// This matches the semantics of the reference Unity compute shader: -// contiguous masked regions are sorted independently. -// -------------------------------------------------------------------- - enum SortKey : uint { Brightness = 0, Hue = 1, @@ -25,10 +10,9 @@ enum SortKey : uint { Blue = 5, }; -// ---- helpers ------------------------------------------------------- - +// --- helpers --- static inline float luminance(float3 rgb) { - return dot(rgb, float3(0.299, 0.587, 0.114)); + return dot(rgb, float3(0.299, 0.587, 0.114)); // RGB -> greyscale } static inline float hue(float3 c) { @@ -64,8 +48,7 @@ static inline float sort_value(float3 rgb, uint key) { return luminance(rgb); } -// ---- parameters passed from the CPU side --------------------------- - +// --- parameters passed from CPU --- struct Params { uint width; uint height; @@ -78,7 +61,7 @@ struct Params { uint invertMask; // 0/1 }; -// ---- create mask (0/1) --------------------------------------------- +// --- create mask (0/1) --- kernel void createMask( texture2d colorTex [[texture(0)]], @@ -95,7 +78,7 @@ kernel void createMask( maskTex.write(m, gid); } -// ---- clear span buffer --------------------------------------------- +// --- clear span buffer --- kernel void clearSpanBuffer( texture2d spanTex [[texture(0)]], @@ -106,8 +89,7 @@ kernel void clearSpanBuffer( spanTex.write(0u, gid); } -// ---- identify spans (horizontal only) ------------------------------ -// One thread per row: write span length at each span start. +// --- identify spans w mask --- kernel void identifySpans( texture2d maskTex [[texture(0)]], @@ -128,8 +110,6 @@ kernel void identifySpans( pos += 1; if (m == 0 || spanLength >= spanLimit) { - // Write at span start. Mirror Unity behavior: if we hit an unmasked pixel, - // spanLength is current count; if we hit limit while still masked, include current pixel. if (spanLength != 0) { uint outLen = (m == 1u) ? (spanLength + 1u) : spanLength; spanTex.write(outLen, uint2(spanStart, row)); @@ -146,82 +126,68 @@ kernel void identifySpans( } } -// ---- per-pixel sort value ------------------------------------------ - -kernel void rgbToSortValue( - texture2d colorTex [[texture(0)]], - texture2d valTex [[texture(1)]], - constant Params ¶ms [[buffer(0)]], - uint2 gid [[thread_position_in_grid]] -) { - if (gid.x >= params.width || gid.y >= params.height) return; - float3 rgb = saturate(colorTex.read(gid).rgb); - float v = sort_value(rgb, params.sortKey); - valTex.write(half(v), gid); -} - -// ---- sort pixels within a span ------------------------------------- -// One thread per pixel, but only span starts do work. -// Writes sorted pixels into `sortedTex` for the span region. +// --- sort pixels within a span --- constant uint MAX_LOCAL_SPAN = 2048; kernel void pixelSortSpan( texture2d colorTex [[texture(0)]], - texture2d valTex [[texture(1)]], texture2d spanTex [[texture(2)]], texture2d sortedTex [[texture(3)]], constant Params ¶ms [[buffer(0)]], uint2 gid [[thread_position_in_grid]] ) { - if (gid.x >= params.width || gid.y >= params.height) return; + uint row = gid.y; + if (gid.x != 0 || row >= params.height) return; - uint spanLength = spanTex.read(gid).x; - if (spanLength == 0) return; + for (uint x = 0; x < params.width; ++x) { + uint spanLength = spanTex.read(uint2(x, row)).x; + if (spanLength == 0) continue; - spanLength = min(spanLength, params.width - gid.x); - spanLength = min(spanLength, max(1u, params.maxSpanLength)); - spanLength = min(spanLength, MAX_LOCAL_SPAN); + spanLength = min(spanLength, params.width - x); + spanLength = min(spanLength, max(1u, params.maxSpanLength)); + spanLength = min(spanLength, MAX_LOCAL_SPAN); - // Cache sort values for this span. - float cache[MAX_LOCAL_SPAN]; - for (uint k = 0; k < spanLength; ++k) { - cache[k] = float(valTex.read(uint2(gid.x + k, gid.y)).x); - } + float cache[MAX_LOCAL_SPAN]; + for (uint k = 0; k < spanLength; ++k) { + float3 rgb = saturate(colorTex.read(uint2(x + k, row)).rgb); + cache[k] = sort_value(rgb, params.sortKey); + } - float minValue = cache[0]; - float maxValue = cache[0]; - uint minIndex = 0; - uint maxIndex = 0; - - uint steps = (spanLength / 2) + 1; - for (uint i = 0; i < steps; ++i) { - for (uint j = 1; j < spanLength; ++j) { - float v = cache[j]; - // Unity checks `v == saturate(v)` to ignore sentinels; equivalent is 0..1. - if (v >= 0.0f && v <= 1.0f) { - if (v < minValue) { minValue = v; minIndex = j; } - if (maxValue < v) { maxValue = v; maxIndex = j; } + float minValue = cache[0]; + float maxValue = cache[0]; + uint minIndex = 0; + uint maxIndex = 0; + + uint steps = (spanLength / 2) + 1; + for (uint i = 0; i < steps; ++i) { + for (uint j = 1; j < spanLength; ++j) { + float v = cache[j]; + if (v >= 0.0f && v <= 1.0f) { + if (v < minValue) { minValue = v; minIndex = j; } + if (maxValue < v) { maxValue = v; maxIndex = j; } + } } - } - uint dstMin = params.reverseSorting ? i : (spanLength - i - 1); - uint dstMax = params.reverseSorting ? (spanLength - i - 1) : i; + uint dstMin = params.reverseSorting ? i : (spanLength - i - 1); + uint dstMax = params.reverseSorting ? (spanLength - i - 1) : i; - float4 cMin = colorTex.read(uint2(gid.x + minIndex, gid.y)); - float4 cMax = colorTex.read(uint2(gid.x + maxIndex, gid.y)); + float4 cMin = colorTex.read(uint2(x + minIndex, row)); + float4 cMax = colorTex.read(uint2(x + maxIndex, row)); - sortedTex.write(cMin, uint2(gid.x + dstMin, gid.y)); - sortedTex.write(cMax, uint2(gid.x + dstMax, gid.y)); + sortedTex.write(cMin, uint2(x + dstMin, row)); + sortedTex.write(cMax, uint2(x + dstMax, row)); - cache[minIndex] = 2.0f; - cache[maxIndex] = -2.0f; - minValue = 1.0f; - maxValue = -1.0f; + cache[minIndex] = 2.0f; + cache[maxIndex] = -2.0f; + minValue = 1.0f; + maxValue = -1.0f; + } + x += (spanLength - 1); } } -// ---- composite sorted pixels onto original ------------------------- +// --- composite sorted pixels onto original --- kernel void composite( texture2d maskTex [[texture(0)]], -- cgit v1.2.3