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

Merge branch 'orca_ui' of git.handmade.network:ilidemi/milepost into orca_ui

This commit is contained in:
Ilia Demianenko 2023-07-15 00:17:53 -07:00
parent 40906bb852 ff1460e629
commit 445325d45f
10 changed files with 4177 additions and 3428 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
scripts/glapi.py
View File

@ -44,10 +44,11 @@ tree = et.parse(args.spec)
gl41 = gather_api(tree, 'gl', 4.1) gl41 = gather_api(tree, 'gl', 4.1)
gl43 = gather_api(tree, 'gl', 4.3) gl43 = gather_api(tree, 'gl', 4.3)
gl44 = gather_api(tree, 'gl', 4.4)
gles30 = gather_api(tree, 'gles2', 3.1) gles30 = gather_api(tree, 'gles2', 3.1)
gles31 = gather_api(tree, 'gles2', 3.2) gles31 = gather_api(tree, 'gles2', 3.2)
glall = list(set().union(gl41, gl43, gles30, gles31)) glall = list(set().union(gl41, gl43, gl44, gles30, gles31))
#--------------------------------------------------------------- #---------------------------------------------------------------
@ -122,6 +123,7 @@ f.write("typedef void*(*mg_gl_load_proc)(const char* name);\n\n")
f.write("void mg_gl_load_gl41(mg_gl_api* api, mg_gl_load_proc loadProc);\n") f.write("void mg_gl_load_gl41(mg_gl_api* api, mg_gl_load_proc loadProc);\n")
f.write("void mg_gl_load_gl43(mg_gl_api* api, mg_gl_load_proc loadProc);\n") f.write("void mg_gl_load_gl43(mg_gl_api* api, mg_gl_load_proc loadProc);\n")
f.write("void mg_gl_load_gl44(mg_gl_api* api, mg_gl_load_proc loadProc);\n")
f.write("void mg_gl_load_gles30(mg_gl_api* api, mg_gl_load_proc loadProc);\n") f.write("void mg_gl_load_gles30(mg_gl_api* api, mg_gl_load_proc loadProc);\n")
f.write("void mg_gl_load_gles31(mg_gl_api* api, mg_gl_load_proc loadProc);\n\n") f.write("void mg_gl_load_gles31(mg_gl_api* api, mg_gl_load_proc loadProc);\n\n")
@ -152,6 +154,7 @@ f.write("mp_thread_local mg_gl_api* __mgGLAPI = 0;\n\n")
emit_loader(f, 'gl41', gl41) emit_loader(f, 'gl41', gl41)
emit_loader(f, 'gl43', gl43) emit_loader(f, 'gl43', gl43)
emit_loader(f, 'gl44', gl44)
emit_loader(f, 'gles30', gles30) emit_loader(f, 'gles30', gles30)
emit_loader(f, 'gles31', gles31) emit_loader(f, 'gles31', gles31)

2160
src/gl_api.h
View File

File diff suppressed because it is too large Load Diff

442
src/gl_canvas.c
View File

