hmn
/
orca
5
Fork 2
Code Issues 21 Pull Requests 1 Packages Projects Releases Wiki Activity

[mtl canvas] Fix shader contents scaling

This commit is contained in:
Martin Fouilleul 2023-04-05 10:31:50 +02:00
parent 4c4f27066b
commit c4e866d9d4
3 changed files with 15 additions and 11 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
examples/polygon/main.c
View File

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

6
src/mtl_renderer.m
View File

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

16
src/mtl_renderer.metal
View File

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