From f35e91a38cd54fcefef0caa19e8b878062c5fc17 Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Tue, 11 Apr 2023 14:51:58 +0200 Subject: [PATCH] [mtl renderer] fixed clipping --- examples/canvas/main.c | 7 ++ src/graphics.c | 148 ++++++++++++++++++---------------------- src/graphics_internal.h | 4 +- src/mtl_renderer.h | 1 + src/mtl_renderer.m | 30 +++++--- src/mtl_renderer.metal | 118 ++++++++++++++++++++++---------- 6 files changed, 176 insertions(+), 132 deletions(-) diff --git a/examples/canvas/main.c b/examples/canvas/main.c index 7632e3d..d241ec4 100644 --- a/examples/canvas/main.c +++ b/examples/canvas/main.c @@ -160,6 +160,11 @@ int main() mg_set_color_rgba(0, 1, 1, 1); mg_clear(); + mg_set_color_rgba(1, 0, 0, 1); + mg_set_width(2); + mg_rectangle_stroke(304, 100, 300, 250); + + mg_clip_push(304, 100, 300, 250); // head mg_set_color_rgba(1, 1, 0, 1); @@ -178,6 +183,8 @@ int main() mg_ellipse_fill(x-70, y-50, 30, 50); mg_ellipse_fill(x+70, y-50, 30, 50); + mg_clip_pop(); + // text mg_set_color_rgba(0, 0, 1, 1); mg_set_font(font); diff --git a/src/graphics.c b/src/graphics.c index 24538c2..6c6d0a9 100644 --- a/src/graphics.c +++ b/src/graphics.c @@ -638,6 +638,43 @@ void mg_matrix_stack_pop(mg_canvas_data* canvas) } } +mp_rect mg_clip_stack_top(mg_canvas_data* canvas) +{ + if(canvas->clipStackSize == 0) + { + return((mp_rect){-FLT_MAX/2, -FLT_MAX/2, FLT_MAX, FLT_MAX}); + } + else + { + return(canvas->clipStack[canvas->clipStackSize-1]); + } +} + +void mg_clip_stack_push(mg_canvas_data* canvas, mp_rect clip) +{ + if(canvas->clipStackSize >= MG_CLIP_STACK_MAX_DEPTH) + { + LOG_ERROR("clip stack overflow\n"); + } + else + { + canvas->clipStack[canvas->clipStackSize] = clip; + canvas->clipStackSize++; + } +} + +void mg_clip_stack_pop(mg_canvas_data* canvas) +{ + if(canvas->clipStackSize == 0) + { + LOG_ERROR("clip stack underflow\n"); + } + else + { + canvas->clipStackSize--; + } +} + void mg_push_command(mg_canvas_data* canvas, mg_primitive primitive) { //NOTE(martin): push primitive and updates current stream, eventually patching a pending jump. @@ -646,6 +683,7 @@ void mg_push_command(mg_canvas_data* canvas, mg_primitive primitive) canvas->primitives[canvas->primitiveCount] = primitive; canvas->primitives[canvas->primitiveCount].attributes = canvas->attributes; canvas->primitives[canvas->primitiveCount].attributes.transform = mg_matrix_stack_top(canvas); + canvas->primitives[canvas->primitiveCount].attributes.clip = mg_clip_stack_top(canvas); canvas->primitiveCount++; } @@ -725,6 +763,7 @@ u32 mg_next_shape(mg_canvas_data* canvas, mg_attributes* attributes) { mg_finalize_shape(canvas); + canvas->clip = attributes->clip; canvas->transform = attributes->transform; canvas->srcRegion = attributes->srcRegion; canvas->shapeExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX}; @@ -2698,71 +2737,6 @@ vec2 mg_canvas_size(void) } //////////////////////////////////////////////////////////// - -mp_rect mg_clip_stack_top(mg_canvas_data* canvas) -{ - if(canvas->clipStackSize == 0) - { - return((mp_rect){-FLT_MAX/2, -FLT_MAX/2, FLT_MAX, FLT_MAX}); - } - else - { - return(canvas->clipStack[canvas->clipStackSize-1]); - } -} - -void mg_clip_stack_push(mg_canvas_data* canvas, mp_rect clip) -{ - if(canvas->clipStackSize >= MG_CLIP_STACK_MAX_DEPTH) - { - LOG_ERROR("clip stack overflow\n"); - } - else - { - canvas->clipStack[canvas->clipStackSize] = clip; - canvas->clipStackSize++; - canvas->clip = clip; - } -} - -void mg_clip_stack_pop(mg_canvas_data* canvas) -{ - if(canvas->clipStackSize == 0) - { - LOG_ERROR("clip stack underflow\n"); - } - else - { - canvas->clipStackSize--; - canvas->clip = mg_clip_stack_top(canvas); - } -} - -void mg_do_clip_push(mg_canvas_data* canvas, mp_rect clip) -{ - //NOTE(martin): transform clip - vec2 p0 = mg_mat2x3_mul(canvas->transform, (vec2){clip.x, clip.y}); - vec2 p1 = mg_mat2x3_mul(canvas->transform, (vec2){clip.x + clip.w, clip.y}); - vec2 p2 = mg_mat2x3_mul(canvas->transform, (vec2){clip.x + clip.w, clip.y + clip.h}); - vec2 p3 = mg_mat2x3_mul(canvas->transform, (vec2){clip.x, clip.y + clip.h}); - - f32 x0 = minimum(p0.x, minimum(p1.x, minimum(p2.x, p3.x))); - f32 y0 = minimum(p0.y, minimum(p1.y, minimum(p2.y, p3.y))); - f32 x1 = maximum(p0.x, maximum(p1.x, maximum(p2.x, p3.x))); - f32 y1 = maximum(p0.y, maximum(p1.y, maximum(p2.y, p3.y))); - - mp_rect current = mg_clip_stack_top(canvas); - - //NOTE(martin): intersect with current clip - x0 = maximum(current.x, x0); - y0 = maximum(current.y, y0); - x1 = minimum(current.x + current.w, x1); - y1 = minimum(current.y + current.h, y1); - - mp_rect r = {x0, y0, maximum(0, x1-x0), maximum(0, y1-y0)}; - mg_clip_stack_push(canvas, r); -} - void mg_draw_batch(mg_canvas_data* canvas, mg_image_data* image) { mg_finalize_shape(canvas); @@ -2857,19 +2831,6 @@ void mg_flush_commands(int primitiveCount, mg_primitive* primitives, mg_path_elt nextIndex = primitive->jump; } } break; - - case MG_CMD_CLIP_PUSH: - { - //TODO(martin): use only aligned rect and avoid this - mp_rect r = {primitive->rect.x, primitive->rect.y, primitive->rect.w, primitive->rect.h}; - mg_do_clip_push(canvas, r); - } break; - - case MG_CMD_CLIP_POP: - { - mg_clip_stack_pop(canvas); - } break; - } } exit_command_loop: ; @@ -2941,8 +2902,32 @@ void mg_clip_push(f32 x, f32 y, f32 w, f32 h) mg_canvas_data* canvas = __mgCurrentCanvas; if(canvas) { - mg_push_command(canvas, (mg_primitive){.cmd = MG_CMD_CLIP_PUSH, - .rect = (mp_rect){x, y, w, h}}); + mp_rect clip = {x, y, w, h}; + + //NOTE(martin): transform clip + mg_mat2x3 transform = mg_matrix_stack_top(canvas); + vec2 p0 = mg_mat2x3_mul(transform, (vec2){clip.x, clip.y}); + vec2 p1 = mg_mat2x3_mul(transform, (vec2){clip.x + clip.w, clip.y}); + vec2 p2 = mg_mat2x3_mul(transform, (vec2){clip.x + clip.w, clip.y + clip.h}); + vec2 p3 = mg_mat2x3_mul(transform, (vec2){clip.x, clip.y + clip.h}); + + f32 x0 = minimum(p0.x, minimum(p1.x, minimum(p2.x, p3.x))); + f32 y0 = minimum(p0.y, minimum(p1.y, minimum(p2.y, p3.y))); + f32 x1 = maximum(p0.x, maximum(p1.x, maximum(p2.x, p3.x))); + f32 y1 = maximum(p0.y, maximum(p1.y, maximum(p2.y, p3.y))); + + mp_rect current = mg_clip_stack_top(canvas); + + //NOTE(martin): intersect with current clip + x0 = maximum(current.x, x0); + y0 = maximum(current.y, y0); + x1 = minimum(current.x + current.w, x1); + y1 = minimum(current.y + current.h, y1); + + mp_rect r = {x0, y0, maximum(0, x1-x0), maximum(0, y1-y0)}; + mg_clip_stack_push(canvas, r); + + canvas->attributes.clip = r; } } @@ -2951,7 +2936,8 @@ void mg_clip_pop() mg_canvas_data* canvas = __mgCurrentCanvas; if(canvas) { - mg_push_command(canvas, (mg_primitive){.cmd = MG_CMD_CLIP_POP}); + mg_clip_stack_pop(canvas); + canvas->attributes.clip = mg_clip_stack_top(canvas); } } diff --git a/src/graphics_internal.h b/src/graphics_internal.h index 330fff6..5e40f9c 100644 --- a/src/graphics_internal.h +++ b/src/graphics_internal.h @@ -178,9 +178,7 @@ typedef struct mg_rounded_rect typedef enum { MG_CMD_FILL, MG_CMD_STROKE, - MG_CMD_JUMP, - MG_CMD_CLIP_PUSH, - MG_CMD_CLIP_POP, + MG_CMD_JUMP } mg_primitive_cmd; typedef struct mg_primitive diff --git a/src/mtl_renderer.h b/src/mtl_renderer.h index cdf0d8a..f2739d2 100644 --- a/src/mtl_renderer.h +++ b/src/mtl_renderer.h @@ -22,6 +22,7 @@ typedef struct mg_mtl_path matrix_float3x3 uvTransform; vector_float4 color; vector_float4 box; + vector_float4 clip; } mg_mtl_path; typedef enum { diff --git a/src/mtl_renderer.m b/src/mtl_renderer.m index 183405b..8cbd90b 100644 --- a/src/mtl_renderer.m +++ b/src/mtl_renderer.m @@ -708,8 +708,10 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend, [mergeEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:4]; [mergeEncoder setBuffer:backend->tileOpCountBuffer offset:0 atIndex:5]; [mergeEncoder setBuffer:backend->screenTilesBuffer offset:0 atIndex:6]; - [mergeEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:7]; - [mergeEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:8]; + [mergeEncoder setBytes:&tileSize length:sizeof(int) atIndex:7]; + [mergeEncoder setBytes:&scale length:sizeof(float) atIndex:8]; + [mergeEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:9]; + [mergeEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:10]; MTLSize mergeGridSize = MTLSizeMake(nTilesX, nTilesY, 1); MTLSize mergeGroupSize = MTLSizeMake(16, 16, 1); @@ -727,9 +729,10 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend, [rasterEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:backend->pathBufferOffset atIndex:2]; [rasterEncoder setBuffer:backend->segmentBuffer offset:0 atIndex:3]; [rasterEncoder setBytes:&tileSize length:sizeof(int) atIndex:4]; - [rasterEncoder setBytes:&backend->msaaCount length:sizeof(int) atIndex:5]; - [rasterEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:6]; - [rasterEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:7]; + [rasterEncoder setBytes:&scale length:sizeof(float) atIndex:5]; + [rasterEncoder setBytes:&backend->msaaCount length:sizeof(int) atIndex:6]; + [rasterEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:7]; + [rasterEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:8]; [rasterEncoder setTexture:backend->outTexture atIndex:0]; @@ -740,8 +743,7 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend, [rasterEncoder setTexture: mtlImage->texture atIndex: 1]; useTexture = 1; } - [rasterEncoder setBytes: &useTexture length:sizeof(int) atIndex: 8]; - + [rasterEncoder setBytes: &useTexture length:sizeof(int) atIndex: 9]; MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1); MTLSize rasterGroupSize = MTLSizeMake(16, 16, 1); @@ -959,10 +961,16 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface, pathCount++; path->cmd = (mg_mtl_cmd)primitive->cmd; - path->box = (vector_float4){maximum(primitive->attributes.clip.x, context.pathScreenExtents.x), - maximum(primitive->attributes.clip.y, context.pathScreenExtents.y), - minimum(primitive->attributes.clip.x + primitive->attributes.clip.w, context.pathScreenExtents.z), - minimum(primitive->attributes.clip.y + primitive->attributes.clip.h, context.pathScreenExtents.w)}; + + path->box = (vector_float4){context.pathScreenExtents.x, + context.pathScreenExtents.y, + context.pathScreenExtents.z, + context.pathScreenExtents.w}; + + path->clip = (vector_float4){primitive->attributes.clip.x, + primitive->attributes.clip.y, + primitive->attributes.clip.x + primitive->attributes.clip.w, + primitive->attributes.clip.y + primitive->attributes.clip.h}; path->color = (vector_float4){primitive->attributes.color.r, primitive->attributes.color.g, diff --git a/src/mtl_renderer.metal b/src/mtl_renderer.metal index 74b3379..9f604bf 100644 --- a/src/mtl_renderer.metal +++ b/src/mtl_renderer.metal @@ -237,10 +237,19 @@ kernel void mtl_path_setup(constant int* pathCount [[buffer(0)]], { const device mg_mtl_path* path = &pathBuffer[pathIndex]; - int2 firstTile = int2(path->box.xy*scale[0])/tileSize[0]; - int2 lastTile = max(firstTile, int2(path->box.zw*scale[0])/tileSize[0]); - int nTilesX = lastTile.x - firstTile.x + 1; - int nTilesY = lastTile.y - firstTile.y + 1; + + //NOTE: we don't clip on the right, since we need those tiles to accurately compute + // the prefix sum of winding increments in the backprop pass. + float4 clippedBox = {max(path->box.x, path->clip.x), + max(path->box.y, path->clip.y), + path->box.z, + min(path->box.w, path->clip.w)}; + + int2 firstTile = int2(clippedBox.xy*scale[0])/tileSize[0]; + int2 lastTile = int2(clippedBox.zw*scale[0])/tileSize[0]; + + int nTilesX = max(0, lastTile.x - firstTile.x + 1); + int nTilesY = max(0, lastTile.y - firstTile.y + 1); int tileCount = nTilesX * nTilesY; int tileQueuesIndex = atomic_fetch_add_explicit(tileQueueCount, tileCount, memory_order_relaxed); @@ -1360,8 +1369,10 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]], device mg_mtl_tile_op* tileOpBuffer [[buffer(4)]], device atomic_int* tileOpCount [[buffer(5)]], device int* screenTilesBuffer [[buffer(6)]], - device char* logBuffer [[buffer(7)]], - device atomic_int* logOffsetBuffer [[buffer(8)]], + constant int* tileSize [[buffer(7)]], + constant float* scale [[buffer(8)]], + device char* logBuffer [[buffer(9)]], + device atomic_int* logOffsetBuffer [[buffer(10)]], uint2 threadCoord [[thread_position_in_grid]], uint2 gridSize [[threads_per_grid]]) { @@ -1380,8 +1391,13 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]], const device mg_mtl_path_queue* pathQueue = &pathQueueBuffer[pathIndex]; int2 pathTileCoord = tileCoord - pathQueue->area.xy; + const device mg_mtl_path* path = &pathBuffer[pathIndex]; + float xMax = min(path->box.z, path->clip.z); + int tileMax = xMax * scale[0] / tileSize[0]; + int pathTileMax = tileMax - pathQueue->area.x; + if( pathTileCoord.x >= 0 - && pathTileCoord.x < pathQueue->area.z + && pathTileCoord.x <= pathTileMax && pathTileCoord.y >= 0 && pathTileCoord.y < pathQueue->area.w) { @@ -1396,7 +1412,7 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]], if(windingOffset & 1) { //NOTE: tile is full covered. Add path start op (with winding offset). - // Additionally if color is opaque, trim tile list. + // Additionally if color is opaque and tile is fully inside clip, trim tile list. int pathOpIndex = atomic_fetch_add_explicit(tileOpCount, 1, memory_order_relaxed); device mg_mtl_tile_op* pathOp = &tileOpBuffer[pathOpIndex]; pathOp->kind = MG_MTL_OP_START; @@ -1404,7 +1420,15 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]], pathOp->index = pathIndex; pathOp->windingOffset = windingOffset; - if(pathBuffer[pathIndex].color.a == 1) + float4 clip = pathBuffer[pathIndex].clip * scale[0]; + float4 tileBox = float4(tileCoord.x, tileCoord.y, tileCoord.x+1, tileCoord.y+1); + tileBox *= tileSize[0]; + + if(pathBuffer[pathIndex].color.a == 1 + && tileBox.x >= clip.x + && tileBox.z < clip.z + && tileBox.y >= clip.y + && tileBox.w < clip.w) { screenTilesBuffer[tileIndex] = pathOpIndex; } @@ -1444,10 +1468,11 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]], const device mg_mtl_path* pathBuffer [[buffer(2)]], const device mg_mtl_segment* segmentBuffer [[buffer(3)]], constant int* tileSize [[buffer(4)]], - constant int* sampleCountBuffer [[buffer(5)]], - device char* logBuffer [[buffer(6)]], - device atomic_int* logOffsetBuffer [[buffer(7)]], - constant int* useTexture [[buffer(8)]], + constant float* scale [[buffer(5)]], + constant int* sampleCountBuffer [[buffer(6)]], + device char* logBuffer [[buffer(7)]], + device atomic_int* logOffsetBuffer [[buffer(8)]], + constant int* useTexture [[buffer(9)]], texture2d outTexture [[texture(0)]], texture2d srcTexture [[texture(1)]], uint2 threadCoord [[thread_position_in_grid]], @@ -1503,23 +1528,32 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]], for(int sampleIndex=0; sampleIndex= clip.x + && sampleCoord.x < clip.z + && sampleCoord.y >= clip.y + && sampleCoord.y < clip.w) { - float4 nextColor = pathColor; - if(useTexture[0]) + bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1)) + ||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0)); + if(filled) { - float3 sampleCoord = float3(sampleCoords[sampleIndex].xy, 1); - float2 uv = (pathBuffer[pathIndex].uvTransform * sampleCoord).xy; + float4 nextColor = pathColor; + if(useTexture[0]) + { + float3 ph = float3(sampleCoords[sampleIndex].xy, 1); + float2 uv = (pathBuffer[pathIndex].uvTransform * ph).xy; - constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); - float4 texColor = srcTexture.sample(smp, uv); - texColor.rgb *= texColor.a; + constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); + float4 texColor = srcTexture.sample(smp, uv); + texColor.rgb *= texColor.a; - nextColor *= texColor; + nextColor *= texColor; + } + color[sampleIndex] = color[sampleIndex]*(1-nextColor.a) + nextColor; } - color[sampleIndex] = color[sampleIndex]*(1-nextColor.a) + nextColor; } winding[sampleIndex] = op->windingOffset; } @@ -1533,6 +1567,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]], { float2 sampleCoord = sampleCoords[sampleIndex]; + //TODO: shouldn't this be redundant with mtl_side_of_segment()? if( (sampleCoord.y > seg->box.y) &&(sampleCoord.y <= seg->box.w) &&(mtl_side_of_segment(sampleCoord, seg) < 0)) @@ -1564,23 +1599,32 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]], for(int sampleIndex=0; sampleIndex= clip.x + && sampleCoord.x < clip.z + && sampleCoord.y >= clip.y + && sampleCoord.y < clip.w) { - float4 nextColor = pathColor; - if(useTexture[0]) + bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1)) + ||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0)); + if(filled) { - float3 sampleCoord = float3(sampleCoords[sampleIndex].xy, 1); - float2 uv = (pathBuffer[pathIndex].uvTransform * sampleCoord).xy; + float4 nextColor = pathColor; + if(useTexture[0]) + { + float3 sampleCoord = float3(sampleCoords[sampleIndex].xy, 1); + float2 uv = (pathBuffer[pathIndex].uvTransform * sampleCoord).xy; - constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); - float4 texColor = srcTexture.sample(smp, uv); - texColor.rgb *= texColor.a; + constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); + float4 texColor = srcTexture.sample(smp, uv); + texColor.rgb *= texColor.a; - nextColor *= texColor; + nextColor *= texColor; + } + color[sampleIndex] = color[sampleIndex]*(1-nextColor.a) + nextColor; } - color[sampleIndex] = color[sampleIndex]*(1-nextColor.a) + nextColor; } pixelColor += color[sampleIndex]; }