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

[mtl canvas, wip] Fixing inclusive/exclusive bound checks during curve side tests / shortcut / rasterization

This commit is contained in:
Martin Fouilleul 2023-04-04 19:02:53 +02:00
parent 9322db8201
commit 8d7fdf3b1a
4 changed files with 555 additions and 230 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

57
examples/polygon/main.c
View File

@ -47,6 +47,7 @@ int main()
mp_window_focus(window);
f64 frameTime = 0;
f32 x = 0, y = 0;
while(!mp_should_quit())
{
@ -63,6 +64,29 @@ int main()
mp_request_quit();
} break;
case MP_EVENT_KEYBOARD_KEY:
{
if(event.key.action == MP_KEY_PRESS)
{
if(event.key.code == MP_KEY_LEFT)
{
x-=1;
}
if(event.key.code == MP_KEY_RIGHT)
{
x+=1;
}
if(event.key.code == MP_KEY_UP)
{
y-=1;
}
if(event.key.code == MP_KEY_DOWN)
{
y+=1;
}
}
} break;
default:
break;
}
@ -95,13 +119,44 @@ int main()
mg_close_path();
mg_set_color_rgba(0, 0, 1, 1);
mg_fill();
*/
mg_move_to(2*400, 2*400);
mg_cubic_to(2*400, 2*200, 2*600, 2*500, 2*600, 2*400);
mg_close_path();
mg_set_color_rgba(0, 0, 1, 1);
mg_fill();
*/
mg_move_to(200, 200);
mg_line_to(300, 300);
mg_line_to(200, 400);
mg_line_to(100, 300);
mg_close_path();
mg_set_color_rgba(1, 0, 0, 1);
// mg_set_width(2);
mg_stroke();
mg_move_to(400, 400);
mg_quadratic_to(600, 601, 800, 400);
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_close_path();
mg_set_color_rgba(0, 0, 1, 1);
mg_stroke();
/*
mg_move_to(x+8, y+8);
mg_line_to(x+33, y+8);
mg_line_to(x+33, y+19);
mg_line_to(x+8, y+19);
mg_close_path();
mg_set_color_rgba(0, 0, 1, 1);
mg_fill();
*/
printf("Milepost vector graphics test program (frame time = %fs, fps = %f)...\n",
frameTime,
1./frameTime);

1
src/mtl_renderer.h
View File

@ -13,6 +13,7 @@
typedef enum {
MG_MTL_FILL,
MG_MTL_STROKE,
} mg_mtl_cmd;
typedef struct mg_mtl_path

364
src/mtl_renderer.m
View File

