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

[mtl canvas] use same ccw test for testing against diagonal and testing against hull

This commit is contained in:
Martin Fouilleul 2023-04-08 10:04:17 +02:00
parent 2b7aadf180
commit effd8f1dd0
6 changed files with 5017 additions and 4000 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
build.sh
View File

@ -1,6 +1,6 @@
#!/bin/bash #!/bin/bash
DEBUG_FLAGS="-g -O2 -DDEBUG -DLOG_COMPILE_DEBUG" DEBUG_FLAGS="-g -DDEBUG -DLOG_COMPILE_DEBUG"
#DEBUG_FLAGS="-O3" #DEBUG_FLAGS="-O3"
#-------------------------------------------------------------- #--------------------------------------------------------------

52
examples/tiger/main.c
View File

@ -84,8 +84,11 @@ int main()
bool tracked = false; bool tracked = false;
vec2 trackPoint = {0}; vec2 trackPoint = {0};
f32 zoom = 1; f32 zoom = 1;
f32 startX = 300, startY = 200; f32 startX = 300, startY = 200;
bool singlePath = false;
int singlePathIndex = 0;
f64 frameTime = 0; f64 frameTime = 0;
@ -135,6 +138,43 @@ int main()
startY = mousePos.y - pinY*zoom; startY = mousePos.y - pinY*zoom;
} break; } break;
case MP_EVENT_KEYBOARD_KEY:
{
if(event.key.action == MP_KEY_PRESS || event.key.action == MP_KEY_REPEAT)
{
switch(event.key.code)
{
case MP_KEY_SPACE:
singlePath = !singlePath;
break;
case MP_KEY_UP:
{
if(event.key.mods & MP_KEYMOD_SHIFT)
{
singlePathIndex++;
}
else
{
zoom += 0.001;
}
} break;
case MP_KEY_DOWN:
{
if(event.key.mods & MP_KEYMOD_SHIFT)
{
singlePathIndex--;
}
else
{
zoom -= 0.001;
}
} break;
}
}
} break;
default: default:
break; break;
} }
@ -155,10 +195,16 @@ int main()
mg_matrix_push((mg_mat2x3){zoom, 0, startX, mg_matrix_push((mg_mat2x3){zoom, 0, startX,
0, zoom, startY}); 0, zoom, startY});
draw_tiger(); draw_tiger(singlePath, singlePathIndex);
if(singlePath)
{
printf("display single path %i\n", singlePathIndex);
}
printf("viewpos = (%f, %f), zoom = %f\n", startX, startY, zoom);
mg_matrix_pop(); mg_matrix_pop();
/*
// text // text
mg_set_color_rgba(0, 0, 1, 1); mg_set_color_rgba(0, 0, 1, 1);
mg_set_font(font); mg_set_font(font);
@ -171,7 +217,7 @@ int main()
1./frameTime); 1./frameTime);
mg_text_outlines(text); mg_text_outlines(text);
mg_fill(); mg_fill();
*/
printf("Milepost vector graphics test program (frame time = %fs, fps = %f)...\n", printf("Milepost vector graphics test program (frame time = %fs, fps = %f)...\n",
frameTime, frameTime,
1./frameTime); 1./frameTime);

41
examples/tiger/svg2mg.py
View File

