[mtl renderer] fixed clipping

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);

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,
+		mp_rect clip = {x, y, w, h};
-	                                        .rect = (mp_rect){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);
 	}
 }
 

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_JUMP
-	           MG_CMD_CLIP_PUSH,
-	           MG_CMD_CLIP_POP,
 	     } mg_primitive_cmd;
 
 typedef struct mg_primitive

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 {

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 setBytes:&tileSize length:sizeof(int) atIndex:7];
-		[mergeEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:8];
+		[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 setBytes:&scale length:sizeof(float) atIndex:5];
-		[rasterEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:6];
+		[rasterEncoder setBytes:&backend->msaaCount length:sizeof(int) atIndex:6];
-		[rasterEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:7];
+		[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),
+			path->box = (vector_float4){context.pathScreenExtents.x,
-		                                minimum(primitive->attributes.clip.x + primitive->attributes.clip.w, context.pathScreenExtents.z),
+		                                context.pathScreenExtents.y,
-		                                minimum(primitive->attributes.clip.y + primitive->attributes.clip.h, context.pathScreenExtents.w)};
+		                                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,

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]);
+	//NOTE: we don't clip on the right, since we need those tiles to accurately compute
-	int nTilesX = lastTile.x - firstTile.x + 1;
+	//      the prefix sum of winding increments in the backprop pass.
-	int nTilesY = lastTile.y - firstTile.y + 1;
+	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)]],
+                      constant int* tileSize [[buffer(7)]],
-                      device atomic_int* logOffsetBuffer [[buffer(8)]],
+                      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)]],
+                       constant float* scale [[buffer(5)]],
-                       device char* logBuffer [[buffer(6)]],
+                       constant int* sampleCountBuffer [[buffer(6)]],
-                       device atomic_int* logOffsetBuffer [[buffer(7)]],
+                       device char* logBuffer [[buffer(7)]],
-                       constant int* useTexture [[buffer(8)]],
+                       device atomic_int* logOffsetBuffer [[buffer(8)]],
+                       constant int* useTexture [[buffer(9)]],
                        texture2d<float, access::write> outTexture [[texture(0)]],
                        texture2d<float> srcTexture [[texture(1)]],
                        uint2 threadCoord [[thread_position_in_grid]],
@@ -1502,6 +1527,14 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 			pathColor.rgb *= pathColor.a;
 
 			for(int sampleIndex=0; sampleIndex<sampleCount; sampleIndex++)
+			{
+				float2 sampleCoord = sampleCoords[sampleIndex];
+				float4 clip = pathBuffer[pathIndex].clip * scale[0];
+
+				if(  sampleCoord.x >= clip.x
+				  && sampleCoord.x < clip.z
+				  && sampleCoord.y >= clip.y
+				  && sampleCoord.y < clip.w)
 				{
 					bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1))
 					            ||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0));
@@ -1510,8 +1543,8 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 						float4 nextColor = pathColor;
 						if(useTexture[0])
 						{
-						float3 sampleCoord = float3(sampleCoords[sampleIndex].xy, 1);
+							float3 ph = float3(sampleCoords[sampleIndex].xy, 1);
-						float2 uv = (pathBuffer[pathIndex].uvTransform * sampleCoord).xy;
+							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);
@@ -1521,6 +1554,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 						}
 						color[sampleIndex] = color[sampleIndex]*(1-nextColor.a) + nextColor;
 					}
+				}
 				winding[sampleIndex] = op->windingOffset;
 			}
 			pathIndex = op->index;
@@ -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))
@@ -1563,6 +1598,14 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 	pathColor.rgb *= pathColor.a;
 
 	for(int sampleIndex=0; sampleIndex<sampleCount; sampleIndex++)
+	{
+		float2 sampleCoord = sampleCoords[sampleIndex];
+		float4 clip = pathBuffer[pathIndex].clip * scale[0];
+
+		if(  sampleCoord.x >= clip.x
+		  && sampleCoord.x < clip.z
+		  && sampleCoord.y >= clip.y
+		  && sampleCoord.y < clip.w)
 		{
 			bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1))
 			            ||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0));
@@ -1582,6 +1625,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 				}
 				color[sampleIndex] = color[sampleIndex]*(1-nextColor.a) + nextColor;
 			}
+		}
 		pixelColor += color[sampleIndex];
 	}
 	pixelColor /= sampleCount;