@ -98,24 +98,19 @@ typedef struct mg_gl_tile_queue
//////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////
typedef struct mg_gl_encoding_context
{
int glEltCount;
mg_gl_path* pathBufferData;
mg_gl_path_elt* elementBufferData;
int pathIndex;
mg_primitive* primitive;
vec4 pathScreenExtents;
vec4 pathUserExtents;
} mg_gl_encoding_context;
enum { enum {
// MG_GL_INPUT_BUFFERS_COUNT = 3, MG_GL_INPUT_BUFFERS_COUNT = 3,
MG_GL_TILE_SIZE = 16, MG_GL_TILE_SIZE = 16,
MG_GL_MSAA_COUNT = 8, MG_GL_MSAA_COUNT = 8,
}; };
typedef struct mg_gl_mapped_buffer
{
GLuint buffer;
int size;
char* contents;
} mg_gl_mapped_buffer;
typedef struct mg_gl_canvas_backend typedef struct mg_gl_canvas_backend
{ {
mg_canvas_backend interface; mg_canvas_backend interface;
@ -133,10 +128,9 @@ typedef struct mg_gl_canvas_backend
GLuint outTexture; GLuint outTexture;
int bufferIndex; int bufferIndex;
//TODO buffer semaphore... GLsync bufferSync[MG_GL_INPUT_BUFFERS_COUNT];
mg_gl_mapped_buffer pathBuffer[MG_GL_INPUT_BUFFERS_COUNT];
GLuint pathBuffer; mg_gl_mapped_buffer elementBuffer[MG_GL_INPUT_BUFFERS_COUNT];
GLuint elementBuffer;
GLuint segmentBuffer; GLuint segmentBuffer;
GLuint segmentCountBuffer; GLuint segmentCountBuffer;
@ -149,12 +143,21 @@ typedef struct mg_gl_canvas_backend
GLuint dummyVertexBuffer; GLuint dummyVertexBuffer;
mg_gl_path* pathBufferData;
mg_gl_path_elt* elementBufferData;
int msaaCount; int msaaCount;
vec2 frameSize; vec2 frameSize;
//encoding context
int pathCount;
int eltCount;
/////////////////
int pathBatchStart;
int eltBatchStart;
mg_primitive* primitive;
vec4 pathScreenExtents;
vec4 pathUserExtents;
} mg_gl_canvas_backend; } mg_gl_canvas_backend;
static void mg_update_path_extents(vec4* extents, vec2 p) static void mg_update_path_extents(vec4* extents, vec2 p)
@ -165,27 +168,28 @@ static void mg_update_path_extents(vec4* extents, vec2 p)
extents->w = maximum(extents->w, p.y); extents->w = maximum(extents->w, p.y);
} }
void mg_gl_canvas_encode_element(mg_gl_encoding_context* context, mg_path_elt_type kind, vec2* p) void mg_gl_canvas_encode_element(mg_gl_canvas_backend* backend, mg_path_elt_type kind, vec2* p)
{ {
mg_gl_path_elt* glElt = &context->elementBufferData[context->glEltCount]; mg_gl_path_elt* elementData = (mg_gl_path_elt*)backend->elementBuffer[backend->bufferIndex].contents;
context->glEltCount++; mg_gl_path_elt* elt = &elementData[backend->eltCount];
backend->eltCount++;
glElt->pathIndex = context->pathIndex; elt->pathIndex = backend->pathCount - backend->pathBatchStart;
int count = 0; int count = 0;
switch(kind) switch(kind)
{ {
case MG_PATH_LINE: case MG_PATH_LINE:
glElt->kind = MG_GL_LINE; elt->kind = MG_GL_LINE;
count = 2; count = 2;
break; break;
case MG_PATH_QUADRATIC: case MG_PATH_QUADRATIC:
glElt->kind = MG_GL_QUADRATIC; elt->kind = MG_GL_QUADRATIC;
count = 3; count = 3;
break; break;
case MG_PATH_CUBIC: case MG_PATH_CUBIC:
glElt->kind = MG_GL_CUBIC; elt->kind = MG_GL_CUBIC;
count = 4; count = 4;
break; break;
@ -195,12 +199,85 @@ void mg_gl_canvas_encode_element(mg_gl_encoding_context* context, mg_path_elt_ty
for(int i=0; i<count; i++) for(int i=0; i<count; i++)
{ {
mg_update_path_extents(&context->pathUserExtents, p[i]); mg_update_path_extents(&backend->pathUserExtents, p[i]);
vec2 screenP = mg_mat2x3_mul(context->primitive->attributes.transform, p[i]); vec2 screenP = mg_mat2x3_mul(backend->primitive->attributes.transform, p[i]);
glElt->p[i] = (vec2){screenP.x, screenP.y}; elt->p[i] = (vec2){screenP.x, screenP.y};
mg_update_path_extents(&context->pathScreenExtents, screenP); mg_update_path_extents(&backend->pathScreenExtents, screenP);
}
}
void mg_gl_canvas_encode_path(mg_gl_canvas_backend* backend, mg_primitive* primitive, f32 scale)
{
mg_gl_path* pathData = (mg_gl_path*)backend->pathBuffer[backend->bufferIndex].contents;
mg_gl_path* path = &pathData[backend->pathCount];
backend->pathCount++;
path->cmd = (mg_gl_cmd)primitive->cmd;
path->box = (vec4){
backend->pathScreenExtents.x,
backend->pathScreenExtents.y,
backend->pathScreenExtents.z,
backend->pathScreenExtents.w};
path->clip = (vec4){
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 = (vec4){
primitive->attributes.color.r,
primitive->attributes.color.g,
primitive->attributes.color.b,
primitive->attributes.color.a};
mp_rect srcRegion = primitive->attributes.srcRegion;
mp_rect destRegion = {
backend->pathUserExtents.x,
backend->pathUserExtents.y,
backend->pathUserExtents.z - backend->pathUserExtents.x,
backend->pathUserExtents.w - backend->pathUserExtents.y};
if(!mg_image_is_nil(primitive->attributes.image))
{
vec2 texSize = mg_image_size(primitive->attributes.image);
mg_mat2x3 srcRegionToImage = {
1/texSize.x, 0, srcRegion.x/texSize.x,
0, 1/texSize.y, srcRegion.y/texSize.y};
mg_mat2x3 destRegionToSrcRegion = {
srcRegion.w/destRegion.w, 0, 0,
0, srcRegion.h/destRegion.h, 0};
mg_mat2x3 userToDestRegion = {
1, 0, -destRegion.x,
0, 1, -destRegion.y};
mg_mat2x3 screenToUser = mg_mat2x3_inv(primitive->attributes.transform);
mg_mat2x3 uvTransform = srcRegionToImage;
uvTransform = mg_mat2x3_mul_m(uvTransform, destRegionToSrcRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, userToDestRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, screenToUser);
//NOTE: mat3 std430 layout is an array of vec3, which are padded to _vec4_ alignment
path->uvTransform[0] = uvTransform.m[0]/scale;
path->uvTransform[1] = uvTransform.m[3]/scale;
path->uvTransform[2] = 0;
path->uvTransform[3] = 0;
path->uvTransform[4] = uvTransform.m[1]/scale;
path->uvTransform[5] = uvTransform.m[4]/scale;
path->uvTransform[6] = 0;
path->uvTransform[7] = 0;
path->uvTransform[8] = uvTransform.m[2];
path->uvTransform[9] = uvTransform.m[5];
path->uvTransform[10] = 1;
path->uvTransform[11] = 0;
} }
} }
@ -392,9 +469,9 @@ void mg_cubic_split(vec2 p[4], f32 t, vec2 outLeft[4], vec2 outRight[4])
outRight[3] = p[3]; outRight[3] = p[3];
} }
void mg_gl_encode_stroke_line(mg_gl_encoding_context* context, vec2* p) void mg_gl_encode_stroke_line(mg_gl_canvas_backend* backend, vec2* p)
{ {
f32 width = context->primitive->attributes.width; f32 width = backend->primitive->attributes.width;
vec2 v = {p[1].x-p[0].x, p[1].y-p[0].y}; vec2 v = {p[1].x-p[0].x, p[1].y-p[0].y};
vec2 n = {v.y, -v.x}; vec2 n = {v.y, -v.x};
@ -406,30 +483,30 @@ void mg_gl_encode_stroke_line(mg_gl_encoding_context* context, vec2* p)
vec2 joint0[2] = {vec2_add(p[0], vec2_mul(-1, offset)), vec2_add(p[0], 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))}; vec2 joint1[2] = {vec2_add(p[1], offset), vec2_add(p[1], vec2_mul(-1, offset))};
mg_gl_canvas_encode_element(context, MG_PATH_LINE, right); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, right);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, left); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, left);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, joint0); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, joint0);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, joint1); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, joint1);
} }
enum { MG_HULL_CHECK_SAMPLE_COUNT = 5 }; enum { MG_HULL_CHECK_SAMPLE_COUNT = 5 };
void mg_gl_encode_stroke_quadratic(mg_gl_encoding_context* context, vec2* p) void mg_gl_encode_stroke_quadratic(mg_gl_canvas_backend* backend, vec2* p)
{ {
f32 width = context->primitive->attributes.width; f32 width = backend->primitive->attributes.width;
f32 tolerance = minimum(context->primitive->attributes.tolerance, 0.5 * width); f32 tolerance = minimum(backend->primitive->attributes.tolerance, 0.5 * width);
//NOTE: check for degenerate line case //NOTE: check for degenerate line case
const f32 equalEps = 1e-3; const f32 equalEps = 1e-3;
if(vec2_close(p[0], p[1], equalEps)) if(vec2_close(p[0], p[1], equalEps))
{ {
mg_gl_encode_stroke_line(context, p+1); mg_gl_encode_stroke_line(backend, p+1);
return; return;
} }
else if(vec2_close(p[1], p[2], equalEps)) else if(vec2_close(p[1], p[2], equalEps))
{ {
mg_gl_encode_stroke_line(context, p); mg_gl_encode_stroke_line(backend, p);
return; return;
} }
@ -444,8 +521,8 @@ void mg_gl_encode_stroke_quadratic(mg_gl_encoding_context* context, vec2* p)
vec2 splitLeft[3]; vec2 splitLeft[3];
vec2 splitRight[3]; vec2 splitRight[3];
mg_quadratic_split(p, 0.5, splitLeft, splitRight); mg_quadratic_split(p, 0.5, splitLeft, splitRight);
mg_gl_encode_stroke_quadratic(context, splitLeft); mg_gl_encode_stroke_quadratic(backend, splitLeft);
mg_gl_encode_stroke_quadratic(context, splitRight); mg_gl_encode_stroke_quadratic(backend, splitRight);
} }
else else
{ {
@ -485,8 +562,8 @@ void mg_gl_encode_stroke_quadratic(mg_gl_encoding_context* context, vec2* p)
vec2 splitLeft[3]; vec2 splitLeft[3];
vec2 splitRight[3]; vec2 splitRight[3];
mg_quadratic_split(p, maxOvershootParameter, splitLeft, splitRight); mg_quadratic_split(p, maxOvershootParameter, splitLeft, splitRight);
mg_gl_encode_stroke_quadratic(context, splitLeft); mg_gl_encode_stroke_quadratic(backend, splitLeft);
mg_gl_encode_stroke_quadratic(context, splitRight); mg_gl_encode_stroke_quadratic(backend, splitRight);
} }
else else
{ {
@ -494,21 +571,21 @@ void mg_gl_encode_stroke_quadratic(mg_gl_encoding_context* context, vec2* p)
leftHull[0] = leftHull[2]; leftHull[0] = leftHull[2];
leftHull[2] = tmp; leftHull[2] = tmp;
mg_gl_canvas_encode_element(context, MG_PATH_QUADRATIC, rightHull); mg_gl_canvas_encode_element(backend, MG_PATH_QUADRATIC, rightHull);
mg_gl_canvas_encode_element(context, MG_PATH_QUADRATIC, leftHull); mg_gl_canvas_encode_element(backend, MG_PATH_QUADRATIC, leftHull);
vec2 joint0[2] = {rightHull[2], leftHull[0]}; vec2 joint0[2] = {rightHull[2], leftHull[0]};
vec2 joint1[2] = {leftHull[2], rightHull[0]}; vec2 joint1[2] = {leftHull[2], rightHull[0]};
mg_gl_canvas_encode_element(context, MG_PATH_LINE, joint0); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, joint0);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, joint1); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, joint1);
} }
} }
} }
void mg_gl_encode_stroke_cubic(mg_gl_encoding_context* context, vec2* p) void mg_gl_encode_stroke_cubic(mg_gl_canvas_backend* backend, vec2* p)
{ {
f32 width = context->primitive->attributes.width; f32 width = backend->primitive->attributes.width;
f32 tolerance = minimum(context->primitive->attributes.tolerance, 0.5 * width); f32 tolerance = minimum(backend->primitive->attributes.tolerance, 0.5 * width);
//NOTE: check degenerate line cases //NOTE: check degenerate line cases
f32 equalEps = 1e-3; f32 equalEps = 1e-3;
@ -518,19 +595,19 @@ void mg_gl_encode_stroke_cubic(mg_gl_encoding_context* context, vec2* p)
||(vec2_close(p[1], p[2], equalEps) && vec2_close(p[2], p[3], equalEps))) ||(vec2_close(p[1], p[2], equalEps) && vec2_close(p[2], p[3], equalEps)))
{ {
vec2 line[2] = {p[0], p[3]}; vec2 line[2] = {p[0], p[3]};
mg_gl_encode_stroke_line(context, line); mg_gl_encode_stroke_line(backend, line);
return; return;
} }
else if(vec2_close(p[0], p[1], equalEps) && vec2_close(p[1], p[3], equalEps)) else if(vec2_close(p[0], p[1], equalEps) && vec2_close(p[1], p[3], equalEps))
{ {
vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[2]))}; vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[2]))};
mg_gl_encode_stroke_line(context, line); mg_gl_encode_stroke_line(backend, line);
return; return;
} }
else if(vec2_close(p[0], p[2], equalEps) && vec2_close(p[2], p[3], equalEps)) else if(vec2_close(p[0], p[2], equalEps) && vec2_close(p[2], p[3], equalEps))
{ {
vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[1]))}; vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[1]))};
mg_gl_encode_stroke_line(context, line); mg_gl_encode_stroke_line(backend, line);
return; return;
} }
@ -545,8 +622,8 @@ void mg_gl_encode_stroke_cubic(mg_gl_encoding_context* context, vec2* p)
vec2 splitLeft[4]; vec2 splitLeft[4];
vec2 splitRight[4]; vec2 splitRight[4];
mg_cubic_split(p, 0.5, splitLeft, splitRight); mg_cubic_split(p, 0.5, splitLeft, splitRight);
mg_gl_encode_stroke_cubic(context, splitLeft); mg_gl_encode_stroke_cubic(backend, splitLeft);
mg_gl_encode_stroke_cubic(context, splitRight); mg_gl_encode_stroke_cubic(backend, splitRight);
} }
else else
{ {
@ -586,8 +663,8 @@ void mg_gl_encode_stroke_cubic(mg_gl_encoding_context* context, vec2* p)
vec2 splitLeft[4]; vec2 splitLeft[4];
vec2 splitRight[4]; vec2 splitRight[4];
mg_cubic_split(p, maxOvershootParameter, splitLeft, splitRight); mg_cubic_split(p, maxOvershootParameter, splitLeft, splitRight);
mg_gl_encode_stroke_cubic(context, splitLeft); mg_gl_encode_stroke_cubic(backend, splitLeft);
mg_gl_encode_stroke_cubic(context, splitRight); mg_gl_encode_stroke_cubic(backend, splitRight);
} }
else else
{ {
@ -598,18 +675,18 @@ void mg_gl_encode_stroke_cubic(mg_gl_encoding_context* context, vec2* p)
leftHull[1] = leftHull[2]; leftHull[1] = leftHull[2];
leftHull[2] = tmp; leftHull[2] = tmp;
mg_gl_canvas_encode_element(context, MG_PATH_CUBIC, rightHull); mg_gl_canvas_encode_element(backend, MG_PATH_CUBIC, rightHull);
mg_gl_canvas_encode_element(context, MG_PATH_CUBIC, leftHull); mg_gl_canvas_encode_element(backend, MG_PATH_CUBIC, leftHull);
vec2 joint0[2] = {rightHull[3], leftHull[0]}; vec2 joint0[2] = {rightHull[3], leftHull[0]};
vec2 joint1[2] = {leftHull[3], rightHull[0]}; vec2 joint1[2] = {leftHull[3], rightHull[0]};
mg_gl_canvas_encode_element(context, MG_PATH_LINE, joint0); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, joint0);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, joint1); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, joint1);
} }
} }
} }
void mg_gl_encode_stroke_element(mg_gl_encoding_context* context, void mg_gl_encode_stroke_element(mg_gl_canvas_backend* backend,
mg_path_elt* element, mg_path_elt* element,
vec2 currentPoint, vec2 currentPoint,
vec2* startTangent, vec2* startTangent,
@ -622,17 +699,17 @@ void mg_gl_encode_stroke_element(mg_gl_encoding_context* context,
switch(element->type) switch(element->type)
{ {
case MG_PATH_LINE: case MG_PATH_LINE:
mg_gl_encode_stroke_line(context, controlPoints); mg_gl_encode_stroke_line(backend, controlPoints);
endPointIndex = 1; endPointIndex = 1;
break; break;
case MG_PATH_QUADRATIC: case MG_PATH_QUADRATIC:
mg_gl_encode_stroke_quadratic(context, controlPoints); mg_gl_encode_stroke_quadratic(backend, controlPoints);
endPointIndex = 2; endPointIndex = 2;
break; break;
case MG_PATH_CUBIC: case MG_PATH_CUBIC:
mg_gl_encode_stroke_cubic(context, controlPoints); mg_gl_encode_stroke_cubic(backend, controlPoints);
endPointIndex = 3; endPointIndex = 3;
break; break;
@ -669,11 +746,11 @@ void mg_gl_encode_stroke_element(mg_gl_encoding_context* context,
DEBUG_ASSERT(startTangent->x != 0 || startTangent->y != 0); DEBUG_ASSERT(startTangent->x != 0 || startTangent->y != 0);
} }
void mg_gl_stroke_cap(mg_gl_encoding_context* context, void mg_gl_stroke_cap(mg_gl_canvas_backend* backend,
vec2 p0, vec2 p0,
vec2 direction) vec2 direction)
{ {
mg_attributes* attributes = &context->primitive->attributes; mg_attributes* attributes = &backend->primitive->attributes;
//NOTE(martin): compute the tangent and normal vectors (multiplied by half width) at the cap point //NOTE(martin): compute the tangent and normal vectors (multiplied by half width) at the cap point
f32 dn = sqrt(Square(direction.x) + Square(direction.y)); f32 dn = sqrt(Square(direction.x) + Square(direction.y));
@ -691,18 +768,18 @@ void mg_gl_stroke_cap(mg_gl_encoding_context* context,
{p0.x - n0.x, p0.y - n0.y}, {p0.x - n0.x, p0.y - n0.y},
{p0.x + n0.x, p0.y + n0.y}}; {p0.x + n0.x, p0.y + n0.y}};
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+1); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+1);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+2); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+2);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+3); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+3);
} }
void mg_gl_stroke_joint(mg_gl_encoding_context* context, void mg_gl_stroke_joint(mg_gl_canvas_backend* backend,
vec2 p0, vec2 p0,
vec2 t0, vec2 t0,
vec2 t1) vec2 t1)
{ {
mg_attributes* attributes = &context->primitive->attributes; mg_attributes* attributes = &backend->primitive->attributes;
//NOTE(martin): compute the normals at the joint point //NOTE(martin): compute the normals at the joint point
f32 norm_t0 = sqrt(Square(t0.x) + Square(t0.y)); f32 norm_t0 = sqrt(Square(t0.x) + Square(t0.y));
@ -750,10 +827,10 @@ void mg_gl_stroke_joint(mg_gl_encoding_context* context,
{p0.x + n1.x*halfW, p0.y + n1.y*halfW}, {p0.x + n1.x*halfW, p0.y + n1.y*halfW},
p0}; p0};
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+1); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+1);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+2); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+2);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+3); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+3);
} }
else else
{ {
@ -763,13 +840,13 @@ void mg_gl_stroke_joint(mg_gl_encoding_context* context,
{p0.x + n1.x*halfW, p0.y + n1.y*halfW}, {p0.x + n1.x*halfW, p0.y + n1.y*halfW},
p0}; p0};
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+1); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+1);
mg_gl_canvas_encode_element(context, MG_PATH_LINE, points+2); mg_gl_canvas_encode_element(backend, MG_PATH_LINE, points+2);
} }
} }
u32 mg_gl_encode_stroke_subpath(mg_gl_encoding_context* context, u32 mg_gl_encode_stroke_subpath(mg_gl_canvas_backend* backend,
mg_path_elt* elements, mg_path_elt* elements,
mg_path_descriptor* path, mg_path_descriptor* path,
u32 startIndex, u32 startIndex,
@ -786,7 +863,7 @@ u32 mg_gl_encode_stroke_subpath(mg_gl_encoding_context* context,
vec2 endTangent = {0, 0}; vec2 endTangent = {0, 0};
//NOTE(martin): encode first element and compute first tangent //NOTE(martin): encode first element and compute first tangent
mg_gl_encode_stroke_element(context, elements + startIndex, currentPoint, &startTangent, &endTangent, &endPoint); mg_gl_encode_stroke_element(backend, elements + startIndex, currentPoint, &startTangent, &endTangent, &endPoint);
firstTangent = startTangent; firstTangent = startTangent;
previousEndTangent = endTangent; previousEndTangent = endTangent;
@ -794,18 +871,18 @@ u32 mg_gl_encode_stroke_subpath(mg_gl_encoding_context* context,
//NOTE(martin): encode subsequent elements along with their joints //NOTE(martin): encode subsequent elements along with their joints
mg_attributes* attributes = &context->primitive->attributes; mg_attributes* attributes = &backend->primitive->attributes;
u32 eltIndex = startIndex + 1; u32 eltIndex = startIndex + 1;
for(; for(;
eltIndex<eltCount && elements[eltIndex].type != MG_PATH_MOVE; eltIndex<eltCount && elements[eltIndex].type != MG_PATH_MOVE;
eltIndex++) eltIndex++)
{ {
mg_gl_encode_stroke_element(context, elements + eltIndex, currentPoint, &startTangent, &endTangent, &endPoint); mg_gl_encode_stroke_element(backend, elements + eltIndex, currentPoint, &startTangent, &endTangent, &endPoint);
if(attributes->joint != MG_JOINT_NONE) if(attributes->joint != MG_JOINT_NONE)
{ {
mg_gl_stroke_joint(context, currentPoint, previousEndTangent, startTangent); mg_gl_stroke_joint(backend, currentPoint, previousEndTangent, startTangent);
} }
previousEndTangent = endTangent; previousEndTangent = endTangent;
currentPoint = endPoint; currentPoint = endPoint;
@ -820,19 +897,19 @@ u32 mg_gl_encode_stroke_subpath(mg_gl_encoding_context* context,
if(attributes->joint != MG_JOINT_NONE) if(attributes->joint != MG_JOINT_NONE)
{ {
//NOTE(martin): add a closing joint if the path is closed //NOTE(martin): add a closing joint if the path is closed
mg_gl_stroke_joint(context, endPoint, endTangent, firstTangent); mg_gl_stroke_joint(backend, endPoint, endTangent, firstTangent);
} }
} }
else if(attributes->cap == MG_CAP_SQUARE) else if(attributes->cap == MG_CAP_SQUARE)
{ {
//NOTE(martin): add start and end cap //NOTE(martin): add start and end cap
mg_gl_stroke_cap(context, startPoint, (vec2){-startTangent.x, -startTangent.y}); mg_gl_stroke_cap(backend, startPoint, (vec2){-startTangent.x, -startTangent.y});
mg_gl_stroke_cap(context, endPoint, endTangent); mg_gl_stroke_cap(backend, endPoint, endTangent);
} }
return(eltIndex); return(eltIndex);
} }
void mg_gl_encode_stroke(mg_gl_encoding_context* context, void mg_gl_encode_stroke(mg_gl_canvas_backend* backend,
mg_path_elt* elements, mg_path_elt* elements,
mg_path_descriptor* path) mg_path_descriptor* path)
{ {
@ -852,16 +929,13 @@ void mg_gl_encode_stroke(mg_gl_encoding_context* context,
} }
if(startIndex < eltCount) if(startIndex < eltCount)
{ {
startIndex = mg_gl_encode_stroke_subpath(context, elements, path, startIndex, startPoint); startIndex = mg_gl_encode_stroke_subpath(backend, elements, path, startIndex, startPoint);
} }
} }
} }
void mg_gl_render_batch(mg_gl_canvas_backend* backend, void mg_gl_render_batch(mg_gl_canvas_backend* backend,
mg_wgl_surface* surface, mg_wgl_surface* surface,
int pathCount,
int eltCount,
mg_image_data* image, mg_image_data* image,
int tileSize, int tileSize,
int nTilesX, int nTilesX,
@ -869,18 +943,20 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend,
vec2 viewportSize, vec2 viewportSize,
f32 scale) f32 scale)
{ {
//NOTE: send the buffers //NOTE: make the buffers visible to gl
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// GLuint pathBuffer = backend->pathBuffer[backend->bufferIndex].buffer;
// this is wrong, we should send the part we just wrote, which doesn't necessarily start at the begining of the buffer GLuint elementBuffer = backend->elementBuffer[backend->bufferIndex].buffer;
// or, alternatively, we should always write at the beginning of our data buffer and send from there...
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->pathBuffer);
glBufferData(GL_SHADER_STORAGE_BUFFER, LAYOUT_PATH_SIZE*pathCount, backend->pathBufferData, GL_STREAM_DRAW);
glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->elementBuffer); int pathBufferOffset = backend->pathBatchStart * sizeof(mg_gl_path);
glBufferData(GL_SHADER_STORAGE_BUFFER, LAYOUT_PATH_ELT_SIZE*eltCount, backend->elementBufferData, GL_STREAM_DRAW); int elementBufferOffset = backend->eltBatchStart * sizeof(mg_gl_path_elt);
int pathCount = backend->pathCount - backend->pathBatchStart;
int eltCount = backend->eltCount - backend->eltBatchStart;
// --> this always bind the begining of the buffer, with the wrong pathCount/eltCount... glBindBuffer(GL_SHADER_STORAGE_BUFFER, pathBuffer);
glFlushMappedBufferRange(GL_SHADER_STORAGE_BUFFER, pathBufferOffset, pathCount*sizeof(mg_gl_path));
glBindBuffer(GL_SHADER_STORAGE_BUFFER, elementBuffer);
glFlushMappedBufferRange(GL_SHADER_STORAGE_BUFFER, elementBufferOffset, eltCount*sizeof(mg_gl_path_elt));
//NOTE: clear counters //NOTE: clear counters
int zero = 0; int zero = 0;
@ -912,11 +988,11 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend,
for(int i=0; i<pathCount; i += maxWorkGroupCount) for(int i=0; i<pathCount; i += maxWorkGroupCount)
{ {
int pathOffset = i*sizeof(mg_gl_path); int pathOffset = pathBufferOffset + i*sizeof(mg_gl_path);
int pathQueueOffset = i*sizeof(mg_gl_path_queue); int pathQueueOffset = i*sizeof(mg_gl_path_queue);
int count = minimum(maxWorkGroupCount, pathCount-i); int count = minimum(maxWorkGroupCount, pathCount-i);
glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 0, backend->pathBuffer, pathOffset, count*sizeof(mg_gl_path)); glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 0, pathBuffer, pathOffset, count*sizeof(mg_gl_path));
glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 1, backend->pathQueueBuffer, pathQueueOffset, count*sizeof(mg_gl_path_queue)); glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 1, backend->pathQueueBuffer, pathQueueOffset, count*sizeof(mg_gl_path_queue));
glDispatchCompute(count, 1, 1); glDispatchCompute(count, 1, 1);
@ -938,10 +1014,10 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend,
for(int i=0; i<eltCount; i += maxWorkGroupCount) for(int i=0; i<eltCount; i += maxWorkGroupCount)
{ {
int offset = i*sizeof(mg_gl_path_elt); int offset = elementBufferOffset + i*sizeof(mg_gl_path_elt);
int count = minimum(maxWorkGroupCount, eltCount-i); int count = minimum(maxWorkGroupCount, eltCount-i);
glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 0, backend->elementBuffer, offset, count*sizeof(mg_gl_path_elt)); glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 0, elementBuffer, offset, count*sizeof(mg_gl_path_elt));
glDispatchCompute(count, 1, 1); glDispatchCompute(count, 1, 1);
glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT); glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT);
@ -966,7 +1042,7 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend,
//NOTE: merge pass //NOTE: merge pass
glUseProgram(backend->merge); glUseProgram(backend->merge);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, backend->pathBuffer); glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 0, pathBuffer, pathBufferOffset, pathCount*sizeof(mg_gl_path));
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, backend->pathQueueBuffer); glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, backend->pathQueueBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, backend->tileQueueBuffer); glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, backend->tileQueueBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, backend->tileOpCountBuffer); glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, backend->tileOpCountBuffer);
@ -993,7 +1069,7 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend,
//NOTE: raster pass //NOTE: raster pass
glUseProgram(backend->raster); glUseProgram(backend->raster);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 0, backend->pathBuffer); glBindBufferRange(GL_SHADER_STORAGE_BUFFER, 0, pathBuffer, pathBufferOffset, pathCount*sizeof(mg_gl_path));
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, backend->segmentBuffer); glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 1, backend->segmentBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, backend->tileOpBuffer); glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 2, backend->tileOpBuffer);
glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, backend->screenTilesBuffer); glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, backend->screenTilesBuffer);
@ -1030,6 +1106,9 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend,
{ {
log_error("gl error %i\n", err); log_error("gl error %i\n", err);
} }
backend->pathBatchStart = backend->pathCount;
backend->eltBatchStart = backend->eltCount;
} }
///////////////////////////////////////////////////////////////////////// /////////////////////////////////////////////////////////////////////////
@ -1046,6 +1125,15 @@ void mg_gl_canvas_render(mg_canvas_backend* interface,
{ {
mg_gl_canvas_backend* backend = (mg_gl_canvas_backend*)interface; mg_gl_canvas_backend* backend = (mg_gl_canvas_backend*)interface;
//NOTE: roll input buffers
backend->bufferIndex = (backend->bufferIndex + 1) % MG_GL_INPUT_BUFFERS_COUNT;
if(backend->bufferSync[backend->bufferIndex] != 0)
{
glClientWaitSync(backend->bufferSync[backend->bufferIndex], GL_SYNC_FLUSH_COMMANDS_BIT, 0xffffffff);
glDeleteSync(backend->bufferSync[backend->bufferIndex]);
backend->bufferSync[backend->bufferIndex] = 0;
}
//TODO update screen tiles buffer size //TODO update screen tiles buffer size
mg_wgl_surface* surface = backend->surface; mg_wgl_surface* surface = backend->surface;
mp_rect frame = surface->interface.getFrame((mg_surface_data*)surface); mp_rect frame = surface->interface.getFrame((mg_surface_data*)surface);
@ -1070,14 +1158,16 @@ void mg_gl_canvas_render(mg_canvas_backend* interface,
glClearColor(clearColor.r, clearColor.g, clearColor.b, clearColor.a); glClearColor(clearColor.r, clearColor.g, clearColor.b, clearColor.a);
glClear(GL_COLOR_BUFFER_BIT); glClear(GL_COLOR_BUFFER_BIT);
backend->pathCount = 0;
backend->pathBatchStart = 0;
backend->eltCount = 0;
backend->eltBatchStart = 0;
//NOTE: encode and render batches //NOTE: encode and render batches
int pathCount = 0;
vec2 currentPos = {0}; vec2 currentPos = {0};
mg_image currentImage = mg_image_nil(); mg_image currentImage = mg_image_nil();
mg_gl_encoding_context context = {.glEltCount = 0, backend->eltCount = 0;
.elementBufferData = backend->elementBufferData,
.pathBufferData = backend->pathBufferData };
for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++) for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++)
{ {
@ -1089,30 +1179,24 @@ void mg_gl_canvas_render(mg_canvas_backend* interface,
mg_gl_render_batch(backend, mg_gl_render_batch(backend,
surface, surface,
pathCount,
context.glEltCount,
imageData, imageData,
tileSize, tileSize,
nTilesX, nTilesX,
nTilesY, nTilesY,
viewportSize, viewportSize,
scale); scale);
pathCount = 0;
context.glEltCount = 0;
} }
currentImage = primitive->attributes.image; currentImage = primitive->attributes.image;
if(primitive->path.count) if(primitive->path.count)
{ {
context.primitive = primitive; backend->primitive = primitive;
context.pathIndex = pathCount; backend->pathScreenExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
context.pathScreenExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX}; backend->pathUserExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
context.pathUserExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
if(primitive->cmd == MG_CMD_STROKE) if(primitive->cmd == MG_CMD_STROKE)
{ {
mg_gl_encode_stroke(&context, pathElements + primitive->path.startIndex, &primitive->path); mg_gl_encode_stroke(backend, pathElements + primitive->path.startIndex, &primitive->path);
} }
else else
{ {
@ -1126,7 +1210,7 @@ void mg_gl_canvas_render(mg_canvas_backend* interface,
if(elt->type != MG_PATH_MOVE) if(elt->type != MG_PATH_MOVE)
{ {
vec2 p[4] = {currentPos, elt->p[0], elt->p[1], elt->p[2]}; vec2 p[4] = {currentPos, elt->p[0], elt->p[1], elt->p[2]};
mg_gl_canvas_encode_element(&context, elt->type, p); mg_gl_canvas_encode_element(backend, elt->type, p);
segCount++; segCount++;
} }
switch(elt->type) switch(elt->type)
@ -1150,75 +1234,13 @@ void mg_gl_canvas_render(mg_canvas_backend* interface,
} }
} }
//NOTE: push path //NOTE: push path
mg_gl_path* path = &context.pathBufferData[pathCount]; mg_gl_canvas_encode_path(backend, primitive, scale);
pathCount++;
path->cmd = (mg_gl_cmd)primitive->cmd;
path->box = (vec4){context.pathScreenExtents.x,
context.pathScreenExtents.y,
context.pathScreenExtents.z,
context.pathScreenExtents.w};
path->clip = (vec4){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 = (vec4){primitive->attributes.color.r,
primitive->attributes.color.g,
primitive->attributes.color.b,
primitive->attributes.color.a};
mp_rect srcRegion = primitive->attributes.srcRegion;
mp_rect destRegion = {context.pathUserExtents.x,
context.pathUserExtents.y,
context.pathUserExtents.z - context.pathUserExtents.x,
context.pathUserExtents.w - context.pathUserExtents.y};
if(!mg_image_is_nil(primitive->attributes.image))
{
vec2 texSize = mg_image_size(primitive->attributes.image);
mg_mat2x3 srcRegionToImage = {1/texSize.x, 0, srcRegion.x/texSize.x,
0, 1/texSize.y, srcRegion.y/texSize.y};
mg_mat2x3 destRegionToSrcRegion = {srcRegion.w/destRegion.w, 0, 0,
0, srcRegion.h/destRegion.h, 0};
mg_mat2x3 userToDestRegion = {1, 0, -destRegion.x,
0, 1, -destRegion.y};
mg_mat2x3 screenToUser = mg_mat2x3_inv(primitive->attributes.transform);
mg_mat2x3 uvTransform = srcRegionToImage;
uvTransform = mg_mat2x3_mul_m(uvTransform, destRegionToSrcRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, userToDestRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, screenToUser);
//NOTE: mat3 std430 layout is an array of vec3, which are padded to _vec4_ alignment
path->uvTransform[0] = uvTransform.m[0]/scale;
path->uvTransform[1] = uvTransform.m[3]/scale;
path->uvTransform[2] = 0;
path->uvTransform[3] = 0;
path->uvTransform[4] = uvTransform.m[1]/scale;
path->uvTransform[5] = uvTransform.m[4]/scale;
path->uvTransform[6] = 0;
path->uvTransform[7] = 0;
path->uvTransform[8] = uvTransform.m[2];
path->uvTransform[9] = uvTransform.m[5];
path->uvTransform[10] = 1;
path->uvTransform[11] = 0;
}
} }
} }
mg_image_data* imageData = mg_image_data_from_handle(currentImage); mg_image_data* imageData = mg_image_data_from_handle(currentImage);
mg_gl_render_batch(backend, mg_gl_render_batch(backend,
surface, surface,
pathCount,
context.glEltCount,
imageData, imageData,
tileSize, tileSize,
nTilesX, nTilesX,
@ -1226,7 +1248,8 @@ void mg_gl_canvas_render(mg_canvas_backend* interface,
viewportSize, viewportSize,
scale); scale);
//TODO add completion handler for rolling input buffers //NOTE: add fence for rolling input buffers
backend->bufferSync[backend->bufferIndex] = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);
} }
//-------------------------------------------------------------------- //--------------------------------------------------------------------
@ -1252,7 +1275,7 @@ mg_image_data* mg_gl_canvas_image_create(mg_canvas_backend* interface, vec2 size
void mg_gl_canvas_image_destroy(mg_canvas_backend* interface, mg_image_data* imageInterface) void mg_gl_canvas_image_destroy(mg_canvas_backend* interface, mg_image_data* imageInterface)
{ {
//TODO: check that this image belongs to this context //TODO: check that this image belongs to this backend
mg_gl_image* image = (mg_gl_image*)imageInterface; mg_gl_image* image = (mg_gl_image*)imageInterface;
glDeleteTextures(1, &image->texture); glDeleteTextures(1, &image->texture);
free(image); free(image);
@ -1263,7 +1286,7 @@ void mg_gl_canvas_image_upload_region(mg_canvas_backend* interface,
mp_rect region, mp_rect region,
u8* pixels) u8* pixels)
{ {
//TODO: check that this image belongs to this context //TODO: check that this image belongs to this backend
mg_gl_image* image = (mg_gl_image*)imageInterface; mg_gl_image* image = (mg_gl_image*)imageInterface;
glBindTexture(GL_TEXTURE_2D, image->texture); glBindTexture(GL_TEXTURE_2D, image->texture);
glTexSubImage2D(GL_TEXTURE_2D, 0, region.x, region.y, region.w, region.h, GL_RGBA, GL_UNSIGNED_BYTE, pixels); glTexSubImage2D(GL_TEXTURE_2D, 0, region.x, region.y, region.w, region.h, GL_RGBA, GL_UNSIGNED_BYTE, pixels);
@ -1444,14 +1467,30 @@ mg_canvas_backend* gl_canvas_backend_create(mg_wgl_surface* surface)
glGenBuffers(1, &backend->dummyVertexBuffer); glGenBuffers(1, &backend->dummyVertexBuffer);
glBindBuffer(GL_ARRAY_BUFFER, backend->dummyVertexBuffer); glBindBuffer(GL_ARRAY_BUFFER, backend->dummyVertexBuffer);
glGenBuffers(1, &backend->pathBuffer); for(int i=0; i<MG_GL_INPUT_BUFFERS_COUNT; i++)
glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->pathBuffer); {
glBufferData(GL_SHADER_STORAGE_BUFFER, MG_GL_PATH_BUFFER_SIZE, 0, GL_STREAM_DRAW); glGenBuffers(1, &backend->pathBuffer[i].buffer);
glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->pathBuffer[i].buffer);
glBufferStorage(GL_SHADER_STORAGE_BUFFER, MG_GL_PATH_BUFFER_SIZE, 0, GL_MAP_WRITE_BIT|GL_MAP_PERSISTENT_BIT);
backend->pathBuffer[i].size = MG_GL_PATH_BUFFER_SIZE;
backend->pathBuffer[i].contents = glMapBufferRange(GL_SHADER_STORAGE_BUFFER,
0,
MG_GL_PATH_BUFFER_SIZE,
GL_MAP_WRITE_BIT
|GL_MAP_PERSISTENT_BIT
|GL_MAP_FLUSH_EXPLICIT_BIT);
//TODO change flags glGenBuffers(1, &backend->elementBuffer[i].buffer);
glGenBuffers(1, &backend->elementBuffer); glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->elementBuffer[i].buffer);
glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->elementBuffer); glBufferStorage(GL_SHADER_STORAGE_BUFFER, MG_GL_ELEMENT_BUFFER_SIZE, 0, GL_MAP_WRITE_BIT|GL_MAP_PERSISTENT_BIT);
glBufferData(GL_SHADER_STORAGE_BUFFER, MG_GL_ELEMENT_BUFFER_SIZE, 0, GL_STREAM_DRAW); backend->elementBuffer[i].size = MG_GL_ELEMENT_BUFFER_SIZE;
backend->elementBuffer[i].contents = glMapBufferRange(GL_SHADER_STORAGE_BUFFER,
0,
MG_GL_ELEMENT_BUFFER_SIZE,
GL_MAP_WRITE_BIT
|GL_MAP_PERSISTENT_BIT
|GL_MAP_FLUSH_EXPLICIT_BIT);
}
glGenBuffers(1, &backend->segmentBuffer); glGenBuffers(1, &backend->segmentBuffer);
glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->segmentBuffer); glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->segmentBuffer);
@ -1489,9 +1528,6 @@ mg_canvas_backend* gl_canvas_backend_create(mg_wgl_surface* surface)
glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->screenTilesBuffer); glBindBuffer(GL_SHADER_STORAGE_BUFFER, backend->screenTilesBuffer);
glBufferData(GL_SHADER_STORAGE_BUFFER, nTilesX*nTilesY*sizeof(int), 0, GL_DYNAMIC_COPY); glBufferData(GL_SHADER_STORAGE_BUFFER, nTilesX*nTilesY*sizeof(int), 0, GL_DYNAMIC_COPY);
backend->pathBufferData = malloc(MG_GL_PATH_BUFFER_SIZE);
backend->elementBufferData = malloc(MG_GL_ELEMENT_BUFFER_SIZE);
if(err) if(err)
{ {
mg_gl_canvas_destroy((mg_canvas_backend*)backend); mg_gl_canvas_destroy((mg_canvas_backend*)backend);

553
src/gl_loader.c
View File

@ -2,7 +2,7 @@
* *
* @file: gl_loader.c * @file: gl_loader.c
* @note: auto-generated by glapi.py from gl.xml * @note: auto-generated by glapi.py from gl.xml
* @date: 22/022023 * @date: 12/072023
* *
*********************************************************/ *********************************************************/
#include"gl_loader.h" #include"gl_loader.h"
@ -1038,6 +1038,557 @@ void mg_gl_load_gl43(mg_gl_api* api, mg_gl_load_proc loadProc)
api->GetPointerv = loadProc("glGetPointerv"); api->GetPointerv = loadProc("glGetPointerv");
} }
void mg_gl_load_gl44(mg_gl_api* api, mg_gl_load_proc loadProc)
{
api->CullFace = loadProc("glCullFace");
api->FrontFace = loadProc("glFrontFace");
api->Hint = loadProc("glHint");
api->LineWidth = loadProc("glLineWidth");
api->PointSize = loadProc("glPointSize");
api->PolygonMode = loadProc("glPolygonMode");
api->Scissor = loadProc("glScissor");
api->TexParameterf = loadProc("glTexParameterf");
api->TexParameterfv = loadProc("glTexParameterfv");
api->TexParameteri = loadProc("glTexParameteri");
api->TexParameteriv = loadProc("glTexParameteriv");
api->TexImage1D = loadProc("glTexImage1D");
api->TexImage2D = loadProc("glTexImage2D");
api->DrawBuffer = loadProc("glDrawBuffer");
api->Clear = loadProc("glClear");
api->ClearColor = loadProc("glClearColor");
api->ClearStencil = loadProc("glClearStencil");
api->ClearDepth = loadProc("glClearDepth");
api->StencilMask = loadProc("glStencilMask");
api->ColorMask = loadProc("glColorMask");
api->DepthMask = loadProc("glDepthMask");
api->Disable = loadProc("glDisable");
api->Enable = loadProc("glEnable");
api->Finish = loadProc("glFinish");
api->Flush = loadProc("glFlush");
api->BlendFunc = loadProc("glBlendFunc");
api->LogicOp = loadProc("glLogicOp");
api->StencilFunc = loadProc("glStencilFunc");
api->StencilOp = loadProc("glStencilOp");
api->DepthFunc = loadProc("glDepthFunc");
api->PixelStoref = loadProc("glPixelStoref");
api->PixelStorei = loadProc("glPixelStorei");
api->ReadBuffer = loadProc("glReadBuffer");
api->ReadPixels = loadProc("glReadPixels");
api->GetBooleanv = loadProc("glGetBooleanv");
api->GetDoublev = loadProc("glGetDoublev");
api->GetError = loadProc("glGetError");
api->GetFloatv = loadProc("glGetFloatv");
api->GetIntegerv = loadProc("glGetIntegerv");
api->GetString = loadProc("glGetString");
api->GetTexImage = loadProc("glGetTexImage");
api->GetTexParameterfv = loadProc("glGetTexParameterfv");
api->GetTexParameteriv = loadProc("glGetTexParameteriv");
api->GetTexLevelParameterfv = loadProc("glGetTexLevelParameterfv");
api->GetTexLevelParameteriv = loadProc("glGetTexLevelParameteriv");
api->IsEnabled = loadProc("glIsEnabled");
api->DepthRange = loadProc("glDepthRange");
api->Viewport = loadProc("glViewport");
api->DrawArrays = loadProc("glDrawArrays");
api->DrawElements = loadProc("glDrawElements");
api->PolygonOffset = loadProc("glPolygonOffset");
api->CopyTexImage1D = loadProc("glCopyTexImage1D");
api->CopyTexImage2D = loadProc("glCopyTexImage2D");
api->CopyTexSubImage1D = loadProc("glCopyTexSubImage1D");
api->CopyTexSubImage2D = loadProc("glCopyTexSubImage2D");
api->TexSubImage1D = loadProc("glTexSubImage1D");
api->TexSubImage2D = loadProc("glTexSubImage2D");
api->BindTexture = loadProc("glBindTexture");
api->DeleteTextures = loadProc("glDeleteTextures");
api->GenTextures = loadProc("glGenTextures");
api->IsTexture = loadProc("glIsTexture");
api->DrawRangeElements = loadProc("glDrawRangeElements");
api->TexImage3D = loadProc("glTexImage3D");
api->TexSubImage3D = loadProc("glTexSubImage3D");
api->CopyTexSubImage3D = loadProc("glCopyTexSubImage3D");
api->ActiveTexture = loadProc("glActiveTexture");
api->SampleCoverage = loadProc("glSampleCoverage");
api->CompressedTexImage3D = loadProc("glCompressedTexImage3D");
api->CompressedTexImage2D = loadProc("glCompressedTexImage2D");
api->CompressedTexImage1D = loadProc("glCompressedTexImage1D");
api->CompressedTexSubImage3D = loadProc("glCompressedTexSubImage3D");
api->CompressedTexSubImage2D = loadProc("glCompressedTexSubImage2D");
api->CompressedTexSubImage1D = loadProc("glCompressedTexSubImage1D");
api->GetCompressedTexImage = loadProc("glGetCompressedTexImage");
api->BlendFuncSeparate = loadProc("glBlendFuncSeparate");
api->MultiDrawArrays = loadProc("glMultiDrawArrays");
api->MultiDrawElements = loadProc("glMultiDrawElements");
api->PointParameterf = loadProc("glPointParameterf");
api->PointParameterfv = loadProc("glPointParameterfv");
api->PointParameteri = loadProc("glPointParameteri");
api->PointParameteriv = loadProc("glPointParameteriv");
api->BlendColor = loadProc("glBlendColor");
api->BlendEquation = loadProc("glBlendEquation");
api->GenQueries = loadProc("glGenQueries");
api->DeleteQueries = loadProc("glDeleteQueries");
api->IsQuery = loadProc("glIsQuery");
api->BeginQuery = loadProc("glBeginQuery");
api->EndQuery = loadProc("glEndQuery");
api->GetQueryiv = loadProc("glGetQueryiv");
api->GetQueryObjectiv = loadProc("glGetQueryObjectiv");
api->GetQueryObjectuiv = loadProc("glGetQueryObjectuiv");
api->BindBuffer = loadProc("glBindBuffer");
api->DeleteBuffers = loadProc("glDeleteBuffers");
api->GenBuffers = loadProc("glGenBuffers");
api->IsBuffer = loadProc("glIsBuffer");
api->BufferData = loadProc("glBufferData");
api->BufferSubData = loadProc("glBufferSubData");
api->GetBufferSubData = loadProc("glGetBufferSubData");
api->MapBuffer = loadProc("glMapBuffer");
api->UnmapBuffer = loadProc("glUnmapBuffer");
api->GetBufferParameteriv = loadProc("glGetBufferParameteriv");
api->GetBufferPointerv = loadProc("glGetBufferPointerv");
api->BlendEquationSeparate = loadProc("glBlendEquationSeparate");
api->DrawBuffers = loadProc("glDrawBuffers");
api->StencilOpSeparate = loadProc("glStencilOpSeparate");
api->StencilFuncSeparate = loadProc("glStencilFuncSeparate");
api->StencilMaskSeparate = loadProc("glStencilMaskSeparate");
api->AttachShader = loadProc("glAttachShader");
api->BindAttribLocation = loadProc("glBindAttribLocation");
api->CompileShader = loadProc("glCompileShader");
api->CreateProgram = loadProc("glCreateProgram");
api->CreateShader = loadProc("glCreateShader");
api->DeleteProgram = loadProc("glDeleteProgram");
api->DeleteShader = loadProc("glDeleteShader");
api->DetachShader = loadProc("glDetachShader");
api->DisableVertexAttribArray = loadProc("glDisableVertexAttribArray");
api->EnableVertexAttribArray = loadProc("glEnableVertexAttribArray");
api->GetActiveAttrib = loadProc("glGetActiveAttrib");
api->GetActiveUniform = loadProc("glGetActiveUniform");
api->GetAttachedShaders = loadProc("glGetAttachedShaders");
api->GetAttribLocation = loadProc("glGetAttribLocation");
api->GetProgramiv = loadProc("glGetProgramiv");
api->GetProgramInfoLog = loadProc("glGetProgramInfoLog");
api->GetShaderiv = loadProc("glGetShaderiv");
api->GetShaderInfoLog = loadProc("glGetShaderInfoLog");
api->GetShaderSource = loadProc("glGetShaderSource");
api->GetUniformLocation = loadProc("glGetUniformLocation");
api->GetUniformfv = loadProc("glGetUniformfv");
api->GetUniformiv = loadProc("glGetUniformiv");
api->GetVertexAttribdv = loadProc("glGetVertexAttribdv");
api->GetVertexAttribfv = loadProc("glGetVertexAttribfv");
api->GetVertexAttribiv = loadProc("glGetVertexAttribiv");
api->GetVertexAttribPointerv = loadProc("glGetVertexAttribPointerv");
api->IsProgram = loadProc("glIsProgram");
api->IsShader = loadProc("glIsShader");
api->LinkProgram = loadProc("glLinkProgram");
api->ShaderSource = loadProc("glShaderSource");
api->UseProgram = loadProc("glUseProgram");
api->Uniform1f = loadProc("glUniform1f");
api->Uniform2f = loadProc("glUniform2f");
api->Uniform3f = loadProc("glUniform3f");
api->Uniform4f = loadProc("glUniform4f");
api->Uniform1i = loadProc("glUniform1i");
api->Uniform2i = loadProc("glUniform2i");
api->Uniform3i = loadProc("glUniform3i");
api->Uniform4i = loadProc("glUniform4i");
api->Uniform1fv = loadProc("glUniform1fv");
api->Uniform2fv = loadProc("glUniform2fv");
api->Uniform3fv = loadProc("glUniform3fv");
api->Uniform4fv = loadProc("glUniform4fv");
api->Uniform1iv = loadProc("glUniform1iv");
api->Uniform2iv = loadProc("glUniform2iv");
api->Uniform3iv = loadProc("glUniform3iv");
api->Uniform4iv = loadProc("glUniform4iv");
api->UniformMatrix2fv = loadProc("glUniformMatrix2fv");
api->UniformMatrix3fv = loadProc("glUniformMatrix3fv");
api->UniformMatrix4fv = loadProc("glUniformMatrix4fv");
api->ValidateProgram = loadProc("glValidateProgram");
api->VertexAttrib1d = loadProc("glVertexAttrib1d");
api->VertexAttrib1dv = loadProc("glVertexAttrib1dv");
api->VertexAttrib1f = loadProc("glVertexAttrib1f");
api->VertexAttrib1fv = loadProc("glVertexAttrib1fv");
api->VertexAttrib1s = loadProc("glVertexAttrib1s");
api->VertexAttrib1sv = loadProc("glVertexAttrib1sv");
api->VertexAttrib2d = loadProc("glVertexAttrib2d");
api->VertexAttrib2dv = loadProc("glVertexAttrib2dv");
api->VertexAttrib2f = loadProc("glVertexAttrib2f");
api->VertexAttrib2fv = loadProc("glVertexAttrib2fv");
api->VertexAttrib2s = loadProc("glVertexAttrib2s");
api->VertexAttrib2sv = loadProc("glVertexAttrib2sv");
api->VertexAttrib3d = loadProc("glVertexAttrib3d");
api->VertexAttrib3dv = loadProc("glVertexAttrib3dv");
api->VertexAttrib3f = loadProc("glVertexAttrib3f");
api->VertexAttrib3fv = loadProc("glVertexAttrib3fv");
api->VertexAttrib3s = loadProc("glVertexAttrib3s");
api->VertexAttrib3sv = loadProc("glVertexAttrib3sv");
api->VertexAttrib4Nbv = loadProc("glVertexAttrib4Nbv");
api->VertexAttrib4Niv = loadProc("glVertexAttrib4Niv");
api->VertexAttrib4Nsv = loadProc("glVertexAttrib4Nsv");
api->VertexAttrib4Nub = loadProc("glVertexAttrib4Nub");
api->VertexAttrib4Nubv = loadProc("glVertexAttrib4Nubv");
api->VertexAttrib4Nuiv = loadProc("glVertexAttrib4Nuiv");
api->VertexAttrib4Nusv = loadProc("glVertexAttrib4Nusv");
api->VertexAttrib4bv = loadProc("glVertexAttrib4bv");
api->VertexAttrib4d = loadProc("glVertexAttrib4d");
api->VertexAttrib4dv = loadProc("glVertexAttrib4dv");
api->VertexAttrib4f = loadProc("glVertexAttrib4f");
api->VertexAttrib4fv = loadProc("glVertexAttrib4fv");
api->VertexAttrib4iv = loadProc("glVertexAttrib4iv");
api->VertexAttrib4s = loadProc("glVertexAttrib4s");
api->VertexAttrib4sv = loadProc("glVertexAttrib4sv");
api->VertexAttrib4ubv = loadProc("glVertexAttrib4ubv");
api->VertexAttrib4uiv = loadProc("glVertexAttrib4uiv");
api->VertexAttrib4usv = loadProc("glVertexAttrib4usv");
api->VertexAttribPointer = loadProc("glVertexAttribPointer");
api->UniformMatrix2x3fv = loadProc("glUniformMatrix2x3fv");
api->UniformMatrix3x2fv = loadProc("glUniformMatrix3x2fv");
api->UniformMatrix2x4fv = loadProc("glUniformMatrix2x4fv");
api->UniformMatrix4x2fv = loadProc("glUniformMatrix4x2fv");
api->UniformMatrix3x4fv = loadProc("glUniformMatrix3x4fv");
api->UniformMatrix4x3fv = loadProc("glUniformMatrix4x3fv");
api->ColorMaski = loadProc("glColorMaski");
api->GetBooleani_v = loadProc("glGetBooleani_v");
api->GetIntegeri_v = loadProc("glGetIntegeri_v");
api->Enablei = loadProc("glEnablei");
api->Disablei = loadProc("glDisablei");
api->IsEnabledi = loadProc("glIsEnabledi");
api->BeginTransformFeedback = loadProc("glBeginTransformFeedback");
api->EndTransformFeedback = loadProc("glEndTransformFeedback");
api->BindBufferRange = loadProc("glBindBufferRange");
api->BindBufferBase = loadProc("glBindBufferBase");
api->TransformFeedbackVaryings = loadProc("glTransformFeedbackVaryings");
api->GetTransformFeedbackVarying = loadProc("glGetTransformFeedbackVarying");
api->ClampColor = loadProc("glClampColor");
api->BeginConditionalRender = loadProc("glBeginConditionalRender");
api->EndConditionalRender = loadProc("glEndConditionalRender");
api->VertexAttribIPointer = loadProc("glVertexAttribIPointer");
api->GetVertexAttribIiv = loadProc("glGetVertexAttribIiv");
api->GetVertexAttribIuiv = loadProc("glGetVertexAttribIuiv");
api->VertexAttribI1i = loadProc("glVertexAttribI1i");
api->VertexAttribI2i = loadProc("glVertexAttribI2i");
api->VertexAttribI3i = loadProc("glVertexAttribI3i");
api->VertexAttribI4i = loadProc("glVertexAttribI4i");
api->VertexAttribI1ui = loadProc("glVertexAttribI1ui");
api->VertexAttribI2ui = loadProc("glVertexAttribI2ui");
api->VertexAttribI3ui = loadProc("glVertexAttribI3ui");
api->VertexAttribI4ui = loadProc("glVertexAttribI4ui");
api->VertexAttribI1iv = loadProc("glVertexAttribI1iv");
api->VertexAttribI2iv = loadProc("glVertexAttribI2iv");
api->VertexAttribI3iv = loadProc("glVertexAttribI3iv");
api->VertexAttribI4iv = loadProc("glVertexAttribI4iv");
api->VertexAttribI1uiv = loadProc("glVertexAttribI1uiv");
api->VertexAttribI2uiv = loadProc("glVertexAttribI2uiv");
api->VertexAttribI3uiv = loadProc("glVertexAttribI3uiv");
api->VertexAttribI4uiv = loadProc("glVertexAttribI4uiv");
api->VertexAttribI4bv = loadProc("glVertexAttribI4bv");
api->VertexAttribI4sv = loadProc("glVertexAttribI4sv");
api->VertexAttribI4ubv = loadProc("glVertexAttribI4ubv");
api->VertexAttribI4usv = loadProc("glVertexAttribI4usv");
api->GetUniformuiv = loadProc("glGetUniformuiv");
api->BindFragDataLocation = loadProc("glBindFragDataLocation");
api->GetFragDataLocation = loadProc("glGetFragDataLocation");
api->Uniform1ui = loadProc("glUniform1ui");
api->Uniform2ui = loadProc("glUniform2ui");
api->Uniform3ui = loadProc("glUniform3ui");
api->Uniform4ui = loadProc("glUniform4ui");
api->Uniform1uiv = loadProc("glUniform1uiv");
api->Uniform2uiv = loadProc("glUniform2uiv");
api->Uniform3uiv = loadProc("glUniform3uiv");
api->Uniform4uiv = loadProc("glUniform4uiv");
api->TexParameterIiv = loadProc("glTexParameterIiv");
api->TexParameterIuiv = loadProc("glTexParameterIuiv");
api->GetTexParameterIiv = loadProc("glGetTexParameterIiv");
api->GetTexParameterIuiv = loadProc("glGetTexParameterIuiv");
api->ClearBufferiv = loadProc("glClearBufferiv");
api->ClearBufferuiv = loadProc("glClearBufferuiv");
api->ClearBufferfv = loadProc("glClearBufferfv");
api->ClearBufferfi = loadProc("glClearBufferfi");
api->GetStringi = loadProc("glGetStringi");
api->IsRenderbuffer = loadProc("glIsRenderbuffer");
api->BindRenderbuffer = loadProc("glBindRenderbuffer");
api->DeleteRenderbuffers = loadProc("glDeleteRenderbuffers");
api->GenRenderbuffers = loadProc("glGenRenderbuffers");
api->RenderbufferStorage = loadProc("glRenderbufferStorage");
api->GetRenderbufferParameteriv = loadProc("glGetRenderbufferParameteriv");
api->IsFramebuffer = loadProc("glIsFramebuffer");
api->BindFramebuffer = loadProc("glBindFramebuffer");
api->DeleteFramebuffers = loadProc("glDeleteFramebuffers");
api->GenFramebuffers = loadProc("glGenFramebuffers");
api->CheckFramebufferStatus = loadProc("glCheckFramebufferStatus");
api->FramebufferTexture1D = loadProc("glFramebufferTexture1D");
api->FramebufferTexture2D = loadProc("glFramebufferTexture2D");
api->FramebufferTexture3D = loadProc("glFramebufferTexture3D");
api->FramebufferRenderbuffer = loadProc("glFramebufferRenderbuffer");
api->GetFramebufferAttachmentParameteriv = loadProc("glGetFramebufferAttachmentParameteriv");
api->GenerateMipmap = loadProc("glGenerateMipmap");
api->BlitFramebuffer = loadProc("glBlitFramebuffer");
api->RenderbufferStorageMultisample = loadProc("glRenderbufferStorageMultisample");
api->FramebufferTextureLayer = loadProc("glFramebufferTextureLayer");
api->MapBufferRange = loadProc("glMapBufferRange");
api->FlushMappedBufferRange = loadProc("glFlushMappedBufferRange");
api->BindVertexArray = loadProc("glBindVertexArray");
api->DeleteVertexArrays = loadProc("glDeleteVertexArrays");
api->GenVertexArrays = loadProc("glGenVertexArrays");
api->IsVertexArray = loadProc("glIsVertexArray");
api->DrawArraysInstanced = loadProc("glDrawArraysInstanced");
api->DrawElementsInstanced = loadProc("glDrawElementsInstanced");
api->TexBuffer = loadProc("glTexBuffer");
api->PrimitiveRestartIndex = loadProc("glPrimitiveRestartIndex");
api->CopyBufferSubData = loadProc("glCopyBufferSubData");
api->GetUniformIndices = loadProc("glGetUniformIndices");
api->GetActiveUniformsiv = loadProc("glGetActiveUniformsiv");
api->GetActiveUniformName = loadProc("glGetActiveUniformName");
api->GetUniformBlockIndex = loadProc("glGetUniformBlockIndex");
api->GetActiveUniformBlockiv = loadProc("glGetActiveUniformBlockiv");
api->GetActiveUniformBlockName = loadProc("glGetActiveUniformBlockName");
api->UniformBlockBinding = loadProc("glUniformBlockBinding");
api->BindBufferRange = loadProc("glBindBufferRange");
api->BindBufferBase = loadProc("glBindBufferBase");
api->GetIntegeri_v = loadProc("glGetIntegeri_v");
api->DrawElementsBaseVertex = loadProc("glDrawElementsBaseVertex");
api->DrawRangeElementsBaseVertex = loadProc("glDrawRangeElementsBaseVertex");
api->DrawElementsInstancedBaseVertex = loadProc("glDrawElementsInstancedBaseVertex");
api->MultiDrawElementsBaseVertex = loadProc("glMultiDrawElementsBaseVertex");
api->ProvokingVertex = loadProc("glProvokingVertex");
api->FenceSync = loadProc("glFenceSync");
api->IsSync = loadProc("glIsSync");
api->DeleteSync = loadProc("glDeleteSync");
api->ClientWaitSync = loadProc("glClientWaitSync");
api->WaitSync = loadProc("glWaitSync");
api->GetInteger64v = loadProc("glGetInteger64v");
api->GetSynciv = loadProc("glGetSynciv");
api->GetInteger64i_v = loadProc("glGetInteger64i_v");
api->GetBufferParameteri64v = loadProc("glGetBufferParameteri64v");
api->FramebufferTexture = loadProc("glFramebufferTexture");
api->TexImage2DMultisample = loadProc("glTexImage2DMultisample");
api->TexImage3DMultisample = loadProc("glTexImage3DMultisample");
api->GetMultisamplefv = loadProc("glGetMultisamplefv");
api->SampleMaski = loadProc("glSampleMaski");
api->BindFragDataLocationIndexed = loadProc("glBindFragDataLocationIndexed");
api->GetFragDataIndex = loadProc("glGetFragDataIndex");
api->GenSamplers = loadProc("glGenSamplers");
api->DeleteSamplers = loadProc("glDeleteSamplers");
api->IsSampler = loadProc("glIsSampler");
api->BindSampler = loadProc("glBindSampler");
api->SamplerParameteri = loadProc("glSamplerParameteri");
api->SamplerParameteriv = loadProc("glSamplerParameteriv");
api->SamplerParameterf = loadProc("glSamplerParameterf");
api->SamplerParameterfv = loadProc("glSamplerParameterfv");
api->SamplerParameterIiv = loadProc("glSamplerParameterIiv");
api->SamplerParameterIuiv = loadProc("glSamplerParameterIuiv");
api->GetSamplerParameteriv = loadProc("glGetSamplerParameteriv");
api->GetSamplerParameterIiv = loadProc("glGetSamplerParameterIiv");
api->GetSamplerParameterfv = loadProc("glGetSamplerParameterfv");
api->GetSamplerParameterIuiv = loadProc("glGetSamplerParameterIuiv");
api->QueryCounter = loadProc("glQueryCounter");
api->GetQueryObjecti64v = loadProc("glGetQueryObjecti64v");
api->GetQueryObjectui64v = loadProc("glGetQueryObjectui64v");
api->VertexAttribDivisor = loadProc("glVertexAttribDivisor");
api->VertexAttribP1ui = loadProc("glVertexAttribP1ui");
api->VertexAttribP1uiv = loadProc("glVertexAttribP1uiv");
api->VertexAttribP2ui = loadProc("glVertexAttribP2ui");
api->VertexAttribP2uiv = loadProc("glVertexAttribP2uiv");
api->VertexAttribP3ui = loadProc("glVertexAttribP3ui");
api->VertexAttribP3uiv = loadProc("glVertexAttribP3uiv");
api->VertexAttribP4ui = loadProc("glVertexAttribP4ui");
api->VertexAttribP4uiv = loadProc("glVertexAttribP4uiv");
api->MinSampleShading = loadProc("glMinSampleShading");
api->BlendEquationi = loadProc("glBlendEquationi");
api->BlendEquationSeparatei = loadProc("glBlendEquationSeparatei");
api->BlendFunci = loadProc("glBlendFunci");
api->BlendFuncSeparatei = loadProc("glBlendFuncSeparatei");
api->DrawArraysIndirect = loadProc("glDrawArraysIndirect");
api->DrawElementsIndirect = loadProc("glDrawElementsIndirect");
api->Uniform1d = loadProc("glUniform1d");
api->Uniform2d = loadProc("glUniform2d");
api->Uniform3d = loadProc("glUniform3d");
api->Uniform4d = loadProc("glUniform4d");
api->Uniform1dv = loadProc("glUniform1dv");
api->Uniform2dv = loadProc("glUniform2dv");
api->Uniform3dv = loadProc("glUniform3dv");
api->Uniform4dv = loadProc("glUniform4dv");
api->UniformMatrix2dv = loadProc("glUniformMatrix2dv");
api->UniformMatrix3dv = loadProc("glUniformMatrix3dv");
api->UniformMatrix4dv = loadProc("glUniformMatrix4dv");
api->UniformMatrix2x3dv = loadProc("glUniformMatrix2x3dv");
api->UniformMatrix2x4dv = loadProc("glUniformMatrix2x4dv");
api->UniformMatrix3x2dv = loadProc("glUniformMatrix3x2dv");
api->UniformMatrix3x4dv = loadProc("glUniformMatrix3x4dv");
api->UniformMatrix4x2dv = loadProc("glUniformMatrix4x2dv");
api->UniformMatrix4x3dv = loadProc("glUniformMatrix4x3dv");
api->GetUniformdv = loadProc("glGetUniformdv");
api->GetSubroutineUniformLocation = loadProc("glGetSubroutineUniformLocation");
api->GetSubroutineIndex = loadProc("glGetSubroutineIndex");
api->GetActiveSubroutineUniformiv = loadProc("glGetActiveSubroutineUniformiv");
api->GetActiveSubroutineUniformName = loadProc("glGetActiveSubroutineUniformName");
api->GetActiveSubroutineName = loadProc("glGetActiveSubroutineName");
api->UniformSubroutinesuiv = loadProc("glUniformSubroutinesuiv");
api->GetUniformSubroutineuiv = loadProc("glGetUniformSubroutineuiv");
api->GetProgramStageiv = loadProc("glGetProgramStageiv");
api->PatchParameteri = loadProc("glPatchParameteri");
api->PatchParameterfv = loadProc("glPatchParameterfv");
api->BindTransformFeedback = loadProc("glBindTransformFeedback");
api->DeleteTransformFeedbacks = loadProc("glDeleteTransformFeedbacks");
api->GenTransformFeedbacks = loadProc("glGenTransformFeedbacks");
api->IsTransformFeedback = loadProc("glIsTransformFeedback");
api->PauseTransformFeedback = loadProc("glPauseTransformFeedback");
api->ResumeTransformFeedback = loadProc("glResumeTransformFeedback");
api->DrawTransformFeedback = loadProc("glDrawTransformFeedback");
api->DrawTransformFeedbackStream = loadProc("glDrawTransformFeedbackStream");
api->BeginQueryIndexed = loadProc("glBeginQueryIndexed");
api->EndQueryIndexed = loadProc("glEndQueryIndexed");
api->GetQueryIndexediv = loadProc("glGetQueryIndexediv");
api->ReleaseShaderCompiler = loadProc("glReleaseShaderCompiler");
api->ShaderBinary = loadProc("glShaderBinary");
api->GetShaderPrecisionFormat = loadProc("glGetShaderPrecisionFormat");
api->DepthRangef = loadProc("glDepthRangef");
api->ClearDepthf = loadProc("glClearDepthf");
api->GetProgramBinary = loadProc("glGetProgramBinary");
api->ProgramBinary = loadProc("glProgramBinary");
api->ProgramParameteri = loadProc("glProgramParameteri");
api->UseProgramStages = loadProc("glUseProgramStages");
api->ActiveShaderProgram = loadProc("glActiveShaderProgram");
api->CreateShaderProgramv = loadProc("glCreateShaderProgramv");
api->BindProgramPipeline = loadProc("glBindProgramPipeline");
api->DeleteProgramPipelines = loadProc("glDeleteProgramPipelines");
api->GenProgramPipelines = loadProc("glGenProgramPipelines");
api->IsProgramPipeline = loadProc("glIsProgramPipeline");
api->GetProgramPipelineiv = loadProc("glGetProgramPipelineiv");
api->ProgramParameteri = loadProc("glProgramParameteri");
api->ProgramUniform1i = loadProc("glProgramUniform1i");
api->ProgramUniform1iv = loadProc("glProgramUniform1iv");
api->ProgramUniform1f = loadProc("glProgramUniform1f");
api->ProgramUniform1fv = loadProc("glProgramUniform1fv");
api->ProgramUniform1d = loadProc("glProgramUniform1d");
api->ProgramUniform1dv = loadProc("glProgramUniform1dv");
api->ProgramUniform1ui = loadProc("glProgramUniform1ui");
api->ProgramUniform1uiv = loadProc("glProgramUniform1uiv");
api->ProgramUniform2i = loadProc("glProgramUniform2i");
api->ProgramUniform2iv = loadProc("glProgramUniform2iv");
api->ProgramUniform2f = loadProc("glProgramUniform2f");
api->ProgramUniform2fv = loadProc("glProgramUniform2fv");
api->ProgramUniform2d = loadProc("glProgramUniform2d");
api->ProgramUniform2dv = loadProc("glProgramUniform2dv");
api->ProgramUniform2ui = loadProc("glProgramUniform2ui");
api->ProgramUniform2uiv = loadProc("glProgramUniform2uiv");
api->ProgramUniform3i = loadProc("glProgramUniform3i");
api->ProgramUniform3iv = loadProc("glProgramUniform3iv");
api->ProgramUniform3f = loadProc("glProgramUniform3f");
api->ProgramUniform3fv = loadProc("glProgramUniform3fv");
api->ProgramUniform3d = loadProc("glProgramUniform3d");
api->ProgramUniform3dv = loadProc("glProgramUniform3dv");
api->ProgramUniform3ui = loadProc("glProgramUniform3ui");
api->ProgramUniform3uiv = loadProc("glProgramUniform3uiv");
api->ProgramUniform4i = loadProc("glProgramUniform4i");
api->ProgramUniform4iv = loadProc("glProgramUniform4iv");
api->ProgramUniform4f = loadProc("glProgramUniform4f");
api->ProgramUniform4fv = loadProc("glProgramUniform4fv");
api->ProgramUniform4d = loadProc("glProgramUniform4d");
api->ProgramUniform4dv = loadProc("glProgramUniform4dv");
api->ProgramUniform4ui = loadProc("glProgramUniform4ui");
api->ProgramUniform4uiv = loadProc("glProgramUniform4uiv");
api->ProgramUniformMatrix2fv = loadProc("glProgramUniformMatrix2fv");
api->ProgramUniformMatrix3fv = loadProc("glProgramUniformMatrix3fv");
api->ProgramUniformMatrix4fv = loadProc("glProgramUniformMatrix4fv");
api->ProgramUniformMatrix2dv = loadProc("glProgramUniformMatrix2dv");
api->ProgramUniformMatrix3dv = loadProc("glProgramUniformMatrix3dv");
api->ProgramUniformMatrix4dv = loadProc("glProgramUniformMatrix4dv");
api->ProgramUniformMatrix2x3fv = loadProc("glProgramUniformMatrix2x3fv");
api->ProgramUniformMatrix3x2fv = loadProc("glProgramUniformMatrix3x2fv");
api->ProgramUniformMatrix2x4fv = loadProc("glProgramUniformMatrix2x4fv");
api->ProgramUniformMatrix4x2fv = loadProc("glProgramUniformMatrix4x2fv");
api->ProgramUniformMatrix3x4fv = loadProc("glProgramUniformMatrix3x4fv");
api->ProgramUniformMatrix4x3fv = loadProc("glProgramUniformMatrix4x3fv");
api->ProgramUniformMatrix2x3dv = loadProc("glProgramUniformMatrix2x3dv");
api->ProgramUniformMatrix3x2dv = loadProc("glProgramUniformMatrix3x2dv");
api->ProgramUniformMatrix2x4dv = loadProc("glProgramUniformMatrix2x4dv");
api->ProgramUniformMatrix4x2dv = loadProc("glProgramUniformMatrix4x2dv");
api->ProgramUniformMatrix3x4dv = loadProc("glProgramUniformMatrix3x4dv");
api->ProgramUniformMatrix4x3dv = loadProc("glProgramUniformMatrix4x3dv");
api->ValidateProgramPipeline = loadProc("glValidateProgramPipeline");
api->GetProgramPipelineInfoLog = loadProc("glGetProgramPipelineInfoLog");
api->VertexAttribL1d = loadProc("glVertexAttribL1d");
api->VertexAttribL2d = loadProc("glVertexAttribL2d");
api->VertexAttribL3d = loadProc("glVertexAttribL3d");
api->VertexAttribL4d = loadProc("glVertexAttribL4d");
api->VertexAttribL1dv = loadProc("glVertexAttribL1dv");
api->VertexAttribL2dv = loadProc("glVertexAttribL2dv");
api->VertexAttribL3dv = loadProc("glVertexAttribL3dv");
api->VertexAttribL4dv = loadProc("glVertexAttribL4dv");
api->VertexAttribLPointer = loadProc("glVertexAttribLPointer");
api->GetVertexAttribLdv = loadProc("glGetVertexAttribLdv");
api->ViewportArrayv = loadProc("glViewportArrayv");
api->ViewportIndexedf = loadProc("glViewportIndexedf");
api->ViewportIndexedfv = loadProc("glViewportIndexedfv");
api->ScissorArrayv = loadProc("glScissorArrayv");
api->ScissorIndexed = loadProc("glScissorIndexed");
api->ScissorIndexedv = loadProc("glScissorIndexedv");
api->DepthRangeArrayv = loadProc("glDepthRangeArrayv");
api->DepthRangeIndexed = loadProc("glDepthRangeIndexed");
api->GetFloati_v = loadProc("glGetFloati_v");
api->GetDoublei_v = loadProc("glGetDoublei_v");
api->DrawArraysInstancedBaseInstance = loadProc("glDrawArraysInstancedBaseInstance");
api->DrawElementsInstancedBaseInstance = loadProc("glDrawElementsInstancedBaseInstance");
api->DrawElementsInstancedBaseVertexBaseInstance = loadProc("glDrawElementsInstancedBaseVertexBaseInstance");
api->GetInternalformativ = loadProc("glGetInternalformativ");
api->GetActiveAtomicCounterBufferiv = loadProc("glGetActiveAtomicCounterBufferiv");
api->BindImageTexture = loadProc("glBindImageTexture");
api->MemoryBarrier = loadProc("glMemoryBarrier");
api->TexStorage1D = loadProc("glTexStorage1D");
api->TexStorage2D = loadProc("glTexStorage2D");
api->TexStorage3D = loadProc("glTexStorage3D");
api->DrawTransformFeedbackInstanced = loadProc("glDrawTransformFeedbackInstanced");
api->DrawTransformFeedbackStreamInstanced = loadProc("glDrawTransformFeedbackStreamInstanced");
api->ClearBufferData = loadProc("glClearBufferData");
api->ClearBufferSubData = loadProc("glClearBufferSubData");
api->DispatchCompute = loadProc("glDispatchCompute");
api->DispatchComputeIndirect = loadProc("glDispatchComputeIndirect");
api->CopyImageSubData = loadProc("glCopyImageSubData");
api->FramebufferParameteri = loadProc("glFramebufferParameteri");
api->GetFramebufferParameteriv = loadProc("glGetFramebufferParameteriv");
api->GetInternalformati64v = loadProc("glGetInternalformati64v");
api->InvalidateTexSubImage = loadProc("glInvalidateTexSubImage");
api->InvalidateTexImage = loadProc("glInvalidateTexImage");
api->InvalidateBufferSubData = loadProc("glInvalidateBufferSubData");
api->InvalidateBufferData = loadProc("glInvalidateBufferData");
api->InvalidateFramebuffer = loadProc("glInvalidateFramebuffer");
api->InvalidateSubFramebuffer = loadProc("glInvalidateSubFramebuffer");
api->MultiDrawArraysIndirect = loadProc("glMultiDrawArraysIndirect");
api->MultiDrawElementsIndirect = loadProc("glMultiDrawElementsIndirect");
api->GetProgramInterfaceiv = loadProc("glGetProgramInterfaceiv");
api->GetProgramResourceIndex = loadProc("glGetProgramResourceIndex");
api->GetProgramResourceName = loadProc("glGetProgramResourceName");
api->GetProgramResourceiv = loadProc("glGetProgramResourceiv");
api->GetProgramResourceLocation = loadProc("glGetProgramResourceLocation");
api->GetProgramResourceLocationIndex = loadProc("glGetProgramResourceLocationIndex");
api->ShaderStorageBlockBinding = loadProc("glShaderStorageBlockBinding");
api->TexBufferRange = loadProc("glTexBufferRange");
api->TexStorage2DMultisample = loadProc("glTexStorage2DMultisample");
api->TexStorage3DMultisample = loadProc("glTexStorage3DMultisample");
api->TextureView = loadProc("glTextureView");
api->BindVertexBuffer = loadProc("glBindVertexBuffer");
api->VertexAttribFormat = loadProc("glVertexAttribFormat");
api->VertexAttribIFormat = loadProc("glVertexAttribIFormat");
api->VertexAttribLFormat = loadProc("glVertexAttribLFormat");
api->VertexAttribBinding = loadProc("glVertexAttribBinding");
api->VertexBindingDivisor = loadProc("glVertexBindingDivisor");
api->DebugMessageControl = loadProc("glDebugMessageControl");
api->DebugMessageInsert = loadProc("glDebugMessageInsert");
api->DebugMessageCallback = loadProc("glDebugMessageCallback");
api->GetDebugMessageLog = loadProc("glGetDebugMessageLog");
api->PushDebugGroup = loadProc("glPushDebugGroup");
api->PopDebugGroup = loadProc("glPopDebugGroup");
api->ObjectLabel = loadProc("glObjectLabel");
api->GetObjectLabel = loadProc("glGetObjectLabel");
api->ObjectPtrLabel = loadProc("glObjectPtrLabel");
api->GetObjectPtrLabel = loadProc("glGetObjectPtrLabel");
api->GetPointerv = loadProc("glGetPointerv");
api->BufferStorage = loadProc("glBufferStorage");
api->ClearTexImage = loadProc("glClearTexImage");
api->ClearTexSubImage = loadProc("glClearTexSubImage");
api->BindBuffersBase = loadProc("glBindBuffersBase");
api->BindBuffersRange = loadProc("glBindBuffersRange");
api->BindTextures = loadProc("glBindTextures");
api->BindSamplers = loadProc("glBindSamplers");
api->BindImageTextures = loadProc("glBindImageTextures");
api->BindVertexBuffers = loadProc("glBindVertexBuffers");
}
void mg_gl_load_gles30(mg_gl_api* api, mg_gl_load_proc loadProc) void mg_gl_load_gles30(mg_gl_api* api, mg_gl_load_proc loadProc)
{ {
api->ActiveTexture = loadProc("glActiveTexture"); api->ActiveTexture = loadProc("glActiveTexture");

3
src/gl_loader.h
View File

@ -2,7 +2,7 @@
* *
* @file: gl_loader.h * @file: gl_loader.h
* @note: auto-generated by glapi.py from gl.xml * @note: auto-generated by glapi.py from gl.xml
* @date: 22/022023 * @date: 12/072023
* *
*********************************************************/ *********************************************************/
#ifndef __GL_LOADER_H__ #ifndef __GL_LOADER_H__
@ -14,6 +14,7 @@ typedef void*(*mg_gl_load_proc)(const char* name);
void mg_gl_load_gl41(mg_gl_api* api, mg_gl_load_proc loadProc); void mg_gl_load_gl41(mg_gl_api* api, mg_gl_load_proc loadProc);
void mg_gl_load_gl43(mg_gl_api* api, mg_gl_load_proc loadProc); void mg_gl_load_gl43(mg_gl_api* api, mg_gl_load_proc loadProc);
void mg_gl_load_gl44(mg_gl_api* api, mg_gl_load_proc loadProc);
void mg_gl_load_gles30(mg_gl_api* api, mg_gl_load_proc loadProc); void mg_gl_load_gles30(mg_gl_api* api, mg_gl_load_proc loadProc);
void mg_gl_load_gles31(mg_gl_api* api, mg_gl_load_proc loadProc); void mg_gl_load_gles31(mg_gl_api* api, mg_gl_load_proc loadProc);

4
src/glsl_shaders/segment_setup.glsl
View File

@ -340,7 +340,7 @@ int quadratic_roots_with_det(float a, float b, float c, float det, out float r[2
if(a == 0) if(a == 0)
{ {
if(b) if(b != 0)
{ {
count = 1; count = 1;
r[0] = -c/b; r[0] = -c/b;
@ -773,7 +773,7 @@ void cubic_setup(vec2 p[4], int pathIndex)
int rootCount = rootCountX + rootCountY; int rootCount = rootCountX + rootCountY;
for(int i=0; i<2; i++) for(int i=0; i<2; i++)
{ {
if(curve.ts[i].y) if(curve.ts[i].y != 0)
{ {
roots[rootCount] = curve.ts[i].x / curve.ts[i].y; roots[rootCount] = curve.ts[i].x / curve.ts[i].y;
rootCount++; rootCount++;

1
src/mtl_renderer.h
View File

@ -34,7 +34,6 @@ typedef enum {
typedef struct mg_mtl_path_elt typedef struct mg_mtl_path_elt
{ {
int pathIndex; int pathIndex;
int localEltIndex;
mg_mtl_seg_kind kind; mg_mtl_seg_kind kind;
vector_float2 p[4]; vector_float2 p[4];
} mg_mtl_path_elt; } mg_mtl_path_elt;

513
src/mtl_renderer.m
View File

@ -17,7 +17,7 @@
#include"mtl_renderer.h" #include"mtl_renderer.h"
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 = 32,
MG_MTL_MSAA_COUNT = 8; MG_MTL_MSAA_COUNT = 8;
typedef struct mg_mtl_canvas_backend typedef struct mg_mtl_canvas_backend
@ -34,8 +34,6 @@ typedef struct mg_mtl_canvas_backend
id<MTLTexture> outTexture; id<MTLTexture> outTexture;
int pathBufferOffset;
int elementBufferOffset;
int bufferIndex; int bufferIndex;
dispatch_semaphore_t bufferSemaphore; dispatch_semaphore_t bufferSemaphore;
@ -57,6 +55,22 @@ typedef struct mg_mtl_canvas_backend
int msaaCount; int msaaCount;
vec2 frameSize; vec2 frameSize;
// encoding context
int eltCap;
int eltCount;
int eltBatchStart;
int pathCap;
int pathCount;
int pathBatchStart;
mg_primitive* primitive;
vec4 pathScreenExtents;
vec4 pathUserExtents;
int maxTileQueueCount;
int maxSegmentCount;
} mg_mtl_canvas_backend; } mg_mtl_canvas_backend;
typedef struct mg_mtl_image_data typedef struct mg_mtl_image_data
@ -84,20 +98,6 @@ void mg_mtl_print_log(int bufferIndex, id<MTLBuffer> logBuffer, id<MTLBuffer> lo
} }
} }
typedef struct mg_mtl_encoding_context
{
int mtlEltCount;
mg_mtl_path* pathBufferData;
mg_mtl_path_elt* elementBufferData;
int pathIndex;
int localEltIndex;
mg_primitive* primitive;
vec4 pathScreenExtents;
vec4 pathUserExtents;
} mg_mtl_encoding_context;
static void mg_update_path_extents(vec4* extents, vec2 p) static void mg_update_path_extents(vec4* extents, vec2 p)
{ {
extents->x = minimum(extents->x, p.x); extents->x = minimum(extents->x, p.x);
@ -106,27 +106,62 @@ static void mg_update_path_extents(vec4* extents, vec2 p)
extents->w = maximum(extents->w, p.y); extents->w = maximum(extents->w, p.y);
} }
void mg_mtl_canvas_encode_element(mg_mtl_encoding_context* context, mg_path_elt_type kind, vec2* p) id<MTLBuffer> mg_mtl_grow_input_buffer(id<MTLDevice> device, id<MTLBuffer> oldBuffer, int oldCopySize, int newSize)
{ {
mg_mtl_path_elt* mtlElt = &context->elementBufferData[context->mtlEltCount]; @autoreleasepool
context->mtlEltCount++; {
MTLResourceOptions bufferOptions = MTLResourceCPUCacheModeWriteCombined
| MTLResourceStorageModeShared;
mtlElt->pathIndex = context->pathIndex; id<MTLBuffer> newBuffer = [device newBufferWithLength: newSize options: bufferOptions];
memcpy([newBuffer contents], [oldBuffer contents], oldCopySize);
[oldBuffer release];
return(newBuffer);
}
}
void mg_mtl_canvas_encode_element(mg_mtl_canvas_backend* backend, mg_path_elt_type kind, vec2* p)
{
int bufferIndex = backend->bufferIndex;
int bufferCap = [backend->elementBuffer[bufferIndex] length] / sizeof(mg_mtl_path_elt);
if(backend->eltCount >= bufferCap)
{
int newBufferCap = (int)(bufferCap * 1.5);
int newBufferSize = newBufferCap * sizeof(mg_mtl_path_elt);
log_info("growing element buffer to %i elements\n", newBufferCap);
backend->elementBuffer[bufferIndex] = mg_mtl_grow_input_buffer(backend->surface->device,
backend->elementBuffer[bufferIndex],
backend->eltCount * sizeof(mg_mtl_path_elt),
newBufferSize);
}
mg_mtl_path_elt* elements = (mg_mtl_path_elt*)[backend->elementBuffer[bufferIndex] contents];
mg_mtl_path_elt* elt = &elements[backend->eltCount];
backend->eltCount++;
elt->pathIndex = backend->pathCount - backend->pathBatchStart;
int count = 0; int count = 0;
switch(kind) switch(kind)
{ {
case MG_PATH_LINE: case MG_PATH_LINE:
mtlElt->kind = MG_MTL_LINE; backend->maxSegmentCount += 1;
elt->kind = MG_MTL_LINE;
count = 2; count = 2;
break; break;
case MG_PATH_QUADRATIC: case MG_PATH_QUADRATIC:
mtlElt->kind = MG_MTL_QUADRATIC; backend->maxSegmentCount += 3;
elt->kind = MG_MTL_QUADRATIC;
count = 3; count = 3;
break; break;
case MG_PATH_CUBIC: case MG_PATH_CUBIC:
mtlElt->kind = MG_MTL_CUBIC; backend->maxSegmentCount += 7;
elt->kind = MG_MTL_CUBIC;
count = 4; count = 4;
break; break;
@ -134,20 +169,93 @@ void mg_mtl_canvas_encode_element(mg_mtl_encoding_context* context, mg_path_elt_
break; break;
} }
mtlElt->localEltIndex = context->localEltIndex;
for(int i=0; i<count; i++) for(int i=0; i<count; i++)
{ {
mg_update_path_extents(&context->pathUserExtents, p[i]); mg_update_path_extents(&backend->pathUserExtents, p[i]);
vec2 screenP = mg_mat2x3_mul(context->primitive->attributes.transform, p[i]); vec2 screenP = mg_mat2x3_mul(backend->primitive->attributes.transform, p[i]);
mtlElt->p[i] = (vector_float2){screenP.x, screenP.y}; elt->p[i] = (vector_float2){screenP.x, screenP.y};
mg_update_path_extents(&context->pathScreenExtents, screenP); mg_update_path_extents(&backend->pathScreenExtents, screenP);
} }
} }
void mg_mtl_encode_path(mg_mtl_canvas_backend* backend, mg_primitive* primitive, float scale)
{
int bufferIndex = backend->bufferIndex;
int bufferCap = [backend->pathBuffer[bufferIndex] length] / sizeof(mg_mtl_path);
if(backend->pathCount >= bufferCap)
{
int newBufferCap = (int)(bufferCap * 1.5);
int newBufferSize = newBufferCap * sizeof(mg_mtl_path);
log_info("growing path buffer to %i elements\n", newBufferCap);
backend->pathBuffer[bufferIndex] = mg_mtl_grow_input_buffer(backend->surface->device,
backend->pathBuffer[bufferIndex],
backend->eltCount * sizeof(mg_mtl_path),
newBufferSize);
}
mg_mtl_path* pathBufferData = (mg_mtl_path*)[backend->pathBuffer[backend->bufferIndex] contents];
mg_mtl_path* path = &(pathBufferData[backend->pathCount]);
backend->pathCount++;
path->cmd = (mg_mtl_cmd)primitive->cmd;
path->box = (vector_float4){backend->pathScreenExtents.x,
backend->pathScreenExtents.y,
backend->pathScreenExtents.z,
backend->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,
primitive->attributes.color.b,
primitive->attributes.color.a};
mp_rect srcRegion = primitive->attributes.srcRegion;
mp_rect destRegion = {backend->pathUserExtents.x,
backend->pathUserExtents.y,
backend->pathUserExtents.z - backend->pathUserExtents.x,
backend->pathUserExtents.w - backend->pathUserExtents.y};
if(!mg_image_is_nil(primitive->attributes.image))
{
vec2 texSize = mg_image_size(primitive->attributes.image);
mg_mat2x3 srcRegionToImage = {1/texSize.x, 0, srcRegion.x/texSize.x,
0, 1/texSize.y, srcRegion.y/texSize.y};
mg_mat2x3 destRegionToSrcRegion = {srcRegion.w/destRegion.w, 0, 0,
0, srcRegion.h/destRegion.h, 0};
mg_mat2x3 userToDestRegion = {1, 0, -destRegion.x,
0, 1, -destRegion.y};
mg_mat2x3 screenToUser = mg_mat2x3_inv(primitive->attributes.transform);
mg_mat2x3 uvTransform = srcRegionToImage;
uvTransform = mg_mat2x3_mul_m(uvTransform, destRegionToSrcRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, userToDestRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, screenToUser);
path->uvTransform = simd_matrix(simd_make_float3(uvTransform.m[0]/scale, uvTransform.m[3]/scale, 0),
simd_make_float3(uvTransform.m[1]/scale, uvTransform.m[4]/scale, 0),
simd_make_float3(uvTransform.m[2], uvTransform.m[5], 1));
}
int nTilesX = ((path->box.z - path->box.x)*scale - 1) / MG_MTL_TILE_SIZE + 1;
int nTilesY = ((path->box.w - path->box.y)*scale - 1) / MG_MTL_TILE_SIZE + 1;
backend->maxTileQueueCount += (nTilesX * nTilesY);
}
bool mg_intersect_hull_legs(vec2 p0, vec2 p1, vec2 p2, vec2 p3, vec2* intersection) bool mg_intersect_hull_legs(vec2 p0, vec2 p1, vec2 p2, vec2 p3, vec2* intersection)
{ {
/*NOTE: check intersection of lines (p0-p1) and (p2-p3) /*NOTE: check intersection of lines (p0-p1) and (p2-p3)
@ -336,9 +444,9 @@ void mg_cubic_split(vec2 p[4], f32 t, vec2 outLeft[4], vec2 outRight[4])
outRight[3] = p[3]; outRight[3] = p[3];
} }
void mg_mtl_render_stroke_line(mg_mtl_encoding_context* context, vec2* p) void mg_mtl_render_stroke_line(mg_mtl_canvas_backend* backend, vec2* p)
{ {
f32 width = context->primitive->attributes.width; f32 width = backend->primitive->attributes.width;
vec2 v = {p[1].x-p[0].x, p[1].y-p[0].y}; vec2 v = {p[1].x-p[0].x, p[1].y-p[0].y};
vec2 n = {v.y, -v.x}; vec2 n = {v.y, -v.x};
@ -350,28 +458,27 @@ void mg_mtl_render_stroke_line(mg_mtl_encoding_context* context, vec2* p)
vec2 joint0[2] = {vec2_add(p[0], vec2_mul(-1, offset)), vec2_add(p[0], 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))}; 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(backend, MG_PATH_LINE, right);
mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, left);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, left); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, joint0);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint0); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, joint1);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint1);
} }
void mg_mtl_render_stroke_quadratic(mg_mtl_encoding_context* context, vec2* p) void mg_mtl_render_stroke_quadratic(mg_mtl_canvas_backend* backend, vec2* p)
{ {
f32 width = context->primitive->attributes.width; f32 width = backend->primitive->attributes.width;
f32 tolerance = minimum(context->primitive->attributes.tolerance, 0.5 * width); f32 tolerance = minimum(backend->primitive->attributes.tolerance, 0.5 * width);
//NOTE: check for degenerate line case //NOTE: check for degenerate line case
const f32 equalEps = 1e-3; const f32 equalEps = 1e-3;
if(vec2_close(p[0], p[1], equalEps)) if(vec2_close(p[0], p[1], equalEps))
{ {
mg_mtl_render_stroke_line(context, p+1); mg_mtl_render_stroke_line(backend, p+1);
return; return;
} }
else if(vec2_close(p[1], p[2], equalEps)) else if(vec2_close(p[1], p[2], equalEps))
{ {
mg_mtl_render_stroke_line(context, p); mg_mtl_render_stroke_line(backend, p);
return; return;
} }
@ -386,8 +493,8 @@ void mg_mtl_render_stroke_quadratic(mg_mtl_encoding_context* context, vec2* p)
vec2 splitLeft[3]; vec2 splitLeft[3];
vec2 splitRight[3]; vec2 splitRight[3];
mg_quadratic_split(p, 0.5, splitLeft, splitRight); mg_quadratic_split(p, 0.5, splitLeft, splitRight);
mg_mtl_render_stroke_quadratic(context, splitLeft); mg_mtl_render_stroke_quadratic(backend, splitLeft);
mg_mtl_render_stroke_quadratic(context, splitRight); mg_mtl_render_stroke_quadratic(backend, splitRight);
} }
else else
{ {
@ -428,8 +535,8 @@ void mg_mtl_render_stroke_quadratic(mg_mtl_encoding_context* context, vec2* p)
vec2 splitLeft[3]; vec2 splitLeft[3];
vec2 splitRight[3]; vec2 splitRight[3];
mg_quadratic_split(p, maxOvershootParameter, splitLeft, splitRight); mg_quadratic_split(p, maxOvershootParameter, splitLeft, splitRight);
mg_mtl_render_stroke_quadratic(context, splitLeft); mg_mtl_render_stroke_quadratic(backend, splitLeft);
mg_mtl_render_stroke_quadratic(context, splitRight); mg_mtl_render_stroke_quadratic(backend, splitRight);
} }
else else
{ {
@ -437,21 +544,21 @@ void mg_mtl_render_stroke_quadratic(mg_mtl_encoding_context* context, vec2* p)
leftHull[0] = leftHull[2]; leftHull[0] = leftHull[2];
leftHull[2] = tmp; leftHull[2] = tmp;
mg_mtl_canvas_encode_element(context, MG_PATH_QUADRATIC, rightHull); mg_mtl_canvas_encode_element(backend, MG_PATH_QUADRATIC, rightHull);
mg_mtl_canvas_encode_element(context, MG_PATH_QUADRATIC, leftHull); mg_mtl_canvas_encode_element(backend, MG_PATH_QUADRATIC, leftHull);
vec2 joint0[2] = {rightHull[2], leftHull[0]}; vec2 joint0[2] = {rightHull[2], leftHull[0]};
vec2 joint1[2] = {leftHull[2], rightHull[0]}; vec2 joint1[2] = {leftHull[2], rightHull[0]};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint0); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, joint0);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint1); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, joint1);
} }
} }
} }
void mg_mtl_render_stroke_cubic(mg_mtl_encoding_context* context, vec2* p) void mg_mtl_render_stroke_cubic(mg_mtl_canvas_backend* backend, vec2* p)
{ {
f32 width = context->primitive->attributes.width; f32 width = backend->primitive->attributes.width;
f32 tolerance = minimum(context->primitive->attributes.tolerance, 0.5 * width); f32 tolerance = minimum(backend->primitive->attributes.tolerance, 0.5 * width);
//NOTE: check degenerate line cases //NOTE: check degenerate line cases
f32 equalEps = 1e-3; f32 equalEps = 1e-3;
@ -461,19 +568,19 @@ void mg_mtl_render_stroke_cubic(mg_mtl_encoding_context* context, vec2* p)
||(vec2_close(p[1], p[2], equalEps) && vec2_close(p[2], p[3], equalEps))) ||(vec2_close(p[1], p[2], equalEps) && vec2_close(p[2], p[3], equalEps)))
{ {
vec2 line[2] = {p[0], p[3]}; vec2 line[2] = {p[0], p[3]};
mg_mtl_render_stroke_line(context, line); mg_mtl_render_stroke_line(backend, line);
return; return;
} }
else if(vec2_close(p[0], p[1], equalEps) && vec2_close(p[1], p[3], equalEps)) else if(vec2_close(p[0], p[1], equalEps) && vec2_close(p[1], p[3], equalEps))
{ {
vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[2]))}; vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[2]))};
mg_mtl_render_stroke_line(context, line); mg_mtl_render_stroke_line(backend, line);
return; return;
} }
else if(vec2_close(p[0], p[2], equalEps) && vec2_close(p[2], p[3], equalEps)) else if(vec2_close(p[0], p[2], equalEps) && vec2_close(p[2], p[3], equalEps))
{ {
vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[1]))}; vec2 line[2] = {p[0], vec2_add(vec2_mul(5./9, p[0]), vec2_mul(4./9, p[1]))};
mg_mtl_render_stroke_line(context, line); mg_mtl_render_stroke_line(backend, line);
return; return;
} }
@ -488,8 +595,8 @@ void mg_mtl_render_stroke_cubic(mg_mtl_encoding_context* context, vec2* p)
vec2 splitLeft[4]; vec2 splitLeft[4];
vec2 splitRight[4]; vec2 splitRight[4];
mg_cubic_split(p, 0.5, splitLeft, splitRight); mg_cubic_split(p, 0.5, splitLeft, splitRight);
mg_mtl_render_stroke_cubic(context, splitLeft); mg_mtl_render_stroke_cubic(backend, splitLeft);
mg_mtl_render_stroke_cubic(context, splitRight); mg_mtl_render_stroke_cubic(backend, splitRight);
} }
else else
{ {
@ -530,8 +637,8 @@ void mg_mtl_render_stroke_cubic(mg_mtl_encoding_context* context, vec2* p)
vec2 splitLeft[4]; vec2 splitLeft[4];
vec2 splitRight[4]; vec2 splitRight[4];
mg_cubic_split(p, maxOvershootParameter, splitLeft, splitRight); mg_cubic_split(p, maxOvershootParameter, splitLeft, splitRight);
mg_mtl_render_stroke_cubic(context, splitLeft); mg_mtl_render_stroke_cubic(backend, splitLeft);
mg_mtl_render_stroke_cubic(context, splitRight); mg_mtl_render_stroke_cubic(backend, splitRight);
} }
else else
{ {
@ -542,18 +649,18 @@ void mg_mtl_render_stroke_cubic(mg_mtl_encoding_context* context, vec2* p)
leftHull[1] = leftHull[2]; leftHull[1] = leftHull[2];
leftHull[2] = tmp; leftHull[2] = tmp;
mg_mtl_canvas_encode_element(context, MG_PATH_CUBIC, rightHull); mg_mtl_canvas_encode_element(backend, MG_PATH_CUBIC, rightHull);
mg_mtl_canvas_encode_element(context, MG_PATH_CUBIC, leftHull); mg_mtl_canvas_encode_element(backend, MG_PATH_CUBIC, leftHull);
vec2 joint0[2] = {rightHull[3], leftHull[0]}; vec2 joint0[2] = {rightHull[3], leftHull[0]};
vec2 joint1[2] = {leftHull[3], rightHull[0]}; vec2 joint1[2] = {leftHull[3], rightHull[0]};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint0); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, joint0);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, joint1); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, joint1);
} }
} }
} }
void mg_mtl_render_stroke_element(mg_mtl_encoding_context* context, void mg_mtl_render_stroke_element(mg_mtl_canvas_backend* backend,
mg_path_elt* element, mg_path_elt* element,
vec2 currentPoint, vec2 currentPoint,
vec2* startTangent, vec2* startTangent,
@ -566,17 +673,17 @@ void mg_mtl_render_stroke_element(mg_mtl_encoding_context* context,
switch(element->type) switch(element->type)
{ {
case MG_PATH_LINE: case MG_PATH_LINE:
mg_mtl_render_stroke_line(context, controlPoints); mg_mtl_render_stroke_line(backend, controlPoints);
endPointIndex = 1; endPointIndex = 1;
break; break;
case MG_PATH_QUADRATIC: case MG_PATH_QUADRATIC:
mg_mtl_render_stroke_quadratic(context, controlPoints); mg_mtl_render_stroke_quadratic(backend, controlPoints);
endPointIndex = 2; endPointIndex = 2;
break; break;
case MG_PATH_CUBIC: case MG_PATH_CUBIC:
mg_mtl_render_stroke_cubic(context, controlPoints); mg_mtl_render_stroke_cubic(backend, controlPoints);
endPointIndex = 3; endPointIndex = 3;
break; break;
@ -613,11 +720,11 @@ void mg_mtl_render_stroke_element(mg_mtl_encoding_context* context,
DEBUG_ASSERT(startTangent->x != 0 || startTangent->y != 0); DEBUG_ASSERT(startTangent->x != 0 || startTangent->y != 0);
} }
void mg_mtl_stroke_cap(mg_mtl_encoding_context* context, void mg_mtl_stroke_cap(mg_mtl_canvas_backend* backend,
vec2 p0, vec2 p0,
vec2 direction) vec2 direction)
{ {
mg_attributes* attributes = &context->primitive->attributes; mg_attributes* attributes = &backend->primitive->attributes;
//NOTE(martin): compute the tangent and normal vectors (multiplied by half width) at the cap point //NOTE(martin): compute the tangent and normal vectors (multiplied by half width) at the cap point
f32 dn = sqrt(Square(direction.x) + Square(direction.y)); f32 dn = sqrt(Square(direction.x) + Square(direction.y));
@ -635,18 +742,18 @@ void mg_mtl_stroke_cap(mg_mtl_encoding_context* context,
{p0.x - n0.x, p0.y - n0.y}, {p0.x - n0.x, p0.y - n0.y},
{p0.x + n0.x, p0.y + n0.y}}; {p0.x + n0.x, p0.y + n0.y}};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+1); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+1);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+2); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+2);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+3); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+3);
} }
void mg_mtl_stroke_joint(mg_mtl_encoding_context* context, void mg_mtl_stroke_joint(mg_mtl_canvas_backend* backend,
vec2 p0, vec2 p0,
vec2 t0, vec2 t0,
vec2 t1) vec2 t1)
{ {
mg_attributes* attributes = &context->primitive->attributes; mg_attributes* attributes = &backend->primitive->attributes;
//NOTE(martin): compute the normals at the joint point //NOTE(martin): compute the normals at the joint point
f32 norm_t0 = sqrt(Square(t0.x) + Square(t0.y)); f32 norm_t0 = sqrt(Square(t0.x) + Square(t0.y));
@ -694,10 +801,10 @@ void mg_mtl_stroke_joint(mg_mtl_encoding_context* context,
{p0.x + n1.x*halfW, p0.y + n1.y*halfW}, {p0.x + n1.x*halfW, p0.y + n1.y*halfW},
p0}; p0};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+1); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+1);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+2); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+2);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+3); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+3);
} }
else else
{ {
@ -707,13 +814,13 @@ void mg_mtl_stroke_joint(mg_mtl_encoding_context* context,
{p0.x + n1.x*halfW, p0.y + n1.y*halfW}, {p0.x + n1.x*halfW, p0.y + n1.y*halfW},
p0}; p0};
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+1); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+1);
mg_mtl_canvas_encode_element(context, MG_PATH_LINE, points+2); mg_mtl_canvas_encode_element(backend, MG_PATH_LINE, points+2);
} }
} }
u32 mg_mtl_render_stroke_subpath(mg_mtl_encoding_context* context, u32 mg_mtl_render_stroke_subpath(mg_mtl_canvas_backend* backend,
mg_path_elt* elements, mg_path_elt* elements,
mg_path_descriptor* path, mg_path_descriptor* path,
u32 startIndex, u32 startIndex,
@ -730,7 +837,7 @@ u32 mg_mtl_render_stroke_subpath(mg_mtl_encoding_context* context,
vec2 endTangent = {0, 0}; vec2 endTangent = {0, 0};
//NOTE(martin): render first element and compute first tangent //NOTE(martin): render first element and compute first tangent
mg_mtl_render_stroke_element(context, elements + startIndex, currentPoint, &startTangent, &endTangent, &endPoint); mg_mtl_render_stroke_element(backend, elements + startIndex, currentPoint, &startTangent, &endTangent, &endPoint);
firstTangent = startTangent; firstTangent = startTangent;
previousEndTangent = endTangent; previousEndTangent = endTangent;
@ -738,18 +845,18 @@ u32 mg_mtl_render_stroke_subpath(mg_mtl_encoding_context* context,
//NOTE(martin): render subsequent elements along with their joints //NOTE(martin): render subsequent elements along with their joints
mg_attributes* attributes = &context->primitive->attributes; mg_attributes* attributes = &backend->primitive->attributes;
u32 eltIndex = startIndex + 1; u32 eltIndex = startIndex + 1;
for(; for(;
eltIndex<eltCount && elements[eltIndex].type != MG_PATH_MOVE; eltIndex<eltCount && elements[eltIndex].type != MG_PATH_MOVE;
eltIndex++) eltIndex++)
{ {
mg_mtl_render_stroke_element(context, elements + eltIndex, currentPoint, &startTangent, &endTangent, &endPoint); mg_mtl_render_stroke_element(backend, elements + eltIndex, currentPoint, &startTangent, &endTangent, &endPoint);
if(attributes->joint != MG_JOINT_NONE) if(attributes->joint != MG_JOINT_NONE)
{ {
mg_mtl_stroke_joint(context, currentPoint, previousEndTangent, startTangent); mg_mtl_stroke_joint(backend, currentPoint, previousEndTangent, startTangent);
} }
previousEndTangent = endTangent; previousEndTangent = endTangent;
currentPoint = endPoint; currentPoint = endPoint;
@ -764,19 +871,19 @@ u32 mg_mtl_render_stroke_subpath(mg_mtl_encoding_context* context,
if(attributes->joint != MG_JOINT_NONE) if(attributes->joint != MG_JOINT_NONE)
{ {
//NOTE(martin): add a closing joint if the path is closed //NOTE(martin): add a closing joint if the path is closed
mg_mtl_stroke_joint(context, endPoint, endTangent, firstTangent); mg_mtl_stroke_joint(backend, endPoint, endTangent, firstTangent);
} }
} }
else if(attributes->cap == MG_CAP_SQUARE) else if(attributes->cap == MG_CAP_SQUARE)
{ {
//NOTE(martin): add start and end cap //NOTE(martin): add start and end cap
mg_mtl_stroke_cap(context, startPoint, (vec2){-startTangent.x, -startTangent.y}); mg_mtl_stroke_cap(backend, startPoint, (vec2){-startTangent.x, -startTangent.y});
mg_mtl_stroke_cap(context, endPoint, endTangent); mg_mtl_stroke_cap(backend, endPoint, endTangent);
} }
return(eltIndex); return(eltIndex);
} }
void mg_mtl_render_stroke(mg_mtl_encoding_context* context, void mg_mtl_render_stroke(mg_mtl_canvas_backend* backend,
mg_path_elt* elements, mg_path_elt* elements,
mg_path_descriptor* path) mg_path_descriptor* path)
{ {
@ -796,16 +903,35 @@ void mg_mtl_render_stroke(mg_mtl_encoding_context* context,
} }
if(startIndex < eltCount) if(startIndex < eltCount)
{ {
startIndex = mg_mtl_render_stroke_subpath(context, elements, path, startIndex, startPoint); startIndex = mg_mtl_render_stroke_subpath(backend, elements, path, startIndex, startPoint);
} }
} }
} }
void mg_mtl_grow_buffer_if_needed(mg_mtl_canvas_backend* backend, id<MTLBuffer>* buffer, u64 wantedSize)
{
u64 bufferSize = [(*buffer) length];
if(bufferSize < wantedSize)
{
int newSize = wantedSize * 1.2;
@autoreleasepool
{
//NOTE: MTLBuffers are retained by the command buffer, so we don't risk deallocating while the buffer is in use
[*buffer release];
*buffer = nil;
id<MTLDevice> device = backend->surface->device;
MTLResourceOptions bufferOptions = MTLResourceStorageModePrivate;
*buffer = [device newBufferWithLength: newSize options: bufferOptions];
}
}
}
void mg_mtl_render_batch(mg_mtl_canvas_backend* backend, void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
mg_mtl_surface* surface, mg_mtl_surface* surface,
int pathCount,
int eltCount,
mg_image_data* image, mg_image_data* image,
int tileSize, int tileSize,
int nTilesX, int nTilesX,
@ -813,10 +939,23 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
vec2 viewportSize, vec2 viewportSize,
f32 scale) f32 scale)
{ {
int pathBufferOffset = backend->pathBatchStart * sizeof(mg_mtl_path);
int elementBufferOffset = backend->eltBatchStart * sizeof(mg_mtl_path_elt);
int pathCount = backend->pathCount - backend->pathBatchStart;
int eltCount = backend->eltCount - backend->eltBatchStart;
//NOTE: update intermediate buffers sizes if needed
mg_mtl_grow_buffer_if_needed(backend, &backend->pathQueueBuffer, pathCount * sizeof(mg_mtl_path_queue));
mg_mtl_grow_buffer_if_needed(backend, &backend->tileQueueBuffer, backend->maxTileQueueCount * sizeof(mg_mtl_tile_queue));
mg_mtl_grow_buffer_if_needed(backend, &backend->segmentBuffer, backend->maxSegmentCount * sizeof(mg_mtl_segment));
mg_mtl_grow_buffer_if_needed(backend, &backend->screenTilesBuffer, nTilesX * nTilesY * sizeof(mg_mtl_screen_tile));
mg_mtl_grow_buffer_if_needed(backend, &backend->tileOpBuffer, backend->maxSegmentCount * 30 * sizeof(mg_mtl_tile_op));
//NOTE: encode GPU commands //NOTE: encode GPU commands
@autoreleasepool @autoreleasepool
{ {
//NOTE: create output texture //NOTE: clear output texture
MTLRenderPassDescriptor* clearDescriptor = [MTLRenderPassDescriptor renderPassDescriptor]; MTLRenderPassDescriptor* clearDescriptor = [MTLRenderPassDescriptor renderPassDescriptor];
clearDescriptor.colorAttachments[0].texture = backend->outTexture; clearDescriptor.colorAttachments[0].texture = backend->outTexture;
clearDescriptor.colorAttachments[0].loadAction = MTLLoadActionClear; clearDescriptor.colorAttachments[0].loadAction = MTLLoadActionClear;
@ -841,13 +980,16 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
pathEncoder.label = @"path pass"; pathEncoder.label = @"path pass";
[pathEncoder setComputePipelineState: backend->pathPipeline]; [pathEncoder setComputePipelineState: backend->pathPipeline];
int tileQueueMax = [backend->tileQueueBuffer length] / sizeof(mg_mtl_tile_queue);
[pathEncoder setBytes:&pathCount length:sizeof(int) atIndex:0]; [pathEncoder setBytes:&pathCount length:sizeof(int) atIndex:0];
[pathEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:backend->pathBufferOffset atIndex:1]; [pathEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:pathBufferOffset atIndex:1];
[pathEncoder setBuffer:backend->pathQueueBuffer offset:0 atIndex:2]; [pathEncoder setBuffer:backend->pathQueueBuffer offset:0 atIndex:2];
[pathEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:3]; [pathEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:3];
[pathEncoder setBuffer:backend->tileQueueCountBuffer offset:0 atIndex:4]; [pathEncoder setBuffer:backend->tileQueueCountBuffer offset:0 atIndex:4];
[pathEncoder setBytes:&tileSize length:sizeof(int) atIndex:5]; [pathEncoder setBytes:&tileQueueMax length:sizeof(int) atIndex:5];
[pathEncoder setBytes:&scale length:sizeof(int) atIndex:6]; [pathEncoder setBytes:&tileSize length:sizeof(int) atIndex:6];
[pathEncoder setBytes:&scale length:sizeof(int) atIndex:7];
MTLSize pathGridSize = MTLSizeMake(pathCount, 1, 1); MTLSize pathGridSize = MTLSizeMake(pathCount, 1, 1);
MTLSize pathGroupSize = MTLSizeMake([backend->pathPipeline maxTotalThreadsPerThreadgroup], 1, 1); MTLSize pathGroupSize = MTLSizeMake([backend->pathPipeline maxTotalThreadsPerThreadgroup], 1, 1);
@ -860,18 +1002,23 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
segmentEncoder.label = @"segment pass"; segmentEncoder.label = @"segment pass";
[segmentEncoder setComputePipelineState: backend->segmentPipeline]; [segmentEncoder setComputePipelineState: backend->segmentPipeline];
int tileOpMax = [backend->tileOpBuffer length] / sizeof(mg_mtl_tile_op);
int segmentMax = [backend->segmentBuffer length] / sizeof(mg_mtl_segment);
[segmentEncoder setBytes:&eltCount length:sizeof(int) atIndex:0]; [segmentEncoder setBytes:&eltCount length:sizeof(int) atIndex:0];
[segmentEncoder setBuffer:backend->elementBuffer[backend->bufferIndex] offset:backend->elementBufferOffset atIndex:1]; [segmentEncoder setBuffer:backend->elementBuffer[backend->bufferIndex] offset:elementBufferOffset atIndex:1];
[segmentEncoder setBuffer:backend->segmentCountBuffer offset:0 atIndex:2]; [segmentEncoder setBuffer:backend->segmentCountBuffer offset:0 atIndex:2];
[segmentEncoder setBuffer:backend->segmentBuffer offset:0 atIndex:3]; [segmentEncoder setBuffer:backend->segmentBuffer offset:0 atIndex:3];
[segmentEncoder setBuffer:backend->pathQueueBuffer offset:0 atIndex:4]; [segmentEncoder setBuffer:backend->pathQueueBuffer offset:0 atIndex:4];
[segmentEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:5]; [segmentEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:5];
[segmentEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:6]; [segmentEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:6];
[segmentEncoder setBuffer:backend->tileOpCountBuffer offset:0 atIndex:7]; [segmentEncoder setBuffer:backend->tileOpCountBuffer offset:0 atIndex:7];
[segmentEncoder setBytes:&tileSize length:sizeof(int) atIndex:8]; [segmentEncoder setBytes:&tileOpMax length:sizeof(int) atIndex:8];
[segmentEncoder setBytes:&scale length:sizeof(int) atIndex:9]; [segmentEncoder setBytes:&segmentMax length:sizeof(int) atIndex:9];
[segmentEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:10]; [segmentEncoder setBytes:&tileSize length:sizeof(int) atIndex:10];
[segmentEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:11]; [segmentEncoder setBytes:&scale length:sizeof(int) atIndex:11];
[segmentEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:12];
[segmentEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:13];
MTLSize segmentGridSize = MTLSizeMake(eltCount, 1, 1); MTLSize segmentGridSize = MTLSizeMake(eltCount, 1, 1);
MTLSize segmentGroupSize = MTLSizeMake([backend->segmentPipeline maxTotalThreadsPerThreadgroup], 1, 1); MTLSize segmentGroupSize = MTLSizeMake([backend->segmentPipeline maxTotalThreadsPerThreadgroup], 1, 1);
@ -901,20 +1048,21 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
[mergeEncoder setComputePipelineState: backend->mergePipeline]; [mergeEncoder setComputePipelineState: backend->mergePipeline];
[mergeEncoder setBytes:&pathCount length:sizeof(int) atIndex:0]; [mergeEncoder setBytes:&pathCount length:sizeof(int) atIndex:0];
[mergeEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:backend->pathBufferOffset atIndex:1]; [mergeEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:pathBufferOffset atIndex:1];
[mergeEncoder setBuffer:backend->pathQueueBuffer offset:0 atIndex:2]; [mergeEncoder setBuffer:backend->pathQueueBuffer offset:0 atIndex:2];
[mergeEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:3]; [mergeEncoder setBuffer:backend->tileQueueBuffer offset:0 atIndex:3];
[mergeEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:4]; [mergeEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:4];
[mergeEncoder setBuffer:backend->tileOpCountBuffer offset:0 atIndex:5]; [mergeEncoder setBuffer:backend->tileOpCountBuffer offset:0 atIndex:5];
[mergeEncoder setBuffer:backend->rasterDispatchBuffer offset:0 atIndex:6]; [mergeEncoder setBuffer:backend->rasterDispatchBuffer offset:0 atIndex:6];
[mergeEncoder setBuffer:backend->screenTilesBuffer offset:0 atIndex:7]; [mergeEncoder setBuffer:backend->screenTilesBuffer offset:0 atIndex:7];
[mergeEncoder setBytes:&tileSize length:sizeof(int) atIndex:8]; [mergeEncoder setBytes:&tileOpMax length:sizeof(int) atIndex:8];
[mergeEncoder setBytes:&scale length:sizeof(float) atIndex:9]; [mergeEncoder setBytes:&tileSize length:sizeof(int) atIndex:9];
[mergeEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:10]; [mergeEncoder setBytes:&scale length:sizeof(float) atIndex:10];
[mergeEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:11]; [mergeEncoder setBuffer:backend->logBuffer[backend->bufferIndex] offset:0 atIndex:11];
[mergeEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:12];
MTLSize mergeGridSize = MTLSizeMake(nTilesX, nTilesY, 1); MTLSize mergeGridSize = MTLSizeMake(nTilesX, nTilesY, 1);
MTLSize mergeGroupSize = MTLSizeMake(16, 16, 1); MTLSize mergeGroupSize = MTLSizeMake(MG_MTL_TILE_SIZE, MG_MTL_TILE_SIZE, 1);
[mergeEncoder dispatchThreads: mergeGridSize threadsPerThreadgroup: mergeGroupSize]; [mergeEncoder dispatchThreads: mergeGridSize threadsPerThreadgroup: mergeGroupSize];
[mergeEncoder endEncoding]; [mergeEncoder endEncoding];
@ -926,7 +1074,7 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
[rasterEncoder setBuffer:backend->screenTilesBuffer offset:0 atIndex:0]; [rasterEncoder setBuffer:backend->screenTilesBuffer offset:0 atIndex:0];
[rasterEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:1]; [rasterEncoder setBuffer:backend->tileOpBuffer offset:0 atIndex:1];
[rasterEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:backend->pathBufferOffset atIndex:2]; [rasterEncoder setBuffer:backend->pathBuffer[backend->bufferIndex] offset:pathBufferOffset 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:&scale length:sizeof(float) atIndex:5]; [rasterEncoder setBytes:&scale length:sizeof(float) atIndex:5];
@ -946,8 +1094,7 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
[rasterEncoder setBytes: &useTexture length:sizeof(int) atIndex: 9]; [rasterEncoder setBytes: &useTexture length:sizeof(int) atIndex: 9];
MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1); MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1);
MTLSize rasterGroupSize = MTLSizeMake(16, 16, 1); MTLSize rasterGroupSize = MTLSizeMake(MG_MTL_TILE_SIZE, MG_MTL_TILE_SIZE, 1);
// [rasterEncoder dispatchThreads: rasterGridSize threadsPerThreadgroup: rasterGroupSize];
[rasterEncoder dispatchThreadgroupsWithIndirectBuffer: backend->rasterDispatchBuffer [rasterEncoder dispatchThreadgroupsWithIndirectBuffer: backend->rasterDispatchBuffer
indirectBufferOffset: 0 indirectBufferOffset: 0
@ -973,6 +1120,12 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
vertexCount: 3 ]; vertexCount: 3 ];
[renderEncoder endEncoding]; [renderEncoder endEncoding];
} }
backend->pathBatchStart = backend->pathCount;
backend->eltBatchStart = backend->eltCount;
backend->maxSegmentCount = 0;
backend->maxTileQueueCount = 0;
} }
void mg_mtl_canvas_resize(mg_mtl_canvas_backend* backend, vec2 size) void mg_mtl_canvas_resize(mg_mtl_canvas_backend* backend, vec2 size)
@ -999,7 +1152,7 @@ void mg_mtl_canvas_resize(mg_mtl_canvas_backend* backend, vec2 size)
MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init];
texDesc.textureType = MTLTextureType2D; texDesc.textureType = MTLTextureType2D;
texDesc.storageMode = MTLStorageModePrivate; texDesc.storageMode = MTLStorageModePrivate;
texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite; texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite | MTLTextureUsageRenderTarget;
texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm; texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm;
texDesc.width = size.x; texDesc.width = size.x;
texDesc.height = size.y; texDesc.height = size.y;
@ -1019,18 +1172,11 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
{ {
mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface;
//NOTE: update rolling buffers //NOTE: update rolling input buffers
dispatch_semaphore_wait(backend->bufferSemaphore, DISPATCH_TIME_FOREVER); dispatch_semaphore_wait(backend->bufferSemaphore, DISPATCH_TIME_FOREVER);
backend->bufferIndex = (backend->bufferIndex + 1) % MG_MTL_INPUT_BUFFERS_COUNT; backend->bufferIndex = (backend->bufferIndex + 1) % MG_MTL_INPUT_BUFFERS_COUNT;
mg_mtl_path_elt* elementBufferData = (mg_mtl_path_elt*)[backend->elementBuffer[backend->bufferIndex] contents]; //NOTE: ensure screen tiles buffer is correct size
mg_mtl_path* pathBufferData = (mg_mtl_path*)[backend->pathBuffer[backend->bufferIndex] contents];
/////////////////////////////////////////////////////////////////////////////////////
//TODO: ensure screen tiles buffer is correct size
/////////////////////////////////////////////////////////////////////////////////////
//NOTE: prepare rendering
mg_mtl_surface* surface = backend->surface; mg_mtl_surface* surface = backend->surface;
mp_rect frame = surface->interface.getFrame((mg_surface_data*)surface); mp_rect frame = surface->interface.getFrame((mg_surface_data*)surface);
@ -1046,6 +1192,7 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
mg_mtl_canvas_resize(backend, viewportSize); mg_mtl_canvas_resize(backend, viewportSize);
} }
//NOTE: acquire metal resources for rendering
mg_mtl_surface_acquire_command_buffer(surface); mg_mtl_surface_acquire_command_buffer(surface);
mg_mtl_surface_acquire_drawable(surface); mg_mtl_surface_acquire_drawable(surface);
@ -1068,17 +1215,17 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
renderEncoder.label = @"clear pass"; renderEncoder.label = @"clear pass";
[renderEncoder endEncoding]; [renderEncoder endEncoding];
} }
backend->pathBufferOffset = 0; backend->pathCount = 0;
backend->elementBufferOffset = 0; backend->pathBatchStart = 0;
backend->eltCount = 0;
backend->eltBatchStart = 0;
backend->maxSegmentCount = 0;
backend->maxTileQueueCount = 0;
//NOTE: encode and render batches //NOTE: encode and render batches
int pathCount = 0;
vec2 currentPos = {0}; vec2 currentPos = {0};
mg_image currentImage = mg_image_nil(); mg_image currentImage = mg_image_nil();
mg_mtl_encoding_context context = {.mtlEltCount = 0,
.elementBufferData = elementBufferData,
.pathBufferData = pathBufferData};
for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++) for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++)
{ {
@ -1090,51 +1237,37 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
mg_mtl_render_batch(backend, mg_mtl_render_batch(backend,
surface, surface,
pathCount,
context.mtlEltCount,
imageData, imageData,
tileSize, tileSize,
nTilesX, nTilesX,
nTilesY, nTilesY,
viewportSize, viewportSize,
scale); scale);
backend->pathBufferOffset += pathCount * sizeof(mg_mtl_path);
backend->elementBufferOffset += context.mtlEltCount * sizeof(mg_mtl_path_elt);
pathCount = 0;
context.mtlEltCount = 0;
context.elementBufferData = (mg_mtl_path_elt*)((char*)elementBufferData + backend->elementBufferOffset);
context.pathBufferData = (mg_mtl_path*)((char*)pathBufferData + backend->pathBufferOffset);
} }
currentImage = primitive->attributes.image; currentImage = primitive->attributes.image;
if(primitive->path.count) if(primitive->path.count)
{ {
context.primitive = primitive; backend->primitive = primitive;
context.pathIndex = pathCount; backend->pathScreenExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
context.pathScreenExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX}; backend->pathUserExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
context.pathUserExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX};
if(primitive->cmd == MG_CMD_STROKE) if(primitive->cmd == MG_CMD_STROKE)
{ {
mg_mtl_render_stroke(&context, pathElements + primitive->path.startIndex, &primitive->path); mg_mtl_render_stroke(backend, pathElements + primitive->path.startIndex, &primitive->path);
} }
else else
{ {
int segCount = 0;
for(int eltIndex = 0; for(int eltIndex = 0;
(eltIndex < primitive->path.count) && (primitive->path.startIndex + eltIndex < eltCount); (eltIndex < primitive->path.count) && (primitive->path.startIndex + eltIndex < eltCount);
eltIndex++) eltIndex++)
{ {
context.localEltIndex = segCount;
mg_path_elt* elt = &pathElements[primitive->path.startIndex + eltIndex]; mg_path_elt* elt = &pathElements[primitive->path.startIndex + eltIndex];
if(elt->type != MG_PATH_MOVE) if(elt->type != MG_PATH_MOVE)
{ {
vec2 p[4] = {currentPos, elt->p[0], elt->p[1], elt->p[2]}; vec2 p[4] = {currentPos, elt->p[0], elt->p[1], elt->p[2]};
mg_mtl_canvas_encode_element(&context, elt->type, p); mg_mtl_canvas_encode_element(backend, elt->type, p);
segCount++;
} }
switch(elt->type) switch(elt->type)
{ {
@ -1156,66 +1289,15 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
} }
} }
} }
//NOTE: push path //NOTE: encode path
mg_mtl_path* path = &context.pathBufferData[pathCount]; mg_mtl_encode_path(backend, primitive, scale);
pathCount++;
path->cmd = (mg_mtl_cmd)primitive->cmd;
path->box = (vector_float4){context.pathScreenExtents.x,
context.pathScreenExtents.y,
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,
primitive->attributes.color.b,
primitive->attributes.color.a};
mp_rect srcRegion = primitive->attributes.srcRegion;
mp_rect destRegion = {context.pathUserExtents.x,
context.pathUserExtents.y,
context.pathUserExtents.z - context.pathUserExtents.x,
context.pathUserExtents.w - context.pathUserExtents.y};
if(!mg_image_is_nil(primitive->attributes.image))
{
vec2 texSize = mg_image_size(primitive->attributes.image);
mg_mat2x3 srcRegionToImage = {1/texSize.x, 0, srcRegion.x/texSize.x,
0, 1/texSize.y, srcRegion.y/texSize.y};
mg_mat2x3 destRegionToSrcRegion = {srcRegion.w/destRegion.w, 0, 0,
0, srcRegion.h/destRegion.h, 0};
mg_mat2x3 userToDestRegion = {1, 0, -destRegion.x,
0, 1, -destRegion.y};
mg_mat2x3 screenToUser = mg_mat2x3_inv(primitive->attributes.transform);
mg_mat2x3 uvTransform = srcRegionToImage;
uvTransform = mg_mat2x3_mul_m(uvTransform, destRegionToSrcRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, userToDestRegion);
uvTransform = mg_mat2x3_mul_m(uvTransform, screenToUser);
path->uvTransform = simd_matrix(simd_make_float3(uvTransform.m[0]/scale, uvTransform.m[3]/scale, 0),
simd_make_float3(uvTransform.m[1]/scale, uvTransform.m[4]/scale, 0),
simd_make_float3(uvTransform.m[2], uvTransform.m[5], 1));
}
} }
} }
mg_image_data* imageData = mg_image_data_from_handle(currentImage); mg_image_data* imageData = mg_image_data_from_handle(currentImage);
mg_mtl_render_batch(backend, mg_mtl_render_batch(backend,
surface, surface,
pathCount,
context.mtlEltCount,
imageData, imageData,
tileSize, tileSize,
nTilesX, nTilesX,
@ -1321,12 +1403,13 @@ void mg_mtl_canvas_image_upload_region(mg_canvas_backend* backendInterface, mg_i
bytesPerRow: 4 * region.w]; bytesPerRow: 4 * region.w];
}} }}
const u32 MG_MTL_PATH_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_path), const u32 MG_MTL_DEFAULT_PATH_BUFFER_LEN = (4<<10),
MG_MTL_ELEMENT_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_path_elt), MG_MTL_DEFAULT_ELT_BUFFER_LEN = (4<<10),
MG_MTL_SEGMENT_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_segment),
MG_MTL_PATH_QUEUE_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_path_queue), MG_MTL_DEFAULT_SEGMENT_BUFFER_LEN = (4<<10),
MG_MTL_TILE_QUEUE_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_tile_queue), MG_MTL_DEFAULT_PATH_QUEUE_BUFFER_LEN = (4<<10),
MG_MTL_TILE_OP_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_tile_op); MG_MTL_DEFAULT_TILE_QUEUE_BUFFER_LEN = (4<<10),
MG_MTL_DEFAULT_TILE_OP_BUFFER_LEN = (4<<14);
mg_canvas_backend* mtl_canvas_backend_create(mg_mtl_surface* surface) mg_canvas_backend* mtl_canvas_backend_create(mg_mtl_surface* surface)
{ {
@ -1407,7 +1490,7 @@ mg_canvas_backend* mtl_canvas_backend_create(mg_mtl_surface* surface)
MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init];
texDesc.textureType = MTLTextureType2D; texDesc.textureType = MTLTextureType2D;
texDesc.storageMode = MTLStorageModePrivate; texDesc.storageMode = MTLStorageModePrivate;
texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite; texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite | MTLTextureUsageRenderTarget;
texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm; texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm;
texDesc.width = backend->frameSize.x; texDesc.width = backend->frameSize.x;
texDesc.height = backend->frameSize.y; texDesc.height = backend->frameSize.y;
@ -1424,30 +1507,30 @@ mg_canvas_backend* mtl_canvas_backend_create(mg_mtl_surface* surface)
for(int i=0; i<MG_MTL_INPUT_BUFFERS_COUNT; i++) for(int i=0; i<MG_MTL_INPUT_BUFFERS_COUNT; i++)
{ {
backend->pathBuffer[i] = [surface->device newBufferWithLength: MG_MTL_PATH_BUFFER_SIZE backend->pathBuffer[i] = [surface->device newBufferWithLength: MG_MTL_DEFAULT_PATH_BUFFER_LEN * sizeof(mg_mtl_path)
options: bufferOptions]; options: bufferOptions];
backend->elementBuffer[i] = [surface->device newBufferWithLength: MG_MTL_ELEMENT_BUFFER_SIZE backend->elementBuffer[i] = [surface->device newBufferWithLength: MG_MTL_DEFAULT_ELT_BUFFER_LEN * sizeof(mg_mtl_path_elt)
options: bufferOptions]; options: bufferOptions];
} }
bufferOptions = MTLResourceStorageModePrivate; bufferOptions = MTLResourceStorageModePrivate;
backend->segmentBuffer = [surface->device newBufferWithLength: MG_MTL_SEGMENT_BUFFER_SIZE backend->segmentBuffer = [surface->device newBufferWithLength: MG_MTL_DEFAULT_SEGMENT_BUFFER_LEN * sizeof(mg_mtl_segment)
options: bufferOptions]; options: bufferOptions];
backend->segmentCountBuffer = [surface->device newBufferWithLength: sizeof(int) backend->segmentCountBuffer = [surface->device newBufferWithLength: sizeof(int)
options: bufferOptions]; options: bufferOptions];
backend->pathQueueBuffer = [surface->device newBufferWithLength: MG_MTL_PATH_QUEUE_BUFFER_SIZE backend->pathQueueBuffer = [surface->device newBufferWithLength: MG_MTL_DEFAULT_PATH_QUEUE_BUFFER_LEN * sizeof(mg_mtl_path_queue)
options: bufferOptions]; options: bufferOptions];
backend->tileQueueBuffer = [surface->device newBufferWithLength: MG_MTL_TILE_QUEUE_BUFFER_SIZE backend->tileQueueBuffer = [surface->device newBufferWithLength: MG_MTL_DEFAULT_TILE_QUEUE_BUFFER_LEN * sizeof(mg_mtl_tile_queue)
options: bufferOptions]; options: bufferOptions];
backend->tileQueueCountBuffer = [surface->device newBufferWithLength: sizeof(int) backend->tileQueueCountBuffer = [surface->device newBufferWithLength: sizeof(int)
options: bufferOptions]; options: bufferOptions];
backend->tileOpBuffer = [surface->device newBufferWithLength: MG_MTL_TILE_OP_BUFFER_SIZE backend->tileOpBuffer = [surface->device newBufferWithLength: MG_MTL_DEFAULT_TILE_OP_BUFFER_LEN * sizeof(mg_mtl_tile_op)
options: bufferOptions]; options: bufferOptions];
backend->tileOpCountBuffer = [surface->device newBufferWithLength: sizeof(int) backend->tileOpCountBuffer = [surface->device newBufferWithLength: sizeof(int)

368
src/mtl_renderer.metal
View File

@ -231,8 +231,9 @@ kernel void mtl_path_setup(constant int* pathCount [[buffer(0)]],
device mg_mtl_path_queue* pathQueueBuffer [[buffer(2)]], device mg_mtl_path_queue* pathQueueBuffer [[buffer(2)]],
device mg_mtl_tile_queue* tileQueueBuffer [[buffer(3)]], device mg_mtl_tile_queue* tileQueueBuffer [[buffer(3)]],
device atomic_int* tileQueueCount [[buffer(4)]], device atomic_int* tileQueueCount [[buffer(4)]],
constant int* tileSize [[buffer(5)]], constant int* tileQueueMax [[buffer(5)]],
constant float* scale [[buffer(6)]], constant int* tileSize [[buffer(6)]],
constant float* scale [[buffer(7)]],
uint pathIndex [[thread_position_in_grid]]) uint pathIndex [[thread_position_in_grid]])
{ {
const device mg_mtl_path* path = &pathBuffer[pathIndex]; const device mg_mtl_path* path = &pathBuffer[pathIndex];
@ -254,16 +255,24 @@ kernel void mtl_path_setup(constant int* pathCount [[buffer(0)]],
int tileQueuesIndex = atomic_fetch_add_explicit(tileQueueCount, tileCount, memory_order_relaxed); int tileQueuesIndex = atomic_fetch_add_explicit(tileQueueCount, tileCount, memory_order_relaxed);
pathQueueBuffer[pathIndex].area = int4(firstTile.x, firstTile.y, nTilesX, nTilesY); if(tileQueuesIndex + tileCount >= tileQueueMax[0])
pathQueueBuffer[pathIndex].tileQueues = tileQueuesIndex;
device mg_mtl_tile_queue* tileQueues = &tileQueueBuffer[tileQueuesIndex];
for(int i=0; i<tileCount; i++)
{ {
atomic_store_explicit(&tileQueues[i].first, -1, memory_order_relaxed); pathQueueBuffer[pathIndex].area = int4(0);
tileQueues[i].last = -1; pathQueueBuffer[pathIndex].tileQueues = 0;
atomic_store_explicit(&tileQueues[i].windingOffset, 0, memory_order_relaxed); }
else
{
pathQueueBuffer[pathIndex].area = int4(firstTile.x, firstTile.y, nTilesX, nTilesY);
pathQueueBuffer[pathIndex].tileQueues = tileQueuesIndex;
device mg_mtl_tile_queue* tileQueues = &tileQueueBuffer[tileQueuesIndex];
for(int i=0; i<tileCount; i++)
{
atomic_store_explicit(&tileQueues[i].first, -1, memory_order_relaxed);
tileQueues[i].last = -1;
atomic_store_explicit(&tileQueues[i].windingOffset, 0, memory_order_relaxed);
}
} }
} }
@ -376,6 +385,9 @@ typedef struct mtl_segment_setup_context
int pathIndex; int pathIndex;
int tileOpMax;
int segmentMax;
} mtl_segment_setup_context; } mtl_segment_setup_context;
void mtl_segment_bin_to_tiles(thread mtl_segment_setup_context* context, device mg_mtl_segment* seg) void mtl_segment_bin_to_tiles(thread mtl_segment_setup_context* context, device mg_mtl_segment* seg)
@ -439,34 +451,38 @@ void mtl_segment_bin_to_tiles(thread mtl_segment_setup_context* context, device
if(crossL || crossR || crossT || crossB || s0Inside || s1Inside) if(crossL || crossR || crossT || crossB || s0Inside || s1Inside)
{ {
int tileOpIndex = atomic_fetch_add_explicit(context->tileOpCount, 1, memory_order_relaxed); 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; if(tileOpIndex < context->tileOpMax)
op->index = segIndex;
op->crossRight = false;
op->next = -1;
int tileIndex = y*pathArea.z + x;
device mg_mtl_tile_queue* tile = &context->tileQueues[tileIndex];
op->next = atomic_exchange_explicit(&tile->first, tileOpIndex, memory_order_relaxed);
if(op->next == -1)
{ {
tile->last = tileOpIndex; device mg_mtl_tile_op* op = &context->tileOpBuffer[tileOpIndex];
}
//NOTE: if the segment crosses the tile's bottom boundary, update the tile's winding offset op->kind = MG_MTL_OP_SEGMENT;
if(crossB) op->index = segIndex;
{ op->crossRight = false;
mtl_log(context->log, "cross bottom boundary, increment "); op->next = -1;
mtl_log_f32(context->log, seg->windingIncrement);
mtl_log(context->log, "\n");
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 int tileIndex = y*pathArea.z + x;
if(crossR) device mg_mtl_tile_queue* tile = &context->tileQueues[tileIndex];
{ op->next = atomic_exchange_explicit(&tile->first, tileOpIndex, memory_order_relaxed);
op->crossRight = true; if(op->next == -1)
{
tile->last = tileOpIndex;
}
//NOTE: if the segment crosses the tile's bottom boundary, update the tile's winding offset
if(crossB)
{
mtl_log(context->log, "cross bottom boundary, increment ");
mtl_log_f32(context->log, seg->windingIncrement);
mtl_log(context->log, "\n");
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
if(crossR)
{
op->crossRight = true;
}
} }
} }
} }
@ -508,54 +524,60 @@ device mg_mtl_segment* mtl_segment_push(thread mtl_segment_setup_context* contex
} break; } break;
} }
device mg_mtl_segment* seg = 0;
int segIndex = atomic_fetch_add_explicit(context->segmentCount, 1, memory_order_relaxed); 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; if(segIndex < context->segmentMax)
bool goingRight = e.x >= s.x;
seg->kind = kind;
seg->pathIndex = context->pathIndex;
seg->windingIncrement = goingUp? 1 : -1;
seg->box = (vector_float4){min(s.x, e.x),
min(s.y, e.y),
max(s.x, e.x),
max(s.y, e.y)};
float dx = c.x - seg->box.x;
float dy = c.y - seg->box.y;
float alpha = (seg->box.w - seg->box.y)/(seg->box.z - seg->box.x);
float ofs = seg->box.w - seg->box.y;
if(goingUp == goingRight)
{ {
if(seg->kind == MG_MTL_LINE) 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 = goingUp? 1 : -1;
seg->box = (vector_float4){min(s.x, e.x),
min(s.y, e.y),
max(s.x, e.x),
max(s.y, e.y)};
float dx = c.x - seg->box.x;
float dy = c.y - seg->box.y;
float alpha = (seg->box.w - seg->box.y)/(seg->box.z - seg->box.x);
float ofs = seg->box.w - seg->box.y;
if(goingUp == goingRight)
{ {
seg->config = MG_MTL_BR; if(seg->kind == MG_MTL_LINE)
} {
else if(dy > alpha*dx) seg->config = MG_MTL_BR;
{ }
seg->config = MG_MTL_TL; else if(dy > alpha*dx)
{
seg->config = MG_MTL_TL;
}
else
{
seg->config = MG_MTL_BR;
}
} }
else else
{ {
seg->config = MG_MTL_BR; if(seg->kind == MG_MTL_LINE)
} {
} seg->config = MG_MTL_TR;
else }
{ else if(dy < ofs - alpha*dx)
if(seg->kind == MG_MTL_LINE) {
{ seg->config = MG_MTL_BL;
seg->config = MG_MTL_TR; }
} else
else if(dy < ofs - alpha*dx) {
{ seg->config = MG_MTL_TR;
seg->config = MG_MTL_BL; }
}
else
{
seg->config = MG_MTL_TR;
} }
} }
return(seg); return(seg);
@ -567,8 +589,11 @@ 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]; if(seg)
mtl_segment_bin_to_tiles(context, seg); {
seg->hullVertex = p[0];
mtl_segment_bin_to_tiles(context, seg);
}
} }
float2 mtl_quadratic_blossom(float2 p[3], float u, float v) float2 mtl_quadratic_blossom(float2 p[3], float u, float v)
@ -636,26 +661,29 @@ void mtl_quadratic_emit(thread mtl_segment_setup_context* context,
{ {
device mg_mtl_segment* seg = mtl_segment_push(context, p, MG_MTL_QUADRATIC); device mg_mtl_segment* seg = mtl_segment_push(context, p, MG_MTL_QUADRATIC);
//NOTE: compute implicit equation matrix if(seg)
float det = p[0].x*(p[1].y-p[2].y) + p[1].x*(p[2].y-p[0].y) + p[2].x*(p[0].y - p[1].y); {
//NOTE: compute implicit equation matrix
float det = p[0].x*(p[1].y-p[2].y) + p[1].x*(p[2].y-p[0].y) + p[2].x*(p[0].y - p[1].y);
float a = p[0].y - p[1].y + 0.5*(p[2].y - p[0].y); float a = p[0].y - p[1].y + 0.5*(p[2].y - p[0].y);
float b = p[1].x - p[0].x + 0.5*(p[0].x - p[2].x); float b = p[1].x - p[0].x + 0.5*(p[0].x - p[2].x);
float c = p[0].x*p[1].y - p[1].x*p[0].y + 0.5*(p[2].x*p[0].y - p[0].x*p[2].y); float c = p[0].x*p[1].y - p[1].x*p[0].y + 0.5*(p[2].x*p[0].y - p[0].x*p[2].y);
float d = p[0].y - p[1].y; float d = p[0].y - p[1].y;
float e = p[1].x - p[0].x; float e = p[1].x - p[0].x;
float f = p[0].x*p[1].y - p[1].x*p[0].y; float f = p[0].x*p[1].y - p[1].x*p[0].y;
float flip = (seg->config == MG_MTL_TL || seg->config == MG_MTL_BL)? -1 : 1; float flip = (seg->config == MG_MTL_TL || seg->config == MG_MTL_BL)? -1 : 1;
float g = flip*(p[2].x*(p[0].y - p[1].y) + p[0].x*(p[1].y - p[2].y) + p[1].x*(p[2].y - p[0].y)); float g = flip*(p[2].x*(p[0].y - p[1].y) + p[0].x*(p[1].y - p[2].y) + p[1].x*(p[2].y - p[0].y));
seg->implicitMatrix = (1/det)*matrix_float3x3({a, d, 0.}, seg->implicitMatrix = (1/det)*matrix_float3x3({a, d, 0.},
{b, e, 0.}, {b, e, 0.},
{c, f, g}); {c, f, g});
seg->hullVertex = p[1]; seg->hullVertex = p[1];
mtl_segment_bin_to_tiles(context, seg); mtl_segment_bin_to_tiles(context, seg);
}
} }
void mtl_quadratic_setup(thread mtl_segment_setup_context* context, thread float2* p) void mtl_quadratic_setup(thread mtl_segment_setup_context* context, thread float2* p)
@ -1044,70 +1072,73 @@ void mtl_cubic_emit(thread mtl_segment_setup_context* context, mtl_cubic_info cu
{ {
device mg_mtl_segment* seg = mtl_segment_push(context, sp, MG_MTL_CUBIC); device mg_mtl_segment* seg = mtl_segment_push(context, sp, MG_MTL_CUBIC);
float2 v0 = p[0]; if(seg)
float2 v1 = p[3];
float2 v2;
matrix_float3x3 K;
float sqrNorm0 = length_squared(p[1]-p[0]);
float sqrNorm1 = length_squared(p[2]-p[3]);
//TODO: should not be the local sub-curve, but the global curve!!!
if(length_squared(p[0]-p[3]) > 1e-5)
{ {
if(sqrNorm0 >= sqrNorm1) float2 v0 = p[0];
{ float2 v1 = p[3];
v2 = p[1]; float2 v2;
K = {curve.K[0].xyz, curve.K[3].xyz, curve.K[1].xyz}; matrix_float3x3 K;
float sqrNorm0 = length_squared(p[1]-p[0]);
float sqrNorm1 = length_squared(p[2]-p[3]);
//TODO: should not be the local sub-curve, but the global curve!!!
if(length_squared(p[0]-p[3]) > 1e-5)
{
if(sqrNorm0 >= sqrNorm1)
{
v2 = p[1];
K = {curve.K[0].xyz, curve.K[3].xyz, curve.K[1].xyz};
}
else
{
v2 = p[2];
K = {curve.K[0].xyz, curve.K[3].xyz, curve.K[2].xyz};
}
} }
else else
{ {
v1 = p[1];
v2 = p[2]; v2 = p[2];
K = {curve.K[0].xyz, curve.K[3].xyz, curve.K[2].xyz}; K = {curve.K[0].xyz, curve.K[1].xyz, curve.K[2].xyz};
} }
} //NOTE: set matrices
else
{
v1 = p[1];
v2 = p[2];
K = {curve.K[0].xyz, curve.K[1].xyz, curve.K[2].xyz};
}
//NOTE: set matrices
//TODO: should we compute matrix relative to a base point to avoid loss of precision //TODO: should we compute matrix relative to a base point to avoid loss of precision
// when computing barycentric matrix? // when computing barycentric matrix?
matrix_float3x3 B = mtl_barycentric_matrix(v0, v1, v2); matrix_float3x3 B = mtl_barycentric_matrix(v0, v1, v2);
seg->implicitMatrix = K*B; seg->implicitMatrix = K*B;
seg->hullVertex = mtl_select_hull_vertex(sp[0], sp[1], sp[2], sp[3], context->log); seg->hullVertex = mtl_select_hull_vertex(sp[0], sp[1], sp[2], sp[3], context->log);
//NOTE: compute sign flip //NOTE: compute sign flip
seg->sign = 1; seg->sign = 1;
if(curve.kind == MTL_CUBIC_SERPENTINE if(curve.kind == MTL_CUBIC_SERPENTINE
|| curve.kind == MTL_CUBIC_CUSP) || curve.kind == MTL_CUBIC_CUSP)
{ {
seg->sign = (curve.d1 < 0)? -1 : 1; seg->sign = (curve.d1 < 0)? -1 : 1;
}
else if(curve.kind == MTL_CUBIC_LOOP)
{
float d1 = curve.d1;
float d2 = curve.d2;
float d3 = curve.d3;
float H0 = d3*d1-square(d2) + d1*d2*s0 - square(d1)*square(s0);
float H1 = d3*d1-square(d2) + d1*d2*s1 - square(d1)*square(s1);
float H = (abs(H0) > abs(H1)) ? H0 : H1;
seg->sign = (H*d1 > 0) ? -1 : 1;
}
if(sp[3].y > sp[0].y)
{
seg->sign *= -1;
}
//NOTE: bin to tiles
mtl_segment_bin_to_tiles(context, seg);
} }
else if(curve.kind == MTL_CUBIC_LOOP)
{
float d1 = curve.d1;
float d2 = curve.d2;
float d3 = curve.d3;
float H0 = d3*d1-square(d2) + d1*d2*s0 - square(d1)*square(s0);
float H1 = d3*d1-square(d2) + d1*d2*s1 - square(d1)*square(s1);
float H = (abs(H0) > abs(H1)) ? H0 : H1;
seg->sign = (H*d1 > 0) ? -1 : 1;
}
if(sp[3].y > sp[0].y)
{
seg->sign *= -1;
}
//NOTE: bin to tiles
mtl_segment_bin_to_tiles(context, seg);
} }
void mtl_cubic_setup(thread mtl_segment_setup_context* context, float2 p[4]) void mtl_cubic_setup(thread mtl_segment_setup_context* context, float2 p[4])
@ -1229,11 +1260,13 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]],
device mg_mtl_tile_queue* tileQueueBuffer [[buffer(5)]], device mg_mtl_tile_queue* tileQueueBuffer [[buffer(5)]],
device mg_mtl_tile_op* tileOpBuffer [[buffer(6)]], device mg_mtl_tile_op* tileOpBuffer [[buffer(6)]],
device atomic_int* tileOpCount [[buffer(7)]], device atomic_int* tileOpCount [[buffer(7)]],
constant int* tileSize [[buffer(8)]], constant int* segmentMax [[buffer(8)]],
constant float* scale [[buffer(9)]], constant int* tileOpMax [[buffer(9)]],
constant int* tileSize [[buffer(10)]],
constant float* scale [[buffer(11)]],
device char* logBuffer [[buffer(10)]], device char* logBuffer [[buffer(12)]],
device atomic_int* logOffsetBuffer [[buffer(11)]], device atomic_int* logOffsetBuffer [[buffer(13)]],
uint eltIndex [[thread_position_in_grid]]) uint eltIndex [[thread_position_in_grid]])
{ {
const device mg_mtl_path_elt* elt = &elementBuffer[eltIndex]; const device mg_mtl_path_elt* elt = &elementBuffer[eltIndex];
@ -1247,10 +1280,12 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]],
.tileQueues = tileQueues, .tileQueues = tileQueues,
.tileOpBuffer = tileOpBuffer, .tileOpBuffer = tileOpBuffer,
.tileOpCount = tileOpCount, .tileOpCount = tileOpCount,
.tileOpMax = tileOpMax[0],
.segmentMax = segmentMax[0],
.tileSize = tileSize[0], .tileSize = tileSize[0],
.log.buffer = logBuffer, .log.buffer = logBuffer,
.log.offset = logOffsetBuffer, .log.offset = logOffsetBuffer,
.log.enabled = false}; .log.enabled = false,};
switch(elt->kind) switch(elt->kind)
{ {
@ -1327,10 +1362,11 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]],
device atomic_int* tileOpCount [[buffer(5)]], device atomic_int* tileOpCount [[buffer(5)]],
device MTLDispatchThreadgroupsIndirectArguments* dispatchBuffer [[buffer(6)]], device MTLDispatchThreadgroupsIndirectArguments* dispatchBuffer [[buffer(6)]],
device mg_mtl_screen_tile* screenTilesBuffer [[buffer(7)]], device mg_mtl_screen_tile* screenTilesBuffer [[buffer(7)]],
constant int* tileSize [[buffer(8)]], constant int* tileOpMax [[buffer(8)]],
constant float* scale [[buffer(9)]], constant int* tileSize [[buffer(9)]],
device char* logBuffer [[buffer(10)]], constant float* scale [[buffer(10)]],
device atomic_int* logOffsetBuffer [[buffer(11)]], device char* logBuffer [[buffer(11)]],
device atomic_int* logOffsetBuffer [[buffer(12)]],
uint2 threadCoord [[thread_position_in_grid]], uint2 threadCoord [[thread_position_in_grid]],
uint2 gridSize [[threads_per_grid]]) uint2 gridSize [[threads_per_grid]])
{ {
@ -1393,6 +1429,12 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]],
//NOTE: tile is full covered. Add path start op (with winding offset). //NOTE: tile is full covered. Add path start op (with winding offset).
// Additionally if color is opaque and tile is fully inside clip, 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); int pathOpIndex = atomic_fetch_add_explicit(tileOpCount, 1, memory_order_relaxed);
if(pathOpIndex >= tileOpMax[0])
{
return;
}
device mg_mtl_tile_op* pathOp = &tileOpBuffer[pathOpIndex]; device mg_mtl_tile_op* pathOp = &tileOpBuffer[pathOpIndex];
pathOp->kind = MG_MTL_OP_CLIP_FILL; pathOp->kind = MG_MTL_OP_CLIP_FILL;
pathOp->next = -1; pathOp->next = -1;
@ -1421,6 +1463,11 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]],
{ {
//NOTE: add path start op (with winding offset) //NOTE: add path start op (with winding offset)
int startOpIndex = atomic_fetch_add_explicit(tileOpCount, 1, memory_order_relaxed); int startOpIndex = atomic_fetch_add_explicit(tileOpCount, 1, memory_order_relaxed);
if(startOpIndex >= tileOpMax[0])
{
return;
}
device mg_mtl_tile_op* startOp = &tileOpBuffer[startOpIndex]; device mg_mtl_tile_op* startOp = &tileOpBuffer[startOpIndex];
startOp->kind = MG_MTL_OP_START; startOp->kind = MG_MTL_OP_START;
startOp->next = -1; startOp->next = -1;
@ -1439,6 +1486,11 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]],
//NOTE: add path end op //NOTE: add path end op
int endOpIndex = atomic_fetch_add_explicit(tileOpCount, 1, memory_order_relaxed); int endOpIndex = atomic_fetch_add_explicit(tileOpCount, 1, memory_order_relaxed);
if(endOpIndex >= tileOpMax[0])
{
return;
}
device mg_mtl_tile_op* endOp = &tileOpBuffer[endOpIndex]; device mg_mtl_tile_op* endOp = &tileOpBuffer[endOpIndex];
endOp->kind = MG_MTL_OP_END; endOp->kind = MG_MTL_OP_END;
endOp->next = -1; endOp->next = -1;
@ -1446,7 +1498,6 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]],
*nextLink = endOpIndex; *nextLink = endOpIndex;
nextLink = &endOp->next; nextLink = &endOp->next;
} }
} }
} }
@ -1611,6 +1662,13 @@ kernel void mtl_raster(const device mg_mtl_screen_tile* screenTilesBuffer [[buff
} }
opIndex = op->next; opIndex = op->next;
} }
/*
if((pixelCoord.x % tileSize[0] == 0) || (pixelCoord.y % tileSize[0] == 0))
{
color = float4(0, 0, 0, 1);
}
//*/
outTexture.write(color, pixelCoord); outTexture.write(color, pixelCoord);
} }
@ -1631,7 +1689,7 @@ vertex vs_out mtl_vertex_shader(ushort vid [[vertex_id]])
return(out); return(out);
} }
fragment float4 mtl_fragment_shader(vs_out i [[stage_in]], texture2d<float> tex [[texture(0)]]) fragment float4 mtl_fragment_shader(vs_out i [[stage_in]], texture2d<float, access::sample> tex [[texture(0)]])
{ {
constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear);
return(tex.sample(smp, i.uv)); return(tex.sample(smp, i.uv));

4
src/wgl_surface.c
View File

@ -254,7 +254,7 @@ mg_surface_data* mg_wgl_surface_create_for_window(mp_window window)
int contextAttrs[] = { int contextAttrs[] = {
WGL_CONTEXT_MAJOR_VERSION_ARB, 4, WGL_CONTEXT_MAJOR_VERSION_ARB, 4,
WGL_CONTEXT_MINOR_VERSION_ARB, 3, WGL_CONTEXT_MINOR_VERSION_ARB, 4,
WGL_CONTEXT_PROFILE_MASK_ARB, WGL_CONTEXT_CORE_PROFILE_BIT_ARB, WGL_CONTEXT_PROFILE_MASK_ARB, WGL_CONTEXT_CORE_PROFILE_BIT_ARB,
0}; 0};
@ -270,7 +270,7 @@ mg_surface_data* mg_wgl_surface_create_for_window(mp_window window)
//NOTE: make gl context current and load api //NOTE: make gl context current and load api
wglMakeCurrent(surface->hDC, surface->glContext); wglMakeCurrent(surface->hDC, surface->glContext);
wglSwapIntervalEXT(1); wglSwapIntervalEXT(1);
mg_gl_load_gl43(&surface->api, mg_wgl_get_proc); mg_gl_load_gl44(&surface->api, mg_wgl_get_proc);
} }
} }
return((mg_surface_data*)surface); return((mg_surface_data*)surface);