@ -54,7 +54,7 @@ class svgContext:
if rel: if rel:
x += self.cp[0] x += self.cp[0]
y += self.cp[1] y += self.cp[1]
print("\tmg_move_to(" + f2s(x) + ", " + f2s(y) + ");") print("\t\tmg_move_to(" + f2s(x) + ", " + f2s(y) + ");")
self.sp = (x, y) self.sp = (x, y)
self.cp = (x, y) self.cp = (x, y)
self.rp = self.cp self.rp = self.cp
@ -67,7 +67,7 @@ class svgContext:
y2 += self.cp[1] y2 += self.cp[1]
x3 += self.cp[0] x3 += self.cp[0]
y3 += self.cp[1] y3 += self.cp[1]
print("\tmg_cubic_to(" + f2s(x1) + ", " + f2s(y1) + ", " + f2s(x2) + ", " + f2s(y2) + ", " + f2s(x3) + ", " + f2s(y3) + ");") print("\t\tmg_cubic_to(" + f2s(x1) + ", " + f2s(y1) + ", " + f2s(x2) + ", " + f2s(y2) + ", " + f2s(x3) + ", " + f2s(y3) + ");")
self.rp = (x2, y2) self.rp = (x2, y2)
self.cp = (x3, y3) self.cp = (x3, y3)
@ -79,7 +79,7 @@ class svgContext:
y3 += self.cp[1] y3 += self.cp[1]
x1 = 2*self.cp[0] - self.rp[0] x1 = 2*self.cp[0] - self.rp[0]
y1 = 2*self.cp[1] - self.rp[1] y1 = 2*self.cp[1] - self.rp[1]
print("\tmg_cubic_to(" + f2s(x1) + ", " + f2s(y1) + ", " + f2s(x2) + ", " + f2s(y2) + ", " + f2s(x3) + ", " + f2s(y3) + ");") print("\t\tmg_cubic_to(" + f2s(x1) + ", " + f2s(y1) + ", " + f2s(x2) + ", " + f2s(y2) + ", " + f2s(x3) + ", " + f2s(y3) + ");")
self.rp = (x2, y2) self.rp = (x2, y2)
self.cp = (x3, y3) self.cp = (x3, y3)
@ -87,7 +87,7 @@ class svgContext:
if rel: if rel:
x1 += self.cp[0] x1 += self.cp[0]
y1 += self.cp[1] y1 += self.cp[1]
print("\tmg_line_to(" + f2s(x1) + ", " + f2s(y1) + ");") print("\t\tmg_line_to(" + f2s(x1) + ", " + f2s(y1) + ");")
self.cp = (x1, y1) self.cp = (x1, y1)
self.rp = self.cp self.rp = self.cp
@ -95,12 +95,12 @@ class svgContext:
if rel: if rel:
y1 += self.cp[1] y1 += self.cp[1]
x1 = self.cp[0] x1 = self.cp[0]
print("\tmg_line_to(" + f2s(x1) + ", " + f2s(y1) + ");") print("\t\tmg_line_to(" + f2s(x1) + ", " + f2s(y1) + ");")
self.cp = (x1, y1) self.cp = (x1, y1)
self.rp = self.cp self.rp = self.cp
def close_path(self): def close_path(self):
print("\tmg_close_path();"); print("\t\tmg_close_path();");
self.cp = self.rp = self.sp self.cp = self.rp = self.sp
def print_path(path, ctx): def print_path(path, ctx):
@ -190,13 +190,17 @@ def parse_color(s):
tree = et.parse('./Ghostscript_Tiger.svg') tree = et.parse('./Ghostscript_Tiger.svg')
ctx = svgContext() ctx = svgContext()
print("void draw_tiger()") print("void draw_tiger(bool singlePath, int singlePathIndex)")
print("{") print("{")
for g in tree.iter('{http://www.w3.org/2000/svg}g'):
pathIndex = 0;
for g in tree.iter('{http://www.w3.org/2000/svg}g'):
for path in g.findall('{http://www.w3.org/2000/svg}path'): for path in g.findall('{http://www.w3.org/2000/svg}path'):
ctx.reset() ctx.reset()
print("\tif(!singlePath || singlePathIndex == "+ str(pathIndex) +")")
print("\t{")
print_path(path, ctx) print_path(path, ctx)
fill = g.get('fill') fill = g.get('fill')
@ -205,22 +209,31 @@ for g in tree.iter('{http://www.w3.org/2000/svg}g'):
if fill != None and fill != "none": if fill != None and fill != "none":
(r, g, b) = parse_color(fill) (r, g, b) = parse_color(fill)
print("\tmg_set_color_rgba(" + f2s(r) + ", " + f2s(g) + ", " + f2s(b) + ", 1);") print("\t\tmg_set_color_rgba(" + f2s(r) + ", " + f2s(g) + ", " + f2s(b) + ", 1);")
print("\tmg_fill();") print("\t\tmg_fill();")
print("\t}")
pathIndex += 1
if stroke != None and stroke != "none": if stroke != None and stroke != "none":
print("\tif(!singlePath || singlePathIndex == "+ str(pathIndex) +")")
print("\t{")
if stroke_width != None: if stroke_width != None:
print("\tmg_set_width(" + stroke_width + ");"); print("\t\tmg_set_width(" + stroke_width + ");");
else: else:
print("\tmg_set_width(1);"); print("\t\tmg_set_width(1);");
(r, g, b) = parse_color(stroke) (r, g, b) = parse_color(stroke)
if fill != None: if fill != None:
ctx.reset() ctx.reset()
print_path(path, ctx) print_path(path, ctx)
print("\tmg_set_color_rgba(" + f2s(r) + ", " + f2s(g) + ", " + f2s(b) + ", 1);") print("\t\tmg_set_color_rgba(" + f2s(r) + ", " + f2s(g) + ", " + f2s(b) + ", 1);")
print("\tmg_stroke();") print("\t\tmg_stroke();")
print("\t}")
pathIndex += 1
if (stroke == None or stroke == 'none') and (fill == None or fill == ''): if (stroke == None or stroke == 'none') and (fill == None or fill == ''):
print("error, group " + g.get("id") + " has no command") print("error, group " + g.get("id") + " has no command")

