From 0d4261d4e8d312f0cf6886d23727986c690de90c Mon Sep 17 00:00:00 2001 From: floppyhammer Date: Wed, 18 May 2022 13:50:00 +0800 Subject: [PATCH] Handle even-odd fill rule for D3D11 --- resources/shaders/gl4/d3d11/bin.cs.glsl | 24 +++++----- resources/shaders/gl4/d3d11/bound.cs.glsl | 8 ++-- resources/shaders/gl4/d3d11/dice.cs.glsl | 20 ++++---- resources/shaders/gl4/d3d11/fill.cs.glsl | 29 +++++++---- resources/shaders/gl4/d3d11/propagate.cs.glsl | 48 ++++++++++++------- resources/shaders/gl4/d3d11/sort.cs.glsl | 12 ++--- resources/shaders/gl4/d3d11/tile.cs.glsl | 24 ++++++++-- resources/shaders/metal/d3d11/bin.cs.metal | 4 +- resources/shaders/metal/d3d11/bound.cs.metal | 2 +- resources/shaders/metal/d3d11/dice.cs.metal | 2 +- resources/shaders/metal/d3d11/fill.cs.metal | 19 +++++++- .../shaders/metal/d3d11/propagate.cs.metal | 34 +++++++++++-- resources/shaders/metal/d3d11/tile.cs.metal | 18 +++++++ shaders/d3d11/bin.cs.glsl | 24 +++++----- shaders/d3d11/bound.cs.glsl | 8 ++-- shaders/d3d11/dice.cs.glsl | 20 ++++---- shaders/d3d11/fill.cs.glsl | 29 +++++++---- shaders/d3d11/propagate.cs.glsl | 48 ++++++++++++------- shaders/d3d11/sort.cs.glsl | 12 ++--- shaders/d3d11/tile.cs.glsl | 24 ++++++++-- 20 files changed, 280 insertions(+), 129 deletions(-) diff --git a/resources/shaders/gl4/d3d11/bin.cs.glsl b/resources/shaders/gl4/d3d11/bin.cs.glsl index c185cc152..f17da3e87 100644 --- a/resources/shaders/gl4/d3d11/bin.cs.glsl +++ b/resources/shaders/gl4/d3d11/bin.cs.glsl @@ -39,18 +39,18 @@ uniform int uMicrolineCount; uniform int uMaxFillCount; -layout(std430, binding = 0)buffer bMicrolines { - restrict readonly uvec4 iMicrolines[]; +restrict readonly layout(std430, binding = 0)buffer bMicrolines { + uvec4 iMicrolines[]; }; -layout(std430, binding = 1)buffer bMetadata { +restrict readonly layout(std430, binding = 1)buffer bMetadata { - restrict readonly ivec4 iMetadata[]; + ivec4 iMetadata[]; }; @@ -58,27 +58,27 @@ layout(std430, binding = 1)buffer bMetadata { -layout(std430, binding = 2)buffer bIndirectDrawParams { - restrict uint iIndirectDrawParams[]; +restrict layout(std430, binding = 2)buffer bIndirectDrawParams { + uint iIndirectDrawParams[]; }; -layout(std430, binding = 3)buffer bFills { - restrict writeonly uint iFills[]; +restrict writeonly layout(std430, binding = 3)buffer bFills { + uint iFills[]; }; -layout(std430, binding = 4)buffer bTiles { +restrict layout(std430, binding = 4)buffer bTiles { - restrict uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 5)buffer bBackdrops { +restrict layout(std430, binding = 5)buffer bBackdrops { - restrict uint iBackdrops[]; + uint iBackdrops[]; }; uint computeTileIndexNoCheck(ivec2 tileCoords, ivec4 pathTileRect, uint pathTileOffset){ diff --git a/resources/shaders/gl4/d3d11/bound.cs.glsl b/resources/shaders/gl4/d3d11/bound.cs.glsl index 57f074446..3d7644a9b 100644 --- a/resources/shaders/gl4/d3d11/bound.cs.glsl +++ b/resources/shaders/gl4/d3d11/bound.cs.glsl @@ -32,20 +32,20 @@ layout(local_size_x = 64)in; uniform int uPathCount; uniform int uTileCount; -layout(std430, binding = 0)buffer bTilePathInfo { +restrict readonly layout(std430, binding = 0)buffer bTilePathInfo { - restrict readonly uvec4 iTilePathInfo[]; + uvec4 iTilePathInfo[]; }; -layout(std430, binding = 1)buffer bTiles { +restrict layout(std430, binding = 1)buffer bTiles { - restrict uint iTiles[]; + uint iTiles[]; }; void main(){ diff --git a/resources/shaders/gl4/d3d11/dice.cs.glsl b/resources/shaders/gl4/d3d11/dice.cs.glsl index 42549df70..2065107af 100644 --- a/resources/shaders/gl4/d3d11/dice.cs.glsl +++ b/resources/shaders/gl4/d3d11/dice.cs.glsl @@ -42,37 +42,37 @@ uniform int uPathCount; uniform int uLastBatchSegmentIndex; uniform int uMaxMicrolineCount; -layout(std430, binding = 0)buffer bComputeIndirectParams { +restrict layout(std430, binding = 0)buffer bComputeIndirectParams { - restrict uint iComputeIndirectParams[]; + uint iComputeIndirectParams[]; }; -layout(std430, binding = 1)buffer bDiceMetadata { +restrict readonly layout(std430, binding = 1)buffer bDiceMetadata { - restrict readonly uvec4 iDiceMetadata[]; + uvec4 iDiceMetadata[]; }; -layout(std430, binding = 2)buffer bPoints { - restrict readonly vec2 iPoints[]; +restrict readonly layout(std430, binding = 2)buffer bPoints { + vec2 iPoints[]; }; -layout(std430, binding = 3)buffer bInputIndices { - restrict readonly uvec2 iInputIndices[]; +restrict readonly layout(std430, binding = 3)buffer bInputIndices { + uvec2 iInputIndices[]; }; -layout(std430, binding = 4)buffer bMicrolines { +restrict layout(std430, binding = 4)buffer bMicrolines { - restrict uvec4 iMicrolines[]; + uvec4 iMicrolines[]; }; void emitMicroline(vec4 microlineSegment, uint pathIndex, uint outputMicrolineIndex){ diff --git a/resources/shaders/gl4/d3d11/fill.cs.glsl b/resources/shaders/gl4/d3d11/fill.cs.glsl index ccc459d93..462c7f731 100644 --- a/resources/shaders/gl4/d3d11/fill.cs.glsl +++ b/resources/shaders/gl4/d3d11/fill.cs.glsl @@ -57,27 +57,33 @@ layout(local_size_x = 16, local_size_y = 4)in; + + + + + + layout(rgba8)uniform image2D uDest; uniform sampler2D uAreaLUT; uniform ivec2 uAlphaTileRange; -layout(std430, binding = 0)buffer bFills { - restrict readonly uint iFills[]; +restrict readonly layout(std430, binding = 0)buffer bFills { + uint iFills[]; }; -layout(std430, binding = 1)buffer bTiles { +restrict layout(std430, binding = 1)buffer bTiles { - restrict uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 2)buffer bAlphaTiles { +restrict readonly layout(std430, binding = 2)buffer bAlphaTiles { - restrict readonly uint iAlphaTiles[]; + uint iAlphaTiles[]; }; @@ -130,10 +136,17 @@ void main(){ int fillIndex = int(iTiles[tileIndex * 4 + 1]); int backdrop = int(iTiles[tileIndex * 4 + 3])>> 24; - vec4 coverages = vec4(backdrop); coverages += accumulateCoverageForFillList(fillIndex, tileSubCoord); - coverages = clamp(abs(coverages), 0.0, 1.0); + + uint tileControlWord = iTiles[tileIndex * 4 + 3]; + int tileCtrl = int((tileControlWord >> 16)& 0xffu); + int maskCtrl =(tileCtrl >> 0)& 0x3; + if((maskCtrl & 0x1)!= 0){ + coverages = clamp(abs(coverages), 0.0, 1.0); + } else { + coverages = clamp(1.0 - abs(1.0 - mod(coverages, 2.0)), 0.0, 1.0); + } int clipTileIndex = int(iAlphaTiles[batchAlphaTileIndex * 2 + 1]); diff --git a/resources/shaders/gl4/d3d11/propagate.cs.glsl b/resources/shaders/gl4/d3d11/propagate.cs.glsl index 744eaff50..608ace5c8 100644 --- a/resources/shaders/gl4/d3d11/propagate.cs.glsl +++ b/resources/shaders/gl4/d3d11/propagate.cs.glsl @@ -32,70 +32,76 @@ layout(local_size_x = 64)in; + + + + + + uniform ivec2 uFramebufferTileSize; uniform int uColumnCount; uniform int uFirstAlphaTileIndex; -layout(std430, binding = 0)buffer bDrawMetadata { +restrict readonly layout(std430, binding = 0)buffer bDrawMetadata { - restrict readonly uvec4 iDrawMetadata[]; + uvec4 iDrawMetadata[]; }; -layout(std430, binding = 1)buffer bClipMetadata { +restrict readonly layout(std430, binding = 1)buffer bClipMetadata { - restrict readonly uvec4 iClipMetadata[]; + uvec4 iClipMetadata[]; }; -layout(std430, binding = 2)buffer bBackdrops { +restrict readonly layout(std430, binding = 2)buffer bBackdrops { - restrict readonly int iBackdrops[]; + int iBackdrops[]; }; -layout(std430, binding = 3)buffer bDrawTiles { +restrict layout(std430, binding = 3)buffer bDrawTiles { - restrict uint iDrawTiles[]; + uint iDrawTiles[]; }; -layout(std430, binding = 4)buffer bClipTiles { +restrict layout(std430, binding = 4)buffer bClipTiles { - restrict uint iClipTiles[]; + uint iClipTiles[]; }; -layout(std430, binding = 5)buffer bZBuffer { +restrict layout(std430, binding = 5)buffer bZBuffer { - restrict int iZBuffer[]; + int iZBuffer[]; }; -layout(std430, binding = 6)buffer bFirstTileMap { - restrict int iFirstTileMap[]; +restrict layout(std430, binding = 6)buffer bFirstTileMap { + int iFirstTileMap[]; }; -layout(std430, binding = 7)buffer bAlphaTiles { +restrict layout(std430, binding = 7)buffer bAlphaTiles { - restrict uint iAlphaTiles[]; + uint iAlphaTiles[]; }; uint calculateTileIndex(uint bufferOffset, uvec4 tileRect, uvec2 tileCoord){ @@ -204,6 +210,16 @@ void main(){ drawTileWord |(uint(drawTileBackdrop)<< 24); + if(drawTileBackdrop != 0){ + int tileCtrl = int((drawTileWord >> 16)& 0xffu); + int maskCtrl =(tileCtrl >> 0)& 0x3; + + if((maskCtrl & 0x2)!= 0 && mod(abs(drawTileBackdrop), 2)== 0){ + zWrite = false; + } + } + + ivec2 tileCoord = ivec2(tileX, tileY)+ ivec2(drawTileRect . xy); int tileMapIndex = tileCoord . y * uFramebufferTileSize . x + tileCoord . x; if(zWrite && drawTileBackdrop != 0 && drawAlphaTileIndex < 0) diff --git a/resources/shaders/gl4/d3d11/sort.cs.glsl b/resources/shaders/gl4/d3d11/sort.cs.glsl index 5ecbcdb47..3a546205c 100644 --- a/resources/shaders/gl4/d3d11/sort.cs.glsl +++ b/resources/shaders/gl4/d3d11/sort.cs.glsl @@ -29,20 +29,20 @@ precision highp float; uniform int uTileCount; -layout(std430, binding = 0)buffer bTiles { +restrict layout(std430, binding = 0)buffer bTiles { - restrict uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 1)buffer bFirstTileMap { - restrict int iFirstTileMap[]; +restrict layout(std430, binding = 1)buffer bFirstTileMap { + int iFirstTileMap[]; }; -layout(std430, binding = 2)buffer bZBuffer { - restrict readonly int iZBuffer[]; +restrict readonly layout(std430, binding = 2)buffer bZBuffer { + int iZBuffer[]; }; layout(local_size_x = 64)in; diff --git a/resources/shaders/gl4/d3d11/tile.cs.glsl b/resources/shaders/gl4/d3d11/tile.cs.glsl index 09503d609..5b01e4086 100644 --- a/resources/shaders/gl4/d3d11/tile.cs.glsl +++ b/resources/shaders/gl4/d3d11/tile.cs.glsl @@ -695,6 +695,12 @@ void computeTileVaryings(vec2 position, + + + + + + uniform int uLoadAction; uniform vec4 uClearColor; uniform vec2 uTileSize; @@ -711,17 +717,17 @@ uniform vec2 uFramebufferSize; uniform ivec2 uFramebufferTileSize; layout(rgba8)uniform image2D uDestImage; -layout(std430, binding = 0)buffer bTiles { +restrict readonly layout(std430, binding = 0)buffer bTiles { - restrict readonly uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 1)buffer bFirstTileMap { - restrict readonly int iFirstTileMap[]; +restrict readonly layout(std430, binding = 1)buffer bFirstTileMap { + int iFirstTileMap[]; }; uint calculateTileIndex(uint bufferOffset, uvec4 tileRect, uvec2 tileCoord){ @@ -772,6 +778,16 @@ void main(){ } else { backdrop = int(tileControlWord)>> 24; + + + if(backdrop != 0){ + int maskCtrl =(tileCtrl >> 0)& 0x3; + + if((maskCtrl & 0x2)!= 0 && mod(abs(backdrop), 2)== 0){ + break; + } + } + maskTileCoord = uvec2(0u); tileCtrl &= ~(0x3 << 0); } diff --git a/resources/shaders/metal/d3d11/bin.cs.metal b/resources/shaders/metal/d3d11/bin.cs.metal index 7008d475c..12f2b978b 100644 --- a/resources/shaders/metal/d3d11/bin.cs.metal +++ b/resources/shaders/metal/d3d11/bin.cs.metal @@ -150,8 +150,8 @@ kernel void main0(constant int& uMaxFillCount [[buffer(2)]], constant int& uMicr 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); + uint pathTileOffset = uint(((device int*)&_360.iMetadata[(pathIndex * 3u) + 1u])[0u]); + uint pathBackdropOffset = uint(((device int*)&_360.iMetadata[(pathIndex * 3u) + 2u])[0u]); int2 tileSize = int2(16); int4 tileLineSegment = int4(floor(lineSegment / float4(tileSize.xyxy))); int2 fromTileCoords = tileLineSegment.xy; diff --git a/resources/shaders/metal/d3d11/bound.cs.metal b/resources/shaders/metal/d3d11/bound.cs.metal index b1a4bd013..432600ecd 100644 --- a/resources/shaders/metal/d3d11/bound.cs.metal +++ b/resources/shaders/metal/d3d11/bound.cs.metal @@ -41,7 +41,7 @@ kernel void main0(constant int& uTileCount [[buffer(0)]], constant int& uPathCou if (_50) { uint midPathIndex = lowPathIndex + ((highPathIndex - lowPathIndex) / 2u); - uint midTileIndex = _64.iTilePathInfo[midPathIndex].z; + uint midTileIndex = ((device uint*)&_64.iTilePathInfo[midPathIndex])[2u]; if (tileIndex < midTileIndex) { highPathIndex = midPathIndex; diff --git a/resources/shaders/metal/d3d11/dice.cs.metal b/resources/shaders/metal/d3d11/dice.cs.metal index 54913e557..47857d868 100644 --- a/resources/shaders/metal/d3d11/dice.cs.metal +++ b/resources/shaders/metal/d3d11/dice.cs.metal @@ -100,7 +100,7 @@ kernel void main0(constant int& uMaxMicrolineCount [[buffer(0)]], constant int& if (_241) { uint midPathIndex = lowPathIndex + ((highPathIndex - lowPathIndex) / 2u); - uint midBatchSegmentIndex = _253.iDiceMetadata[midPathIndex].z; + uint midBatchSegmentIndex = ((device uint*)&_253.iDiceMetadata[midPathIndex])[2u]; if (batchSegmentIndex < midBatchSegmentIndex) { highPathIndex = midPathIndex; diff --git a/resources/shaders/metal/d3d11/fill.cs.metal b/resources/shaders/metal/d3d11/fill.cs.metal index c3f47df7c..a2fbd0400 100644 --- a/resources/shaders/metal/d3d11/fill.cs.metal +++ b/resources/shaders/metal/d3d11/fill.cs.metal @@ -23,6 +23,13 @@ struct bTiles constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(16u, 4u, 1u); +// Implementation of the GLSL mod() function, which is slightly different than Metal fmod() +template +inline Tx mod(Tx x, Ty y) +{ + return x - y * floor(x / y); +} + static inline __attribute__((always_inline)) float4 computeCoverage(thread const float2& from, thread const float2& to, thread const texture2d areaLUT, thread const sampler areaLUTSmplr) { @@ -87,7 +94,17 @@ kernel void main0(constant int2& uAlphaTileRange [[buffer(1)]], const device bFi int2 param_1 = tileSubCoord; float4 _334 = accumulateCoverageForFillList(param, param_1, v_148, uAreaLUT, uAreaLUTSmplr); coverages += _334; - coverages = fast::clamp(abs(coverages), float4(0.0), float4(1.0)); + uint tileControlWord = _294.iTiles[(tileIndex * 4u) + 3u]; + int tileCtrl = int((tileControlWord >> uint(16)) & 255u); + int maskCtrl = (tileCtrl >> 0) & 3; + if ((maskCtrl & 1) != 0) + { + coverages = fast::clamp(abs(coverages), float4(0.0), float4(1.0)); + } + else + { + coverages = fast::clamp(float4(1.0) - abs(float4(1.0) - mod(coverages, float4(2.0))), float4(0.0), float4(1.0)); + } int clipTileIndex = int(_284.iAlphaTiles[(batchAlphaTileIndex * 2u) + 1u]); if (clipTileIndex >= 0) { diff --git a/resources/shaders/metal/d3d11/propagate.cs.metal b/resources/shaders/metal/d3d11/propagate.cs.metal index 92d18dabf..4bc4bbda1 100644 --- a/resources/shaders/metal/d3d11/propagate.cs.metal +++ b/resources/shaders/metal/d3d11/propagate.cs.metal @@ -50,13 +50,20 @@ struct bFirstTileMap constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u); +// Implementation of the GLSL mod() function, which is slightly different than Metal fmod() +template +inline Tx mod(Tx x, Ty y) +{ + return x - y * floor(x / y); +} + static inline __attribute__((always_inline)) uint calculateTileIndex(thread const uint& bufferOffset, thread const uint4& tileRect, thread const uint2& tileCoord) { return (bufferOffset + (tileCoord.y * (tileRect.z - tileRect.x))) + tileCoord.x; } -kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirstAlphaTileIndex [[buffer(8)]], constant int2& uFramebufferTileSize [[buffer(9)]], const device bBackdrops& _59 [[buffer(1)]], const device bDrawMetadata& _85 [[buffer(2)]], const device bClipMetadata& _126 [[buffer(3)]], device bDrawTiles& _175 [[buffer(4)]], device bClipTiles& _252 [[buffer(5)]], device bZBuffer& _302 [[buffer(6)]], device bAlphaTiles& _310 [[buffer(7)]], device bFirstTileMap& _395 [[buffer(10)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirstAlphaTileIndex [[buffer(8)]], constant int2& uFramebufferTileSize [[buffer(9)]], const device bBackdrops& _59 [[buffer(1)]], const device bDrawMetadata& _85 [[buffer(2)]], const device bClipMetadata& _126 [[buffer(3)]], device bDrawTiles& _175 [[buffer(4)]], device bClipTiles& _252 [[buffer(5)]], device bZBuffer& _302 [[buffer(6)]], device bAlphaTiles& _310 [[buffer(7)]], device bFirstTileMap& _427 [[buffer(10)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint columnIndex = gl_GlobalInvocationID.x; if (int(columnIndex) >= uColumnCount) @@ -157,16 +164,35 @@ kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirst } _175.iDrawTiles[(drawTileIndex * 4u) + 2u] = (uint(drawAlphaTileIndex) & 16777215u) | (uint(drawBackdropDelta) << uint(24)); _175.iDrawTiles[(drawTileIndex * 4u) + 3u] = drawTileWord | (uint(drawTileBackdrop) << uint(24)); + if (drawTileBackdrop != 0) + { + int tileCtrl = int((drawTileWord >> uint(16)) & 255u); + int maskCtrl = (tileCtrl >> 0) & 3; + bool _365 = (maskCtrl & 2) != 0; + bool _376; + if (_365) + { + _376 = mod(float(abs(drawTileBackdrop)), 2.0) == 0.0; + } + else + { + _376 = _365; + } + if (_376) + { + zWrite = false; + } + } int2 tileCoord_1 = int2(tileX, int(tileY)) + int2(drawTileRect.xy); int tileMapIndex = (tileCoord_1.y * uFramebufferTileSize.x) + tileCoord_1.x; if ((zWrite && (drawTileBackdrop != 0)) && (drawAlphaTileIndex < 0)) { - int _383 = atomic_fetch_max_explicit((device atomic_int*)&_302.iZBuffer[tileMapIndex + 8], int(drawTileIndex), memory_order_relaxed); + int _415 = atomic_fetch_max_explicit((device atomic_int*)&_302.iZBuffer[tileMapIndex + 8], int(drawTileIndex), memory_order_relaxed); } if ((drawTileBackdrop != 0) || (drawAlphaTileIndex >= 0)) { - int _400 = atomic_exchange_explicit((device atomic_int*)&_395.iFirstTileMap[tileMapIndex], int(drawTileIndex), memory_order_relaxed); - int nextTileIndex = _400; + int _432 = atomic_exchange_explicit((device atomic_int*)&_427.iFirstTileMap[tileMapIndex], int(drawTileIndex), memory_order_relaxed); + int nextTileIndex = _432; _175.iDrawTiles[(drawTileIndex * 4u) + 0u] = uint(nextTileIndex); } currentBackdrop += drawBackdropDelta; diff --git a/resources/shaders/metal/d3d11/tile.cs.metal b/resources/shaders/metal/d3d11/tile.cs.metal index 6817716be..cbcd7c458 100644 --- a/resources/shaders/metal/d3d11/tile.cs.metal +++ b/resources/shaders/metal/d3d11/tile.cs.metal @@ -729,6 +729,24 @@ kernel void main0(constant int2& uFramebufferTileSize [[buffer(3)]], constant in else { backdrop = int(tileControlWord) >> 24; + if (backdrop != 0) + { + int maskCtrl = (tileCtrl >> 0) & 3; + bool _1751 = (maskCtrl & 2) != 0; + bool _1759; + if (_1751) + { + _1759 = mod(float(abs(backdrop)), 2.0) == 0.0; + } + else + { + _1759 = _1751; + } + if (_1759) + { + break; + } + } maskTileCoord = uint2(0u); tileCtrl &= (-4); } diff --git a/shaders/d3d11/bin.cs.glsl b/shaders/d3d11/bin.cs.glsl index d34260eb4..af134cd91 100644 --- a/shaders/d3d11/bin.cs.glsl +++ b/shaders/d3d11/bin.cs.glsl @@ -37,18 +37,18 @@ uniform int uMicrolineCount; // How many slots we have allocated for fills. uniform int uMaxFillCount; -layout(std430, binding = 0) buffer bMicrolines { - restrict readonly uvec4 iMicrolines[]; +restrict readonly layout(std430, binding = 0) buffer bMicrolines { + uvec4 iMicrolines[]; }; -layout(std430, binding = 1) buffer bMetadata { +restrict readonly layout(std430, binding = 1) buffer bMetadata { // [0]: tile rect // [1].x: tile offset // [1].y: path ID // [1].z: z write flag // [1].w: clip path ID // [2].x: backdrop offset - restrict readonly ivec4 iMetadata[]; + ivec4 iMetadata[]; }; // [0]: vertexCount (6) @@ -56,27 +56,27 @@ layout(std430, binding = 1) buffer bMetadata { // [2]: vertexStart (0) // [3]: baseInstance (0) // [4]: alpha tile count -layout(std430, binding = 2) buffer bIndirectDrawParams { - restrict uint iIndirectDrawParams[]; +restrict layout(std430, binding = 2) buffer bIndirectDrawParams { + uint iIndirectDrawParams[]; }; -layout(std430, binding = 3) buffer bFills { - restrict writeonly uint iFills[]; +restrict writeonly layout(std430, binding = 3) buffer bFills { + uint iFills[]; }; -layout(std430, binding = 4) buffer bTiles { +restrict layout(std430, binding = 4) buffer bTiles { // [0]: next tile ID (initialized to -1) // [1]: first fill ID (initialized to -1) // [2]: backdrop delta upper 8 bits, alpha tile ID lower 24 (initialized to 0, -1 respectively) // [3]: color/ctrl/backdrop word - restrict uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 5) buffer bBackdrops { +restrict layout(std430, binding = 5) buffer bBackdrops { // [0]: backdrop // [1]: tile X offset // [2]: path ID - restrict uint iBackdrops[]; + uint iBackdrops[]; }; uint computeTileIndexNoCheck(ivec2 tileCoords, ivec4 pathTileRect, uint pathTileOffset) { diff --git a/shaders/d3d11/bound.cs.glsl b/shaders/d3d11/bound.cs.glsl index 6f29b2f25..616859342 100644 --- a/shaders/d3d11/bound.cs.glsl +++ b/shaders/d3d11/bound.cs.glsl @@ -30,20 +30,20 @@ layout(local_size_x = 64) in; uniform int uPathCount; uniform int uTileCount; -layout(std430, binding = 0) buffer bTilePathInfo { +restrict readonly layout(std430, binding = 0) buffer bTilePathInfo { // x: tile upper left, 16-bit packed x/y // y: tile lower right, 16-bit packed x/y // z: first tile index in this path // w: color/ctrl/backdrop word - restrict readonly uvec4 iTilePathInfo[]; + uvec4 iTilePathInfo[]; }; -layout(std430, binding = 1) buffer bTiles { +restrict layout(std430, binding = 1) buffer bTiles { // [0]: next tile ID (initialized to -1) // [1]: first fill ID (initialized to -1) // [2]: backdrop delta upper 8 bits, alpha tile ID lower 24 (initialized to 0, -1 respectively) // [3]: color/ctrl/backdrop word - restrict uint iTiles[]; + uint iTiles[]; }; void main() { diff --git a/shaders/d3d11/dice.cs.glsl b/shaders/d3d11/dice.cs.glsl index cc1ab72ca..a13eb217a 100644 --- a/shaders/d3d11/dice.cs.glsl +++ b/shaders/d3d11/dice.cs.glsl @@ -40,37 +40,37 @@ uniform int uPathCount; uniform int uLastBatchSegmentIndex; uniform int uMaxMicrolineCount; -layout(std430, binding = 0) buffer bComputeIndirectParams { +restrict layout(std430, binding = 0) buffer bComputeIndirectParams { // [0]: number of x workgroups // [1]: number of y workgroups (always 1) // [2]: number of z workgroups (always 1) // [3]: number of output microlines - restrict uint iComputeIndirectParams[]; + uint iComputeIndirectParams[]; }; // Indexed by batch path index. -layout(std430, binding = 1) buffer bDiceMetadata { +restrict readonly layout(std430, binding = 1) buffer bDiceMetadata { // x: global path ID // y: first global segment index // z: first batch segment index // w: unused - restrict readonly uvec4 iDiceMetadata[]; + uvec4 iDiceMetadata[]; }; -layout(std430, binding = 2) buffer bPoints { - restrict readonly vec2 iPoints[]; +restrict readonly layout(std430, binding = 2) buffer bPoints { + vec2 iPoints[]; }; -layout(std430, binding = 3) buffer bInputIndices { - restrict readonly uvec2 iInputIndices[]; +restrict readonly layout(std430, binding = 3) buffer bInputIndices { + uvec2 iInputIndices[]; }; -layout(std430, binding = 4) buffer bMicrolines { +restrict layout(std430, binding = 4) buffer bMicrolines { // x: from (X, Y) whole pixels, packed signed 16-bit // y: to (X, Y) whole pixels, packed signed 16-bit // z: (from X, from Y, to X, to Y) fractional pixels, packed unsigned 8-bit (0.8 fixed point) // w: path ID - restrict uvec4 iMicrolines[]; + uvec4 iMicrolines[]; }; void emitMicroline(vec4 microlineSegment, uint pathIndex, uint outputMicrolineIndex) { diff --git a/shaders/d3d11/fill.cs.glsl b/shaders/d3d11/fill.cs.glsl index c3e2ff351..e49961fac 100644 --- a/shaders/d3d11/fill.cs.glsl +++ b/shaders/d3d11/fill.cs.glsl @@ -27,27 +27,33 @@ layout(local_size_x = 16, local_size_y = 4) in; #define TILE_FIELD_BACKDROP_ALPHA_TILE_ID 2 #define TILE_FIELD_CONTROL 3 +#define TILE_CTRL_MASK_MASK 0x3 +#define TILE_CTRL_MASK_WINDING 0x1 +#define TILE_CTRL_MASK_EVEN_ODD 0x2 + +#define TILE_CTRL_MASK_0_SHIFT 0 + layout(rgba8) uniform image2D uDest; uniform sampler2D uAreaLUT; uniform ivec2 uAlphaTileRange; -layout(std430, binding = 0) buffer bFills { - restrict readonly uint iFills[]; +restrict readonly layout(std430, binding = 0) buffer bFills { + uint iFills[]; }; -layout(std430, binding = 1) buffer bTiles { +restrict layout(std430, binding = 1) buffer bTiles { // [0]: path ID // [1]: next tile ID // [2]: first fill ID // [3]: backdrop delta upper 8 bits, alpha tile ID lower 24 bits // [4]: color/ctrl/backdrop word - restrict uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 2) buffer bAlphaTiles { +restrict readonly layout(std430, binding = 2) buffer bAlphaTiles { // [0]: alpha tile index // [1]: clip tile index - restrict readonly uint iAlphaTiles[]; + uint iAlphaTiles[]; }; #include "fill_compute.inc.glsl" @@ -74,10 +80,17 @@ void main() { int fillIndex = int(iTiles[tileIndex * 4 + TILE_FIELD_FIRST_FILL_ID]); int backdrop = int(iTiles[tileIndex * 4 + TILE_FIELD_CONTROL]) >> 24; - // TODO(pcwalton): Handle even-odd fill rule. vec4 coverages = vec4(backdrop); coverages += accumulateCoverageForFillList(fillIndex, tileSubCoord); - coverages = clamp(abs(coverages), 0.0, 1.0); + + uint tileControlWord = iTiles[tileIndex * 4 + TILE_FIELD_CONTROL]; + int tileCtrl = int((tileControlWord >> 16) & 0xffu); + int maskCtrl = (tileCtrl >> TILE_CTRL_MASK_0_SHIFT) & TILE_CTRL_MASK_MASK; + if ((maskCtrl & TILE_CTRL_MASK_WINDING) != 0) { + coverages = clamp(abs(coverages), 0.0, 1.0); + } else { + coverages = clamp(1.0 - abs(1.0 - mod(coverages, 2.0)), 0.0, 1.0); + } // Handle clip if necessary. int clipTileIndex = int(iAlphaTiles[batchAlphaTileIndex * 2 + 1]); diff --git a/shaders/d3d11/propagate.cs.glsl b/shaders/d3d11/propagate.cs.glsl index 5a382c792..1b74d75e8 100644 --- a/shaders/d3d11/propagate.cs.glsl +++ b/shaders/d3d11/propagate.cs.glsl @@ -30,70 +30,76 @@ layout(local_size_x = 64) in; #define FILL_INDIRECT_DRAW_PARAMS_ALPHA_TILE_COUNT_INDEX 4 #define FILL_INDIRECT_DRAW_PARAMS_SIZE 8 +#define TILE_CTRL_MASK_MASK 0x3 +#define TILE_CTRL_MASK_WINDING 0x1 +#define TILE_CTRL_MASK_EVEN_ODD 0x2 + +#define TILE_CTRL_MASK_0_SHIFT 0 + uniform ivec2 uFramebufferTileSize; uniform int uColumnCount; uniform int uFirstAlphaTileIndex; -layout(std430, binding = 0) buffer bDrawMetadata { +restrict readonly layout(std430, binding = 0) buffer bDrawMetadata { // [0]: tile rect // [1].x: tile offset // [1].y: path ID // [1].z: Z write enabled? // [1].w: clip path ID, or ~0 // [2].x: backdrop column offset - restrict readonly uvec4 iDrawMetadata[]; + uvec4 iDrawMetadata[]; }; -layout(std430, binding = 1) buffer bClipMetadata { +restrict readonly layout(std430, binding = 1) buffer bClipMetadata { // [0]: tile rect // [1].x: tile offset // [1].y: unused // [1].z: unused // [1].w: unused - restrict readonly uvec4 iClipMetadata[]; + uvec4 iClipMetadata[]; }; -layout(std430, binding = 2) buffer bBackdrops { +restrict readonly layout(std430, binding = 2) buffer bBackdrops { // [0]: backdrop // [1]: tile X offset // [2]: path ID - restrict readonly int iBackdrops[]; + int iBackdrops[]; }; -layout(std430, binding = 3) buffer bDrawTiles { +restrict layout(std430, binding = 3) buffer bDrawTiles { // [0]: next tile ID // [1]: first fill ID // [2]: backdrop delta upper 8 bits, alpha tile ID lower 24 // [3]: color/ctrl/backdrop word - restrict uint iDrawTiles[]; + uint iDrawTiles[]; }; -layout(std430, binding = 4) buffer bClipTiles { +restrict layout(std430, binding = 4) buffer bClipTiles { // [0]: next tile ID // [1]: first fill ID // [2]: backdrop delta upper 8 bits, alpha tile ID lower 24 // [3]: color/ctrl/backdrop word - restrict uint iClipTiles[]; + uint iClipTiles[]; }; -layout(std430, binding = 5) buffer bZBuffer { +restrict layout(std430, binding = 5) buffer bZBuffer { // [0]: vertexCount (6) // [1]: instanceCount (of fills) // [2]: vertexStart (0) // [3]: baseInstance (0) // [4]: alpha tile count // [8..]: z-buffer - restrict int iZBuffer[]; + int iZBuffer[]; }; -layout(std430, binding = 6) buffer bFirstTileMap { - restrict int iFirstTileMap[]; +restrict layout(std430, binding = 6) buffer bFirstTileMap { + int iFirstTileMap[]; }; -layout(std430, binding = 7) buffer bAlphaTiles { +restrict layout(std430, binding = 7) buffer bAlphaTiles { // [0]: alpha tile index // [1]: clip tile index - restrict uint iAlphaTiles[]; + uint iAlphaTiles[]; }; uint calculateTileIndex(uint bufferOffset, uvec4 tileRect, uvec2 tileCoord) { @@ -201,6 +207,16 @@ void main() { iDrawTiles[drawTileIndex * 4 + TILE_FIELD_CONTROL] = drawTileWord | (uint(drawTileBackdrop) << 24); + // Even-Odd fill rule will make some solid tiles invisible, so we shouldn't write them into the Z buffer. + if (drawTileBackdrop != 0) { + int tileCtrl = int((drawTileWord >> 16) & 0xffu); + int maskCtrl = (tileCtrl >> TILE_CTRL_MASK_0_SHIFT) & TILE_CTRL_MASK_MASK; + + if ((maskCtrl & TILE_CTRL_MASK_EVEN_ODD) != 0 && mod(abs(drawTileBackdrop), 2) == 0) { + zWrite = false; + } + } + // Write to Z-buffer if necessary. ivec2 tileCoord = ivec2(tileX, tileY) + ivec2(drawTileRect.xy); int tileMapIndex = tileCoord.y * uFramebufferTileSize.x + tileCoord.x; diff --git a/shaders/d3d11/sort.cs.glsl b/shaders/d3d11/sort.cs.glsl index a7436cb19..1bb7f0c40 100644 --- a/shaders/d3d11/sort.cs.glsl +++ b/shaders/d3d11/sort.cs.glsl @@ -27,20 +27,20 @@ precision highp sampler2D; uniform int uTileCount; -layout(std430, binding = 0) buffer bTiles { +restrict layout(std430, binding = 0) buffer bTiles { // [0]: next tile ID // [1]: first fill ID // [2]: backdrop delta upper 8 bits, alpha tile ID lower 24 // [3]: color/ctrl/backdrop word - restrict uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 1) buffer bFirstTileMap { - restrict int iFirstTileMap[]; +restrict layout(std430, binding = 1) buffer bFirstTileMap { + int iFirstTileMap[]; }; -layout(std430, binding = 2) buffer bZBuffer { - restrict readonly int iZBuffer[]; +restrict readonly layout(std430, binding = 2) buffer bZBuffer { + int iZBuffer[]; }; layout(local_size_x = 64) in; diff --git a/shaders/d3d11/tile.cs.glsl b/shaders/d3d11/tile.cs.glsl index 945b7fe87..8c3d7a2de 100644 --- a/shaders/d3d11/tile.cs.glsl +++ b/shaders/d3d11/tile.cs.glsl @@ -31,6 +31,12 @@ layout(local_size_x = 16, local_size_y = 4) in; #define TILE_FIELD_BACKDROP_ALPHA_TILE_ID 2 #define TILE_FIELD_CONTROL 3 +#define TILE_CTRL_MASK_MASK 0x3 +#define TILE_CTRL_MASK_WINDING 0x1 +#define TILE_CTRL_MASK_EVEN_ODD 0x2 + +#define TILE_CTRL_MASK_0_SHIFT 0 + uniform int uLoadAction; uniform vec4 uClearColor; uniform vec2 uTileSize; @@ -47,17 +53,17 @@ uniform vec2 uFramebufferSize; uniform ivec2 uFramebufferTileSize; layout(rgba8) uniform image2D uDestImage; -layout(std430, binding = 0) buffer bTiles { +restrict readonly layout(std430, binding = 0) buffer bTiles { // [0]: path ID // [1]: next tile ID // [2]: first fill ID // [3]: backdrop delta upper 8 bits, alpha tile ID lower 24 bits // [4]: color/ctrl/backdrop word - restrict readonly uint iTiles[]; + uint iTiles[]; }; -layout(std430, binding = 1) buffer bFirstTileMap { - restrict readonly int iFirstTileMap[]; +restrict readonly layout(std430, binding = 1) buffer bFirstTileMap { + int iFirstTileMap[]; }; uint calculateTileIndex(uint bufferOffset, uvec4 tileRect, uvec2 tileCoord) { @@ -108,6 +114,16 @@ void main() { } else { // We have no alpha mask. Clear the mask bits so we don't try to look one up. backdrop = int(tileControlWord) >> 24; + + // Handle solid tiles affected by the even-odd fill rule. + if (backdrop != 0) { + int maskCtrl = (tileCtrl >> TILE_CTRL_MASK_0_SHIFT) & TILE_CTRL_MASK_MASK; + + if ((maskCtrl & TILE_CTRL_MASK_EVEN_ODD) != 0 && mod(abs(backdrop), 2) == 0) { + break; + } + } + maskTileCoord = uvec2(0u); tileCtrl &= ~(TILE_CTRL_MASK_MASK << TILE_CTRL_MASK_0_SHIFT); }