@ -83,6 +83,233 @@ void mg_mtl_print_log(int bufferIndex, id<MTLBuffer> logBuffer, id<MTLBuffer> lo
}
typedef struct mg_mtl_encoding_context
{
int mtlEltCount;
mg_mtl_path_elt* elementBufferData;
int pathIndex;
mg_primitive* primitive;
vec4 pathExtents;
} mg_mtl_encoding_context;
void mg_mtl_canvas_encode_element(mg_mtl_encoding_context* context, mg_path_elt_type kind, vec2* p)
{
mg_mtl_path_elt* mtlElt = &context->elementBufferData[context->mtlEltCount];
context->mtlEltCount++;
mtlElt->pathIndex = context->pathIndex;
int count = 0;
switch(kind)
{
case MG_PATH_LINE:
mtlElt->kind = MG_MTL_LINE;
count = 2;
break;
case MG_PATH_QUADRATIC:
mtlElt->kind = MG_MTL_QUADRATIC;
count = 3;
break;
case MG_PATH_CUBIC:
mtlElt->kind = MG_MTL_CUBIC;
count = 4;
break;
default:
break;
}
for(int i=0; i<count; i++)
{
mg_update_path_extents(&context->pathExtents, p[i]);
vec2 screenP = mg_mat2x3_mul(context->primitive->attributes.transform, p[i]);
mtlElt->p[i] = (vector_float2){screenP.x, screenP.y};
}
}
void mg_mtl_canvas_stroke_line(mg_mtl_encoding_context* context, vec2* p)
{
f32 width = context->primitive->attributes.width;
vec2 v = {p[1].x-p[0].x, p[1].y-p[0].y};
vec2 n = {v.y, -v.x};
f32 norm = sqrt(n.x*n.x + n.y*n.y);
vec2 offset = vec2_mul(0.5*width/norm, n);
vec2 left[2] = {vec2_add(p[0], offset), vec2_add(p[1], offset)};
vec2 right[2] = {vec2_add(p[1], vec2_mul(-1, offset)), vec2_add(p[0], vec2_mul(-1, offset))};
vec2 joint0[2] = {vec2_add(p[0], vec2_mul(-1, offset)), vec2_add(p[0], offset)};
vec2 joint1[2] = {vec2_add(p[1], offset), vec2_add(p[1], vec2_mul(-1, offset))};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, right);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, left);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint0);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint1);
}
void mg_mtl_canvas_stroke_quadratic(mg_mtl_encoding_context* context, vec2* p)
{
f32 width = context->primitive->attributes.width;
f32 tolerance = minimum(context->primitive->attributes.tolerance, 0.5 * width);
vec2 leftHull[3];
vec2 rightHull[3];
if( !mg_offset_hull(3, p, leftHull, width/2)
|| !mg_offset_hull(3, p, rightHull, -width/2))
{
//TODO split and recurse
//NOTE: offsetting the hull failed, split the curve
vec2 splitLeft[3];
vec2 splitRight[3];
mg_quadratic_split(p, 0.5, splitLeft, splitRight);
mg_mtl_canvas_stroke_quadratic(context, splitLeft);
mg_mtl_canvas_stroke_quadratic(context, splitRight);
}
else
{
const int CHECK_SAMPLE_COUNT = 5;
f32 checkSamples[CHECK_SAMPLE_COUNT] = {1./6, 2./6, 3./6, 4./6, 5./6};
f32 d2LowBound = Square(0.5 * width - tolerance);
f32 d2HighBound = Square(0.5 * width + tolerance);
f32 maxOvershoot = 0;
f32 maxOvershootParameter = 0;
for(int i=0; i<CHECK_SAMPLE_COUNT; i++)
{
f32 t = checkSamples[i];
vec2 c = mg_quadratic_get_point(p, t);
vec2 cp = mg_quadratic_get_point(leftHull, t);
vec2 cn = mg_quadratic_get_point(rightHull, t);
f32 positiveDistSquare = Square(c.x - cp.x) + Square(c.y - cp.y);
f32 negativeDistSquare = Square(c.x - cn.x) + Square(c.y - cn.y);
f32 positiveOvershoot = maximum(positiveDistSquare - d2HighBound, d2LowBound - positiveDistSquare);
f32 negativeOvershoot = maximum(negativeDistSquare - d2HighBound, d2LowBound - negativeDistSquare);
f32 overshoot = maximum(positiveOvershoot, negativeOvershoot);
if(overshoot > maxOvershoot)
{
maxOvershoot = overshoot;
maxOvershootParameter = t;
}
}
if(maxOvershoot > 0)
{
vec2 splitLeft[3];
vec2 splitRight[3];
mg_quadratic_split(p, maxOvershootParameter, splitLeft, splitRight);
mg_mtl_canvas_stroke_quadratic(context, splitLeft);
mg_mtl_canvas_stroke_quadratic(context, splitRight);
}
else
{
vec2 tmp = leftHull[0];
leftHull[0] = leftHull[2];
leftHull[2] = tmp;
mg_mtl_canvas_encode_element(context, MG_PATH_QUADRATIC, rightHull);
mg_mtl_canvas_encode_element(context, MG_PATH_QUADRATIC, leftHull);
vec2 joint0[2] = {rightHull[2], leftHull[0]};
vec2 joint1[2] = {leftHull[2], rightHull[0]};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint0);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint1);
}
}
}
void mg_mtl_canvas_stroke_cubic(mg_mtl_encoding_context* context, vec2* p)
{
f32 width = context->primitive->attributes.width;
f32 tolerance = minimum(context->primitive->attributes.tolerance, 0.5 * width);
vec2 leftHull[4];
vec2 rightHull[4];
if( !mg_offset_hull(4, p, leftHull, width/2)
|| !mg_offset_hull(4, p, rightHull, -width/2))
{
//TODO split and recurse
//NOTE: offsetting the hull failed, split the curve
vec2 splitLeft[4];
vec2 splitRight[4];
mg_cubic_split(p, 0.5, splitLeft, splitRight);
mg_mtl_canvas_stroke_cubic(context, splitLeft);
mg_mtl_canvas_stroke_cubic(context, splitRight);
}
else
{
const int CHECK_SAMPLE_COUNT = 5;
f32 checkSamples[CHECK_SAMPLE_COUNT] = {1./6, 2./6, 3./6, 4./6, 5./6};
f32 d2LowBound = Square(0.5 * width - tolerance);
f32 d2HighBound = Square(0.5 * width + tolerance);
f32 maxOvershoot = 0;
f32 maxOvershootParameter = 0;
for(int i=0; i<CHECK_SAMPLE_COUNT; i++)
{
f32 t = checkSamples[i];
vec2 c = mg_cubic_get_point(p, t);
vec2 cp = mg_cubic_get_point(leftHull, t);
vec2 cn = mg_cubic_get_point(rightHull, t);
f32 positiveDistSquare = Square(c.x - cp.x) + Square(c.y - cp.y);
f32 negativeDistSquare = Square(c.x - cn.x) + Square(c.y - cn.y);
f32 positiveOvershoot = maximum(positiveDistSquare - d2HighBound, d2LowBound - positiveDistSquare);
f32 negativeOvershoot = maximum(negativeDistSquare - d2HighBound, d2LowBound - negativeDistSquare);
f32 overshoot = maximum(positiveOvershoot, negativeOvershoot);
if(overshoot > maxOvershoot)
{
maxOvershoot = overshoot;
maxOvershootParameter = t;
}
}
if(maxOvershoot > 0)
{
vec2 splitLeft[4];
vec2 splitRight[4];
mg_cubic_split(p, maxOvershootParameter, splitLeft, splitRight);
mg_mtl_canvas_stroke_cubic(context, splitLeft);
mg_mtl_canvas_stroke_cubic(context, splitRight);
}
else
{
vec2 tmp = leftHull[0];
leftHull[0] = leftHull[3];
leftHull[3] = tmp;
tmp = leftHull[1];
leftHull[1] = leftHull[2];
leftHull[2] = tmp;
mg_mtl_canvas_encode_element(context, MG_PATH_CUBIC, rightHull);
mg_mtl_canvas_encode_element(context, MG_PATH_CUBIC, leftHull);
vec2 joint0[2] = {rightHull[3], leftHull[0]};
vec2 joint1[2] = {leftHull[3], rightHull[0]};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint0);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint1);
}
}
}
void mg_mtl_canvas_render(mg_canvas_backend* interface,
mg_color clearColor,
u32 primitiveCount,
@ -101,90 +328,83 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
//NOTE: fill renderer input buffers
int pathCount = 0;
int mtlEltCount = 0;
vec2 currentPos = {0};
mg_mtl_encoding_context context = {.mtlEltCount = 0,
.elementBufferData = elementBufferData};
for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++)
{
mg_primitive* primitive = &primitives[primitiveIndex];
if(primitive->cmd == MG_CMD_FILL && primitive->path.count)
if(primitive->path.count)
{
vec4 pathExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
context.primitive = primitive;
context.pathIndex = primitiveIndex;
context.pathExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
for(int eltIndex = 0;
(eltIndex < primitive->path.count) && (primitive->path.startIndex + eltIndex < eltCount);
eltIndex++)
{
mg_path_elt* elt = &pathElements[primitive->path.startIndex + eltIndex];
if(elt->type == MG_PATH_MOVE)
if(elt->type != MG_PATH_MOVE)
{
currentPos = elt->p[0];
vec2 p[4] = {currentPos, elt->p[0], elt->p[1], elt->p[2]};
if(primitive->cmd == MG_CMD_FILL)
{
mg_mtl_canvas_encode_element(&context, elt->type, p);
}
else if(primitive->cmd == MG_CMD_STROKE)
{
switch(elt->type)
{
case MG_PATH_LINE:
mg_mtl_canvas_stroke_line(&context, p);
break;
case MG_PATH_QUADRATIC:
mg_mtl_canvas_stroke_quadratic(&context, p);
break;
case MG_PATH_CUBIC:
mg_mtl_canvas_stroke_cubic(&context, p);
break;
default:
break;
}
}
}
else if(elt->type == MG_PATH_LINE)
switch(elt->type)
{
/////////////////////////////////////////////////////////////////////////////////////
//TODO: order control points so that we can collapse all elements into same codepath
/////////////////////////////////////////////////////////////////////////////////////
case MG_PATH_MOVE:
currentPos = elt->p[0];
break;
//NOTE: transform and push path elt + update primitive bounding box
vec2 p0 = mg_mat2x3_mul(primitive->attributes.transform, currentPos);
vec2 p1 = mg_mat2x3_mul(primitive->attributes.transform, elt->p[0]);
currentPos = elt->p[0];
case MG_PATH_LINE:
currentPos = elt->p[0];
break;
mg_update_path_extents(&pathExtents, p0);
mg_update_path_extents(&pathExtents, p1);
case MG_PATH_QUADRATIC:
currentPos = elt->p[1];
break;
mg_mtl_path_elt* mtlElt = &elementBufferData[mtlEltCount];
mtlEltCount++;
mtlElt->pathIndex = primitiveIndex;
mtlElt->kind = (mg_mtl_seg_kind)elt->type;
mtlElt->p[0] = (vector_float2){p0.x, p0.y};
mtlElt->p[1] = (vector_float2){p1.x, p1.y};
case MG_PATH_CUBIC:
currentPos = elt->p[2];
break;
}
else if(elt->type == MG_PATH_QUADRATIC)
{
vec2 p0 = mg_mat2x3_mul(primitive->attributes.transform, currentPos);
vec2 p1 = mg_mat2x3_mul(primitive->attributes.transform, elt->p[0]);
vec2 p2 = mg_mat2x3_mul(primitive->attributes.transform, elt->p[1]);
currentPos = elt->p[1];
}
mg_update_path_extents(&pathExtents, p0);
mg_update_path_extents(&pathExtents, p1);
mg_update_path_extents(&pathExtents, p2);
mg_mtl_path_elt* mtlElt = &elementBufferData[mtlEltCount];
mtlEltCount++;
mtlElt->pathIndex = primitiveIndex;
mtlElt->kind = (mg_mtl_seg_kind)elt->type;
mtlElt->p[0] = (vector_float2){p0.x, p0.y};
mtlElt->p[1] = (vector_float2){p1.x, p1.y};
mtlElt->p[2] = (vector_float2){p2.x, p2.y};
}
else if(elt->type == MG_PATH_CUBIC)
{
vec2 p0 = mg_mat2x3_mul(primitive->attributes.transform, currentPos);
vec2 p1 = mg_mat2x3_mul(primitive->attributes.transform, elt->p[0]);
vec2 p2 = mg_mat2x3_mul(primitive->attributes.transform, elt->p[1]);
vec2 p3 = mg_mat2x3_mul(primitive->attributes.transform, elt->p[2]);
currentPos = elt->p[2];
mg_update_path_extents(&pathExtents, p0);
mg_update_path_extents(&pathExtents, p1);
mg_update_path_extents(&pathExtents, p2);
mg_update_path_extents(&pathExtents, p3);
mg_mtl_path_elt* mtlElt = &elementBufferData[mtlEltCount];
mtlEltCount++;
mtlElt->pathIndex = primitiveIndex;
mtlElt->kind = (mg_mtl_seg_kind)elt->type;
mtlElt->p[0] = (vector_float2){p0.x, p0.y};
mtlElt->p[1] = (vector_float2){p1.x, p1.y};
mtlElt->p[2] = (vector_float2){p2.x, p2.y};
mtlElt->p[3] = (vector_float2){p3.x, p3.y};
}
if(primitive->cmd == MG_CMD_STROKE)
{
f32 margin = maximum(primitive->attributes.width, primitive->attributes.maxJointExcursion);
context.pathExtents.x -= margin;
context.pathExtents.y -= margin;
context.pathExtents.z += margin;
context.pathExtents.w += margin;
}
//NOTE: push path
@ -192,10 +412,10 @@ 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, pathExtents.x),
maximum(primitive->attributes.clip.y, pathExtents.y),
minimum(primitive->attributes.clip.x + primitive->attributes.clip.w, pathExtents.z),
minimum(primitive->attributes.clip.y + primitive->attributes.clip.h, pathExtents.w)};
path->box = (vector_float4){maximum(primitive->attributes.clip.x, context.pathExtents.x),
maximum(primitive->attributes.clip.y, context.pathExtents.y),
minimum(primitive->attributes.clip.x + primitive->attributes.clip.w, context.pathExtents.z),
minimum(primitive->attributes.clip.y + primitive->attributes.clip.h, context.pathExtents.w)};
path->color = (vector_float4){primitive->attributes.color.r,
primitive->attributes.color.g,
@ -269,7 +489,7 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
[segmentEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:9];
[segmentEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:10];
MTLSize segmentGridSize = MTLSizeMake(mtlEltCount, 1, 1);
MTLSize segmentGridSize = MTLSizeMake(context.mtlEltCount, 1, 1);
MTLSize segmentGroupSize = MTLSizeMake([backend->segmentPipeline maxTotalThreadsPerThreadgroup], 1, 1);
[segmentEncoder dispatchThreads: segmentGridSize threadsPerThreadgroup: segmentGroupSize];
@ -282,6 +502,8 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
[backpropEncoder setBuffer:backend->pathQueueBuffer offset:0 atIndex:0];
[backpropEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:1];
[backpropEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:2];
[backpropEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:3];
MTLSize backpropGroupSize = MTLSizeMake([backend->backpropPipeline maxTotalThreadsPerThreadgroup], 1, 1);
MTLSize backpropGridSize = MTLSizeMake(pathCount*backpropGroupSize.width, 1, 1);
@ -301,6 +523,8 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
[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];
MTLSize mergeGridSize = MTLSizeMake(nTilesX, nTilesY, 1);
MTLSize mergeGroupSize = MTLSizeMake(16, 16, 1);
@ -535,7 +759,7 @@ mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface)
bufferOptions = MTLResourceStorageModeShared;
for(int i=0; i<MG_MTL_INPUT_BUFFERS_COUNT; i++)
{
backend->logBuffer[i] = [metalSurface->device newBufferWithLength: 4<<20
backend->logBuffer[i] = [metalSurface->device newBufferWithLength: 1<<20
options: bufferOptions];
backend->logOffsetBuffer[i] = [metalSurface->device newBufferWithLength: sizeof(int)

363
src/mtl_renderer.metal
View File

@ -11,14 +11,14 @@ using namespace metal;
typedef struct mtl_log_context
{
device char* buffer;
device atomic_int* offset;
device volatile atomic_int* offset;
bool enabled;
} mtl_log_context;
int strlen(const constant char* msg)
{
int count = 0;
while(msg[count] != 0)
while(msg[count] != '\0')
{
count++;
}
@ -28,7 +28,7 @@ int strlen(const constant char* msg)
int strlen(const thread char* msg)
{
int count = 0;
while(msg[count] != 0)
while(msg[count] != '\0')
{
count++;
}
@ -42,10 +42,11 @@ void mtl_log(mtl_log_context context, const constant char* msg)
int len = strlen(msg);
int offset = atomic_fetch_add_explicit(context.offset, len+1, memory_order_relaxed);
for(int i=0; i<len+1; i++)
for(int i=0; i<len; i++)
{
context.buffer[offset+i] = msg[i];
}
context.buffer[offset+len] = '\0';
}
}
@ -56,10 +57,11 @@ void mtl_log(mtl_log_context context, const thread char* msg)
int len = strlen(msg);
int offset = atomic_fetch_add_explicit(context.offset, len+1, memory_order_relaxed);
for(int i=0; i<len+1; i++)
for(int i=0; i<len; i++)
{
context.buffer[offset+i] = msg[i];
}
context.buffer[offset+len] = '\0';
}
}
@ -112,6 +114,10 @@ void mtl_log_f32(mtl_log_context context, float value)
{
decimal /= 10;
}
if(decimal < 0)
{
decimal *= -1;
}
const int bufSize = 64;
char buffer[bufSize];
@ -138,6 +144,27 @@ void mtl_log_f32(mtl_log_context context, float value)
mtl_log(context, buffer);
}
void log_cubic_bezier(thread float2* p, mtl_log_context logCtx)
{
mtl_log(logCtx, "(");
mtl_log_f32(logCtx, p[0].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[0].y);
mtl_log(logCtx, ") (");
mtl_log_f32(logCtx, p[1].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[1].y);
mtl_log(logCtx, ") (");
mtl_log_f32(logCtx, p[2].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[2].y);
mtl_log(logCtx, ") (");
mtl_log_f32(logCtx, p[3].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[3].y);
mtl_log(logCtx, ")\n");
}
kernel void mtl_path_setup(constant int* pathCount [[buffer(0)]],
const device mg_mtl_path* pathBuffer [[buffer(1)]],
device mg_mtl_path_queue* pathQueueBuffer [[buffer(2)]],
@ -169,31 +196,34 @@ kernel void mtl_path_setup(constant int* pathCount [[buffer(0)]],
}
}
bool mtl_is_left_of_segment(float2 p, const device mg_mtl_segment* seg)
int mtl_side_of_segment(float2 p, const device mg_mtl_segment* seg, mtl_log_context log = {.enabled = false})
{
//NOTE: test is p is on the left of a curve segment.
/*WARN: if p is outside the bounding box of segment, we still consider it left from
the segment if it is left of its diagonal. This is done so that we can test
if tile corners are on the same side of the curve during tiling (corner are
not necessarily inside the bounding box, even if the tile itself overlaps
the curve).
During fine rasterization, this function need to be guarded by a the following
check: if(p.y >= seg->box.y && p.y < seg->box.w) {...}
*/
bool isLeft = false;
//NOTE: if point is left of curve bounding box, it is left of curve
if(p.x < seg->box.x)
int side = 0;
if(p.y > seg->box.w || p.y <= seg->box.y)
{
isLeft = true;
if(p.x > seg->box.x && p.x <= seg->box.z)
{
if(p.y > seg->box.w)
{
side = (seg->config == MG_MTL_TL || seg->config == MG_MTL_BR)? -1 : 1;
}
else
{
side = (seg->config == MG_MTL_TL || seg->config == MG_MTL_BR)? 1 : -1;
}
}
}
else if(p.x < seg->box.z)
else if(p.x > seg->box.z)
{
/*NOTE: if point and curve are on opposite sides of diagonal and on the left of diagonal,
it is left from the curve
otherwise if point and curve are on the same side of diagonal, do implicit test
*/
side = 1;
}
else if(p.x <= seg->box.x)
{
side = -1;
}
else
{
// eval based on diagonal
float alpha = (seg->box.w - seg->box.y)/(seg->box.z - seg->box.x);
float ofs = seg->box.w - seg->box.y;
float dx = p.x - seg->box.x;
@ -202,42 +232,38 @@ bool mtl_is_left_of_segment(float2 p, const device mg_mtl_segment* seg)
if( (seg->config == MG_MTL_BR && dy > alpha*dx)
||(seg->config == MG_MTL_TR && dy < ofs - alpha*dx))
{
isLeft = true;
side = -1;
}
else if( !(seg->config == MG_MTL_TL && dy < alpha*dx)
&& !(seg->config == MG_MTL_BL && dy > ofs - alpha*dx))
else if( (seg->config == MG_MTL_TL && dy < alpha*dx)
||(seg->config == MG_MTL_BL && dy > ofs - alpha*dx))
{
//NOTE: for lines, we only have config BR or TR, so the test is always negative
side = 1;
}
else
{
switch(seg->kind)
{
case MG_MTL_LINE:
side = 1;
break;
if(seg->kind == MG_MTL_QUADRATIC)
{
float3 ph = {p.x, p.y, 1};
float3 klm = seg->implicitMatrix * ph;
if((klm.x*klm.x - klm.y)*klm.z < 0)
case MG_MTL_QUADRATIC:
{
isLeft = true;
}
}
else if(seg->kind == MG_MTL_CUBIC)
{
/*
//DEBUG: behave as a straight line segment
if((seg->config == MG_MTL_BL || seg->config == MG_MTL_TL))
float3 ph = {p.x, p.y, 1};
float3 klm = seg->implicitMatrix * ph;
side = ((klm.x*klm.x - klm.y)*klm.z < 0)? -1 : 1;
} break;
case MG_MTL_CUBIC:
{
isLeft = true;
}
/*/
float3 ph = {p.x, p.y, 1};
float3 klm = seg->implicitMatrix * ph;
if(klm.x*klm.x*klm.x - klm.y*klm.z < 0)
{
isLeft = true;
}
//*/
float3 ph = {p.x, p.y, 1};
float3 klm = seg->implicitMatrix * ph;
side = (klm.x*klm.x*klm.x - klm.y*klm.z < 0)? -1 : 1;
} break;
}
}
}
return(isLeft);
return(side);
}
typedef struct mtl_segment_setup_context
@ -275,30 +301,76 @@ void mtl_segment_bin_to_tiles(thread mtl_segment_setup_context* context, device
float(x + pathArea.x + 1),
float(y + pathArea.y + 1)} * float(tileSize);
//NOTE: select two corners of tile box to test against the curve
float2 testPoint0;
float2 testPoint1;
if(seg->config == MG_MTL_BL || seg->config == MG_MTL_TR)
float2 bl = {tileBox.x, tileBox.y};
float2 br = {tileBox.z, tileBox.y};
float2 tr = {tileBox.z, tileBox.w};
float2 tl = {tileBox.x, tileBox.w};
int sbl = mtl_side_of_segment(bl, seg, context->log);
int sbr = mtl_side_of_segment(br, seg, context->log);
int str = mtl_side_of_segment(tr, seg, context->log);
int stl = mtl_side_of_segment(tl, seg, context->log);
bool crossL = (stl*sbl < 0);
bool crossR = (str*sbr < 0);
bool crossT = (stl*str < 0);
bool crossB = (sbl*sbr < 0);
mtl_log(context->log, "tile ");
mtl_log_i32(context->log, x);
mtl_log(context->log, ", ");
mtl_log_i32(context->log, y);
mtl_log(context->log, "\n");
mtl_log(context->log, "crossL ");
mtl_log_i32(context->log, crossL);
mtl_log(context->log, "\n");
mtl_log(context->log, "crossR ");
mtl_log_i32(context->log, crossR);
mtl_log(context->log, "\n");
mtl_log(context->log, "crossT ");
mtl_log_i32(context->log, crossT);
mtl_log(context->log, "\n");
mtl_log(context->log, "crossB ");
mtl_log_i32(context->log, crossB);
mtl_log(context->log, "\n");
float2 s0, s1;
if(seg->config == MG_MTL_TL||seg->config == MG_MTL_BR)
{
testPoint0 = (float2){tileBox.x, tileBox.y},
testPoint1 = (float2){tileBox.z, tileBox.w};
s0 = seg->box.xy;
s1 = seg->box.zw;
}
else
{
testPoint0 = (float2){tileBox.z, tileBox.y};
testPoint1 = (float2){tileBox.x, tileBox.w};
s0 = seg->box.xw;
s1 = seg->box.zy;
}
bool test0 = mtl_is_left_of_segment(testPoint0, seg);
bool test1 = mtl_is_left_of_segment(testPoint1, seg);
bool s0Inside = s0.x >= tileBox.x
&& s0.x < tileBox.z
&& s0.y >= tileBox.y
&& s0.y < tileBox.w;
//NOTE: the curve overlaps the tile only if test points are on opposite sides of segment
if(test0 != test1)
bool s1Inside = s1.x >= tileBox.x
&& s1.x < tileBox.z
&& s1.y >= tileBox.y
&& s1.y < tileBox.w;
mtl_log(context->log, "s0Inside ");
mtl_log_i32(context->log, s0Inside ? 1 : 0);
mtl_log(context->log, "\n");
mtl_log(context->log, "s1Inside ");
mtl_log_i32(context->log, s1Inside ? 1 : 0);
mtl_log(context->log, "\n");
if(crossL || crossR || crossT || crossB || s0Inside || s1Inside)
{
int tileOpIndex = atomic_fetch_add_explicit(context->tileOpCount, 1, memory_order_relaxed);
device mg_mtl_tile_op* op = &context->tileOpBuffer[tileOpIndex];
op->kind = MG_MTL_OP_SEGMENT;
op->index = segIndex;
op->crossRight = false;
op->next = -1;
int tileIndex = y*pathArea.z + x;
@ -310,43 +382,16 @@ void mtl_segment_bin_to_tiles(thread mtl_segment_setup_context* context, device
}
//NOTE: if the segment crosses the tile's bottom boundary, update the tile's winding offset
// testPoint0 is always a bottom point. We select the other one and check if they are on
// opposite sides of the curve.
// We also need to check that the endpoints of the curve are on opposite sides of the bottom
// boundary.
float2 testPoint3;
if(seg->config == MG_MTL_BL || seg->config == MG_MTL_TR)
{
testPoint3 = (float2){tileBox.z, tileBox.y};
}
else
{
testPoint3 = (float2){tileBox.x, tileBox.y};
}
bool test3 = mtl_is_left_of_segment(testPoint3, seg);
if( test0 != test3
&& seg->box.y < testPoint0.y
&& seg->box.w > testPoint0.y)
if(crossB)
{
atomic_fetch_add_explicit(&tile->windingOffset, seg->windingIncrement, memory_order_relaxed);
}
//NOTE: if the segment crosses the right boundary, mark it. We reuse one of the previous tests
float2 top = {tileBox.z, tileBox.w};
bool testTop = mtl_is_left_of_segment(top, seg);
bool testBottom = (seg->config == MG_MTL_BL || seg->config == MG_MTL_TR)? test3 : test0;
if(testTop != testBottom
&& seg->box.x <= top.x
&& seg->box.z > top.x)
if(crossR)
{
op->crossRight = true;
}
else
{
op->crossRight = false;
}
}
}
}
@ -387,9 +432,12 @@ device mg_mtl_segment* mtl_segment_push(thread mtl_segment_setup_context* contex
int segIndex = atomic_fetch_add_explicit(context->segmentCount, 1, memory_order_relaxed);
device mg_mtl_segment* seg = &context->segmentBuffer[segIndex];
bool goingUp = e.y >= s.y;
bool goingRight = e.x >= s.x;
seg->kind = kind;
seg->pathIndex = context->pathIndex;
seg->windingIncrement = (e.y > s.y)? 1 : -1;
seg->windingIncrement = goingUp? 1 : -1;
seg->box = (vector_float4){min(s.x, e.x),
min(s.y, e.y),
@ -401,25 +449,13 @@ device mg_mtl_segment* mtl_segment_push(thread mtl_segment_setup_context* contex
float alpha = (seg->box.w - seg->box.y)/(seg->box.z - seg->box.x);
float ofs = seg->box.w - seg->box.y;
//TODO: check that it works for line segments!
if( (e.x > s.x && e.y < s.y)
||(e.x <= s.x && e.y > s.y))
if(goingUp == goingRight)
{
if(dy < ofs - alpha*dx)
if(seg->kind == MG_MTL_LINE)
{
seg->config = MG_MTL_BL;
seg->config = MG_MTL_BR;
}
else
{
seg->config = MG_MTL_TR;
}
}
else if( (e.x > s.x && e.y >= s.y)
||(e.x <= s.x && e.y <= s.y))
{
//NOTE: it is important to include horizontal segments here, so that the mtl_is_left_of_segment() test
// becomes x > seg->box.x, in order to correctly detect right-crossing horizontal segments
if(dy > alpha*dx)
else if(dy > alpha*dx)
{
seg->config = MG_MTL_TL;
}
@ -428,6 +464,21 @@ device mg_mtl_segment* mtl_segment_push(thread mtl_segment_setup_context* contex
seg->config = MG_MTL_BR;
}
}
else
{
if(seg->kind == MG_MTL_LINE)
{
seg->config = MG_MTL_TR;
}
else if(dy < ofs - alpha*dx)
{
seg->config = MG_MTL_BL;
}
else
{
seg->config = MG_MTL_TR;
}
}
return(seg);
}
@ -545,27 +596,6 @@ int mtl_quadratic_roots(float a, float b, float c, thread float* r)
return(count);
}
void log_cubic_bezier(thread float2* p, mtl_log_context logCtx)
{
mtl_log(logCtx, "(");
mtl_log_f32(logCtx, p[0].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[0].y);
mtl_log(logCtx, ") (");
mtl_log_f32(logCtx, p[1].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[1].y);
mtl_log(logCtx, ") (");
mtl_log_f32(logCtx, p[2].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[2].y);
mtl_log(logCtx, ") (");
mtl_log_f32(logCtx, p[3].x);
mtl_log(logCtx, ", ");
mtl_log_f32(logCtx, p[3].y);
mtl_log(logCtx, ")\n");
}
void mtl_cubic_slice(float2 p[4], float s0, float s1, float2 sp[4])
{
float sr = (s1 - s0)/(1-s0);
@ -665,7 +695,7 @@ typedef struct mtl_cubic_info
} mtl_cubic_info;
mtl_cubic_info mtl_cubic_classify(thread float2* p)
mtl_cubic_info mtl_cubic_classify(thread float2* p, mtl_log_context log = {.enabled = false})
{
mtl_cubic_info result = {MTL_CUBIC_ERROR};
matrix_float4x4 F;
@ -787,7 +817,7 @@ mtl_cubic_info mtl_cubic_classify(thread float2* p)
F = (matrix_float4x4){{td*te, -se*td-sd*te, sd*se, 0},
{square(td)*te, -se*square(td)-2*sd*te*td, te*square(sd)+2*se*td*sd, -square(sd)*se},
{td*square(te), -sd*square(td)-2*se*td*te, td*square(se)+2*sd*te*se, -sd*square(se)},
{td*square(te), -sd*square(te)-2*se*td*te, td*square(se)+2*sd*te*se, -sd*square(se)},
{1, 0, 0, 0}};
//NOTE: if necessary, flip sign of k and l to ensure the interior is west from the curve
@ -885,6 +915,7 @@ void mtl_cubic_emit(thread mtl_segment_setup_context* context, float2 p[4], mtl_
//NOTE: set implicit matrix and bin segment
seg->implicitMatrix = K*B;
mtl_segment_bin_to_tiles(context, seg);
}
@ -899,7 +930,7 @@ void mtl_cubic_setup(thread mtl_segment_setup_context* context, float2 p[4])
float2 sp[4];
mtl_cubic_slice(p, splits[sliceIndex], splits[sliceIndex+1], sp);
mtl_cubic_info curve = mtl_cubic_classify(sp);
mtl_cubic_info curve = mtl_cubic_classify(sp, context->log);
switch(curve.kind)
{
case MTL_CUBIC_ERROR:
@ -908,13 +939,13 @@ void mtl_cubic_setup(thread mtl_segment_setup_context* context, float2 p[4])
case MTL_CUBIC_DEGENERATE_LINE:
{
float2 l[2] = {p[0], p[1]};
float2 l[2] = {sp[0], sp[1]};
mtl_line_setup(context, l);
} break;
case MTL_CUBIC_DEGENERATE_QUADRATIC:
{
float2 q[3] = {p[0], curve.quadPoint, p[3]};
float2 q[3] = {sp[0], curve.quadPoint, sp[3]};
mtl_quadratic_setup(context, q);
} break;
@ -945,10 +976,6 @@ void mtl_cubic_setup(thread mtl_segment_setup_context* context, float2 p[4])
case MTL_CUBIC_CUSP:
case MTL_CUBIC_SERPENTINE:
{
if(sliceIndex == 2)
{
log_cubic_bezier(sp, context->log);
}
mtl_cubic_emit(context, sp, curve);
} break;
}
@ -973,6 +1000,7 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]],
const device mg_mtl_path_queue* pathQueue = &pathQueueBuffer[elt->pathIndex];
device mg_mtl_tile_queue* tileQueues = &tileQueueBuffer[pathQueue->tileQueues];
mtl_segment_setup_context setupCtx = {.pathIndex = elt->pathIndex,
.segmentCount = segmentCount,
.segmentBuffer = segmentBuffer,
@ -982,7 +1010,8 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]],
.tileOpCount = tileOpCount,
.tileSize = tileSize[0],
.log.buffer = logBuffer,
.log.offset = logOffsetBuffer};
.log.offset = logOffsetBuffer,
.log.enabled = (eltIndex == 1)};
switch(elt->kind)
{
@ -1009,9 +1038,13 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]],
kernel void mtl_backprop(const device mg_mtl_path_queue* pathQueueBuffer [[buffer(0)]],
device mg_mtl_tile_queue* tileQueueBuffer [[buffer(1)]],
device char* logBuffer [[buffer(2)]],
device atomic_int* logOffsetBuffer [[buffer(3)]],
uint pathIndex [[threadgroup_position_in_grid]],
uint localID [[thread_position_in_threadgroup]])
{
// mtl_log_context log = {.buffer = logBuffer, .offset = logOffsetBuffer, .enabled = false};
threadgroup atomic_int nextRowIndex;
if(localID == 0)
{
@ -1048,6 +1081,8 @@ 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)]],
uint2 threadCoord [[thread_position_in_grid]],
uint2 gridSize [[threads_per_grid]])
{
@ -1056,6 +1091,11 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]],
device int* nextLink = &screenTilesBuffer[tileIndex];
*nextLink = -1;
/*
mtl_log_context log = {.buffer = logBuffer,
.offset = logOffsetBuffer,
.enabled = true};
*/
for(int pathIndex = 0; pathIndex < pathCount[0]; pathIndex++)
{
const device mg_mtl_path_queue* pathQueue = &pathQueueBuffer[pathIndex];
@ -1134,13 +1174,6 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
int nTilesX = (int(gridSize.x) + tileSize[0] - 1)/tileSize[0];
int tileIndex = tileCoord.y * nTilesX + tileCoord.x;
if( (pixelCoord.x % tileSize[0] == 0)
||(pixelCoord.y % tileSize[0] == 0))
{
outTexture.write(float4(0, 0, 0, 1), uint2(pixelCoord));
return;
}
float4 color = float4(0, 0, 0, 0);
int pathIndex = 0;
int winding = 0;
@ -1152,12 +1185,15 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
if(op->kind == MG_MTL_OP_START)
{
if(winding & 1)
bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding & 1))
||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding != 0));
if(filled)
{
float4 pathColor = pathBuffer[pathIndex].color;
pathColor.rgb *= pathColor.a;
color = color*(1-pathColor.a) + pathColor;
}
pathIndex = op->index;
winding = op->windingOffset;
}
@ -1165,23 +1201,22 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
{
const device mg_mtl_segment* seg = &segmentBuffer[op->index];
if(pixelCoord.y >= seg->box.y && pixelCoord.y < seg->box.w)
if( (pixelCoord.y > seg->box.y)
&&(pixelCoord.y <= seg->box.w)
&&(mtl_side_of_segment(float2(pixelCoord), seg) < 0))
{
if(mtl_is_left_of_segment(float2(pixelCoord), seg))
{
winding += seg->windingIncrement;
}
winding += seg->windingIncrement;
}
if(op->crossRight)
{
if( (seg->config == MG_MTL_BR || seg->config == MG_MTL_TL)
&&(pixelCoord.y >= seg->box.w))
&&(pixelCoord.y > seg->box.w))
{
winding += seg->windingIncrement;
}
else if( (seg->config == MG_MTL_BL || seg->config == MG_MTL_TR)
&&(pixelCoord.y >= seg->box.y))
&&(pixelCoord.y > seg->box.y))
{
winding -= seg->windingIncrement;
}
@ -1189,13 +1224,23 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
}
opIndex = op->next;
}
if(winding & 1)
bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding & 1))
||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding != 0));
if(filled)
{
float4 pathColor = pathBuffer[pathIndex].color;
pathColor.rgb *= pathColor.a;
color = color*(1-pathColor.a) + pathColor;
}
if( (pixelCoord.x % tileSize[0] == 0)
||(pixelCoord.y % tileSize[0] == 0))
{
outTexture.write(float4(0, 0, 0, 1), uint2(pixelCoord));
return;
}
outTexture.write(color, uint2(pixelCoord));
}