956
examples/tiger/tiger.c
View File

File diff suppressed because it is too large Load Diff

9
src/mtl_renderer.m
View File

@ -19,7 +19,8 @@
#define LOG_SUBSYSTEM "Graphics" #define LOG_SUBSYSTEM "Graphics"
const int MG_MTL_INPUT_BUFFERS_COUNT = 3, const int MG_MTL_INPUT_BUFFERS_COUNT = 3,
MG_MTL_TILE_SIZE = 16; MG_MTL_TILE_SIZE = 16,
MG_MTL_MSAA_COUNT = 1;
typedef struct mg_mtl_canvas_backend typedef struct mg_mtl_canvas_backend
{ {
@ -52,6 +53,8 @@ typedef struct mg_mtl_canvas_backend
id<MTLBuffer> tileOpCountBuffer; id<MTLBuffer> tileOpCountBuffer;
id<MTLBuffer> screenTilesBuffer; id<MTLBuffer> screenTilesBuffer;
int msaaCount;
} mg_mtl_canvas_backend; } mg_mtl_canvas_backend;
@ -807,6 +810,9 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
[rasterEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:0 atIndex:2]; [rasterEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:0 atIndex:2];
[rasterEncoder setBuffer:backend->segmentBuffer offset:0 atIndex:3]; [rasterEncoder setBuffer:backend->segmentBuffer offset:0 atIndex:3];
[rasterEncoder setBytes:&tileSize length:sizeof(int) atIndex:4]; [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 setTexture:backend->outTexture atIndex:0]; [rasterEncoder setTexture:backend->outTexture atIndex:0];
@ -900,6 +906,7 @@ mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface)
backend = malloc_type(mg_mtl_canvas_backend); backend = malloc_type(mg_mtl_canvas_backend);
memset(backend, 0, sizeof(mg_mtl_canvas_backend)); memset(backend, 0, sizeof(mg_mtl_canvas_backend));
backend->msaaCount = MG_MTL_MSAA_COUNT;
backend->surface = surface; backend->surface = surface;
//NOTE(martin): setup interface functions //NOTE(martin): setup interface functions

135
src/mtl_renderer.metal
View File

