diff --git a/src/graphics.c b/src/graphics.c index f24fbcd..1f405a4 100644 --- a/src/graphics.c +++ b/src/graphics.c @@ -816,6 +816,7 @@ u32 mg_next_shape(mg_canvas_data* canvas, mg_attributes* attributes) canvas->nextShapeIndex++; *(mg_color*)(((char*)layout->colorBuffer) + index*layout->colorStride) = attributes->color; + *(bool*)(((char*)layout->texturedBuffer) + index*layout->texturedStride) = !mg_image_is_nil(attributes->image); return(index); } diff --git a/src/graphics_internal.h b/src/graphics_internal.h index 9aa2d2f..90fdfd4 100644 --- a/src/graphics_internal.h +++ b/src/graphics_internal.h @@ -90,6 +90,9 @@ typedef struct mg_vertex_layout char* colorBuffer; u32 colorStride; + char* texturedBuffer; + u32 texturedStride; + char* shapeIndexBuffer; u32 shapeIndexStride; diff --git a/src/mtl_canvas.m b/src/mtl_canvas.m index ccb867a..6da7310 100644 --- a/src/mtl_canvas.m +++ b/src/mtl_canvas.m @@ -104,6 +104,10 @@ void mg_mtl_canvas_update_vertex_layout(mg_mtl_canvas_backend* backend) .colorBuffer = shapeBase + offsetof(mg_shape, color), .colorStride = sizeof(mg_shape), + + .texturedBuffer = shapeBase + offsetof(mg_shape, textured), + .texturedStride = sizeof(mg_shape), + .clipBuffer = shapeBase + offsetof(mg_shape, clip), .clipStride = sizeof(mg_shape), .uvTransformBuffer = shapeBase + offsetof(mg_shape, uvTransform), diff --git a/src/mtl_shader.h b/src/mtl_shader.h index b4c2a20..e1bea72 100644 --- a/src/mtl_shader.h +++ b/src/mtl_shader.h @@ -30,7 +30,7 @@ typedef struct mg_shape vector_float4 color; vector_float4 clip; float uvTransform[6]; - + bool textured; } mg_shape; typedef struct mg_triangle_data @@ -64,17 +64,15 @@ typedef struct mg_triangle_data using namespace metal; #endif +#define MG_TILE_CMD_MASK (1<<31) + typedef enum mg_tile_cmd_kind { - mg_cmd_triangle, - mg_cmd_color + mg_cmd_triangle = 0, + mg_cmd_color = 1<<31, } mg_tile_cmd_kind; -typedef struct mg_tile_cmd -{ - mg_tile_cmd_kind kind; - int triangleIndex; -} mg_tile_cmd; +typedef int mg_tile_cmd; typedef struct mg_tile_elt { @@ -89,6 +87,7 @@ typedef struct mg_tile atomic_int eltCount; atomic_int partial; atomic_int flipCount; + bool textured; } mg_tile; diff --git a/src/mtl_shader.metal b/src/mtl_shader.metal index 03ed9aa..3a1bd82 100644 --- a/src/mtl_shader.metal +++ b/src/mtl_shader.metal @@ -87,6 +87,7 @@ kernel void ShapeSetup(constant mg_shape* shapeBuffer [[buffer(0)]], for(int i=0; iflipCount, memory_order_relaxed) & 0x01) { - if(tile->color.a == 1) + if(tile->color.a == 1 && !tile->textured) { - //NOTE: tile is full covered by a solid color, reset counter + //NOTE: tile is full covered by a solid color, reset counter and push a color command int firstEltIndex = *(device int*)&tile->firstElt; const device mg_tile_elt* elt = &eltBuffer[firstEltIndex]; count = 0; - tileArray[count].kind = mg_cmd_color; - tileArray[count].triangleIndex = elt->triangleIndex; + tileArray[count] = mg_cmd_color | elt->triangleIndex; count++; continue; } @@ -328,8 +328,7 @@ kernel void GatherKernel(const device mg_shape_queue* shapeQueueBuffer [[buffer( elt = &eltBuffer[eltIndex]; eltIndex = elt->next; - tileArray[count].kind = mg_cmd_triangle; - tileArray[count].triangleIndex = elt->triangleIndex; + tileArray[count] = elt->triangleIndex; count++; } } @@ -420,10 +419,13 @@ kernel void RenderKernel(const device uint* tileCounters [[buffer(0)]], for(uint tileArrayIndex=0; tileArrayIndex < tileCounter; tileArrayIndex++) { - const device mg_tile_cmd* cmd = &tileArrayBuffer[RENDERER_TILE_BUFFER_COUNT * tileIndex + tileArrayIndex]; - const device mg_triangle_data* triangle = &triangleArray[cmd->triangleIndex]; + mg_tile_cmd cmd = tileArrayBuffer[RENDERER_TILE_BUFFER_COUNT * tileIndex + tileArrayIndex]; - if(cmd->kind == mg_cmd_color) + int cmdKind = cmd & MG_TILE_CMD_MASK; + int triangleIndex = cmd & ~(MG_TILE_CMD_MASK); + const device mg_triangle_data* triangle = &triangleArray[triangleIndex]; + + if(cmdKind == mg_cmd_color) { for(int sampleIndex=0; sampleIndexshapeIndex; } } - else if(cmd->kind == mg_cmd_triangle) + else { int2 p0 = triangle->p0; int2 p1 = triangle->p1;