diff --git a/examples/polygon/main.c b/examples/polygon/main.c index c865373..9aebb5f 100644 --- a/examples/polygon/main.c +++ b/examples/polygon/main.c @@ -145,8 +145,8 @@ int main() mg_set_color_rgba(0, 0, 1, 1); mg_stroke(); - mg_move_to(x+2*400, y+2*400); - mg_cubic_to(x+2*400, y+2*200, x+2*600, y+2*500, x+2*600, y+2*400); + mg_move_to(x+400, y+300); + mg_cubic_to(x+400, y+100, x+600, y+400, x+600, y+300); mg_close_path(); mg_set_color_rgba(0, 0, 1, 1); mg_stroke(); diff --git a/src/mtl_renderer.m b/src/mtl_renderer.m index 5589ce0..3c86418 100644 --- a/src/mtl_renderer.m +++ b/src/mtl_renderer.m @@ -689,6 +689,7 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface, [pathEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:3]; [pathEncoder setBuffer:backend->tileQueueCountBuffer offset:0 atIndex:4]; [pathEncoder setBytes:&tileSize length:sizeof(int) atIndex:5]; + [pathEncoder setBytes:&scale length:sizeof(int) atIndex:6]; MTLSize pathGridSize = MTLSizeMake(pathCount, 1, 1); MTLSize pathGroupSize = MTLSizeMake([backend->pathPipeline maxTotalThreadsPerThreadgroup], 1, 1); @@ -710,8 +711,9 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface, [segmentEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:6]; [segmentEncoder setBuffer:backend->tileOpCountBuffer offset:0 atIndex:7]; [segmentEncoder setBytes:&tileSize length:sizeof(int) atIndex:8]; - [segmentEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:9]; - [segmentEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:10]; + [segmentEncoder setBytes:&scale length:sizeof(int) atIndex:9]; + [segmentEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:10]; + [segmentEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:11]; MTLSize segmentGridSize = MTLSizeMake(context.mtlEltCount, 1, 1); MTLSize segmentGroupSize = MTLSizeMake([backend->segmentPipeline maxTotalThreadsPerThreadgroup], 1, 1); diff --git a/src/mtl_renderer.metal b/src/mtl_renderer.metal index c32d055..d5eff37 100644 --- a/src/mtl_renderer.metal +++ b/src/mtl_renderer.metal @@ -171,12 +171,13 @@ kernel void mtl_path_setup(constant int* pathCount [[buffer(0)]], device mg_mtl_tile_queue* tileQueueBuffer [[buffer(3)]], device atomic_int* tileQueueCount [[buffer(4)]], constant int* tileSize [[buffer(5)]], + constant float* scale [[buffer(6)]], uint pathIndex [[thread_position_in_grid]]) { const device mg_mtl_path* path = &pathBuffer[pathIndex]; - int2 firstTile = int2(path->box.xy)/tileSize[0]; - int2 lastTile = max(firstTile, int2(path->box.zw)/tileSize[0]); + 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; int tileCount = nTilesX * nTilesY; @@ -965,9 +966,10 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]], device mg_mtl_tile_op* tileOpBuffer [[buffer(6)]], device atomic_int* tileOpCount [[buffer(7)]], constant int* tileSize [[buffer(8)]], + constant float* scale [[buffer(9)]], - device char* logBuffer [[buffer(9)]], - device atomic_int* logOffsetBuffer [[buffer(10)]], + device char* logBuffer [[buffer(10)]], + device atomic_int* logOffsetBuffer [[buffer(11)]], uint eltIndex [[thread_position_in_grid]]) { const device mg_mtl_path_elt* elt = &elementBuffer[eltIndex]; @@ -991,19 +993,19 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]], { case MG_MTL_LINE: { - float2 p[2] = {elt->p[0], elt->p[1]}; + float2 p[2] = {elt->p[0]*scale[0], elt->p[1]*scale[0]}; mtl_line_setup(&setupCtx, p); } break; case MG_MTL_QUADRATIC: { - float2 p[3] = {elt->p[0], elt->p[1], elt->p[2]}; + float2 p[3] = {elt->p[0]*scale[0], elt->p[1]*scale[0], elt->p[2]*scale[0]}; mtl_quadratic_setup(&setupCtx, p); } break; case MG_MTL_CUBIC: { - float2 p[4] = {elt->p[0], elt->p[1], elt->p[2], elt->p[3]}; + float2 p[4] = {elt->p[0]*scale[0], elt->p[1]*scale[0], elt->p[2]*scale[0], elt->p[3]*scale[0]}; mtl_cubic_setup(&setupCtx, p); } break;