@ -290,40 +290,6 @@ int mtl_side_of_segment(float2 p, const device mg_mtl_segment* seg, mtl_log_cont
} }
else 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;
float dy = p.y - seg->box.y;
if( (seg->config == MG_MTL_BR && dy >= alpha*dx)
||(seg->config == MG_MTL_TR && dy <= ofs - alpha*dx))
{
side = -1;
}
else if( (seg->config == MG_MTL_TL && dy < alpha*dx)
||(seg->config == MG_MTL_BL && dy > ofs - alpha*dx))
{
side = 1;
}
else
{
switch(seg->kind)
{
case MG_MTL_LINE:
side = 1;
break;
case MG_MTL_QUADRATIC:
{
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:
{
/*
float2 a, b, c; float2 a, b, c;
switch(seg->config) switch(seg->config)
{ {
@ -349,29 +315,37 @@ int mtl_side_of_segment(float2 p, const device mg_mtl_segment* seg, mtl_log_cont
} }
c = seg->hullVertex; c = seg->hullVertex;
if(ccw(b, c, p) > 0 && ccw(c, a, p) > 0) if(ccw(a, b, p) < 0)
{ {
float3 ph = {p.x, p.y, 1}; // other side of the diagonal
float3 klm = seg->implicitMatrix * ph; side = (seg->config == MG_MTL_BR || seg->config == MG_MTL_TR) ? -1 : 1;
side = (seg->sign*(klm.x*klm.x*klm.x - klm.y*klm.z) < 0)? -1 : 1; }
else if(ccw(b, c, p) < 0 || ccw(c, a, p) < 0)
{
// same side of the diagonal, but outside curve hull
side = (seg->config == MG_MTL_BL || seg->config == MG_MTL_TL) ? -1 : 1;
} }
else else
{ {
side = (seg->config == MG_MTL_BL || seg->config == MG_MTL_TL) ? -1 : 1; // inside curve hull
} switch(seg->kind)
*/
float3 ph = {p.x, p.y, 1};
float3 hullCoords = seg->hullMatrix * ph;
if(all(hullCoords > 0))
{ {
float3 klm = seg->implicitMatrix * ph; case MG_MTL_LINE:
side = (seg->sign*(klm.x*klm.x*klm.x - klm.y*klm.z) < 0)? -1 : 1; side = 1;
} break;
else
{
side = (seg->config == MG_MTL_BL || seg->config == MG_MTL_TL) ? -1 : 1;
}
case MG_MTL_QUADRATIC:
{
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:
{
float3 ph = {p.x, p.y, 1};
float3 klm = seg->implicitMatrix * ph;
side = (seg->sign*(klm.x*klm.x*klm.x - klm.y*klm.z) < 0)? -1 : 1;
} break; } break;
} }
} }
@ -584,6 +558,7 @@ device mg_mtl_segment* mtl_segment_push(thread mtl_segment_setup_context* contex
void mtl_line_setup(thread mtl_segment_setup_context* context, float2 p[2]) void mtl_line_setup(thread mtl_segment_setup_context* context, float2 p[2])
{ {
device mg_mtl_segment* seg = mtl_segment_push(context, p, MG_MTL_LINE); device mg_mtl_segment* seg = mtl_segment_push(context, p, MG_MTL_LINE);
seg->hullVertex = p[0];
mtl_segment_bin_to_tiles(context, seg); mtl_segment_bin_to_tiles(context, seg);
} }
@ -669,6 +644,8 @@ void mtl_quadratic_emit(thread mtl_segment_setup_context* context,
{b, e, 0.}, {b, e, 0.},
{c, f, g}); {c, f, g});
seg->hullVertex = p[1];
mtl_segment_bin_to_tiles(context, seg); mtl_segment_bin_to_tiles(context, seg);
} }
@ -1483,10 +1460,18 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
const device mg_mtl_path* pathBuffer [[buffer(2)]], const device mg_mtl_path* pathBuffer [[buffer(2)]],
const device mg_mtl_segment* segmentBuffer [[buffer(3)]], const device mg_mtl_segment* segmentBuffer [[buffer(3)]],
constant int* tileSize [[buffer(4)]], constant int* tileSize [[buffer(4)]],
constant int* sampleCountBuffer [[buffer(5)]],
device char* logBuffer [[buffer(6)]],
device atomic_int* logOffsetBuffer [[buffer(7)]],
texture2d<float, access::write> outTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(0)]],
uint2 threadCoord [[thread_position_in_grid]], uint2 threadCoord [[thread_position_in_grid]],
uint2 gridSize [[threads_per_grid]]) uint2 gridSize [[threads_per_grid]])
{ {
/*
mtl_log_context log = {.buffer = logBuffer,
.offset = logOffsetBuffer,
.enabled = true};
*/
float2 pixelCoord = float2(threadCoord); float2 pixelCoord = float2(threadCoord);
int2 tileCoord = int2(threadCoord) / tileSize[0]; int2 tileCoord = int2(threadCoord) / tileSize[0];
int nTilesX = (int(gridSize.x) + tileSize[0] - 1)/tileSize[0]; int nTilesX = (int(gridSize.x) + tileSize[0] - 1)/tileSize[0];
@ -1495,18 +1480,30 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
int pathIndex = 0; int pathIndex = 0;
int opIndex = screenTilesBuffer[tileIndex]; int opIndex = screenTilesBuffer[tileIndex];
const int MG_MTL_SAMPLE_COUNT = 8; const int MG_MTL_MAX_SAMPLE_COUNT = 8;
const float2 sampleCoords[MG_MTL_SAMPLE_COUNT] = {pixelCoord + float2(1, 3)/16, float2 sampleCoords[MG_MTL_MAX_SAMPLE_COUNT];
pixelCoord + float2(-1, -3)/16, int sampleCount = sampleCountBuffer[0];
pixelCoord + float2(5, -1)/16,
pixelCoord + float2(-3, 5)/16,
pixelCoord + float2(-5, -5)/16,
pixelCoord + float2(-7, 1)/16,
pixelCoord + float2(3, -7)/16,
pixelCoord + float2(7, 7)/16};
float4 color[MG_MTL_SAMPLE_COUNT] = {0}; if(sampleCount == 8)
int winding[MG_MTL_SAMPLE_COUNT] = {0}; {
sampleCount = 8;
sampleCoords[0] = pixelCoord + float2(1, 3)/16;
sampleCoords[1] = pixelCoord + float2(-1, -3)/16;
sampleCoords[2] = pixelCoord + float2(5, -1)/16;
sampleCoords[3] = pixelCoord + float2(-3, 5)/16;
sampleCoords[4] = pixelCoord + float2(-5, -5)/16;
sampleCoords[5] = pixelCoord + float2(-7, 1)/16;
sampleCoords[6] = pixelCoord + float2(3, -7)/16;
sampleCoords[7] = pixelCoord + float2(7, 7)/16;
}
else
{
sampleCount = 1;
sampleCoords[0] = pixelCoord;
}
float4 color[MG_MTL_MAX_SAMPLE_COUNT] = {0};
int winding[MG_MTL_MAX_SAMPLE_COUNT] = {0};
while(opIndex != -1) while(opIndex != -1)
{ {
@ -1517,7 +1514,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
float4 pathColor = pathBuffer[pathIndex].color; float4 pathColor = pathBuffer[pathIndex].color;
pathColor.rgb *= pathColor.a; pathColor.rgb *= pathColor.a;
for(int sampleIndex=0; sampleIndex<MG_MTL_SAMPLE_COUNT; sampleIndex++) for(int sampleIndex=0; sampleIndex<sampleCount; sampleIndex++)
{ {
bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1)) bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1))
||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0)); ||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0));
@ -1533,7 +1530,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
{ {
const device mg_mtl_segment* seg = &segmentBuffer[op->index]; const device mg_mtl_segment* seg = &segmentBuffer[op->index];
for(int sampleIndex=0; sampleIndex<MG_MTL_SAMPLE_COUNT; sampleIndex++) for(int sampleIndex=0; sampleIndex<sampleCount; sampleIndex++)
{ {
float2 sampleCoord = sampleCoords[sampleIndex]; float2 sampleCoord = sampleCoords[sampleIndex];
@ -1566,7 +1563,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
float4 pathColor = pathBuffer[pathIndex].color; float4 pathColor = pathBuffer[pathIndex].color;
pathColor.rgb *= pathColor.a; pathColor.rgb *= pathColor.a;
for(int sampleIndex=0; sampleIndex<MG_MTL_SAMPLE_COUNT; sampleIndex++) for(int sampleIndex=0; sampleIndex<sampleCount; sampleIndex++)
{ {
bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1)) bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1))
||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0)); ||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0));
@ -1576,16 +1573,16 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
} }
pixelColor += color[sampleIndex]; pixelColor += color[sampleIndex];
} }
pixelColor /= MG_MTL_SAMPLE_COUNT; pixelColor /= sampleCount;
/* /*
if( (pixelCoord.x % tileSize[0] == 0) if( (int(pixelCoord.x) % tileSize[0] == 0)
||(pixelCoord.y % tileSize[0] == 0)) ||(int(pixelCoord.y) % tileSize[0] == 0))
{ {
outTexture.write(float4(0, 0, 0, 1), uint2(pixelCoord)); outTexture.write(float4(0, 0, 0, 1), uint2(pixelCoord));
return; return;
} }
*/ //*/
outTexture.write(pixelColor, uint2(pixelCoord)); outTexture.write(pixelColor, uint2(pixelCoord));
} }