[mtl canvas] prevent tiles entirely covered by textured shapes to be treated as solid (as the texture could be non-uniformly transparent)

This commit is contained in:
Martin Fouilleul 2023-03-22 12:02:13 +01:00
parent 26f669f3c1
commit b0be21795e
5 changed files with 27 additions and 18 deletions

View File

@ -816,6 +816,7 @@ u32 mg_next_shape(mg_canvas_data* canvas, mg_attributes* attributes)
canvas->nextShapeIndex++; canvas->nextShapeIndex++;
*(mg_color*)(((char*)layout->colorBuffer) + index*layout->colorStride) = attributes->color; *(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); return(index);
} }

View File

@ -90,6 +90,9 @@ typedef struct mg_vertex_layout
char* colorBuffer; char* colorBuffer;
u32 colorStride; u32 colorStride;
char* texturedBuffer;
u32 texturedStride;
char* shapeIndexBuffer; char* shapeIndexBuffer;
u32 shapeIndexStride; u32 shapeIndexStride;

View File

@ -104,6 +104,10 @@ void mg_mtl_canvas_update_vertex_layout(mg_mtl_canvas_backend* backend)
.colorBuffer = shapeBase + offsetof(mg_shape, color), .colorBuffer = shapeBase + offsetof(mg_shape, color),
.colorStride = sizeof(mg_shape), .colorStride = sizeof(mg_shape),
.texturedBuffer = shapeBase + offsetof(mg_shape, textured),
.texturedStride = sizeof(mg_shape),
.clipBuffer = shapeBase + offsetof(mg_shape, clip), .clipBuffer = shapeBase + offsetof(mg_shape, clip),
.clipStride = sizeof(mg_shape), .clipStride = sizeof(mg_shape),
.uvTransformBuffer = shapeBase + offsetof(mg_shape, uvTransform), .uvTransformBuffer = shapeBase + offsetof(mg_shape, uvTransform),

View File

@ -30,7 +30,7 @@ typedef struct mg_shape
vector_float4 color; vector_float4 color;
vector_float4 clip; vector_float4 clip;
float uvTransform[6]; float uvTransform[6];
bool textured;
} mg_shape; } mg_shape;
typedef struct mg_triangle_data typedef struct mg_triangle_data
@ -64,17 +64,15 @@ typedef struct mg_triangle_data
using namespace metal; using namespace metal;
#endif #endif
#define MG_TILE_CMD_MASK (1<<31)
typedef enum mg_tile_cmd_kind typedef enum mg_tile_cmd_kind
{ {
mg_cmd_triangle, mg_cmd_triangle = 0,
mg_cmd_color mg_cmd_color = 1<<31,
} mg_tile_cmd_kind; } mg_tile_cmd_kind;
typedef struct mg_tile_cmd typedef int mg_tile_cmd;
{
mg_tile_cmd_kind kind;
int triangleIndex;
} mg_tile_cmd;
typedef struct mg_tile_elt typedef struct mg_tile_elt
{ {
@ -89,6 +87,7 @@ typedef struct mg_tile
atomic_int eltCount; atomic_int eltCount;
atomic_int partial; atomic_int partial;
atomic_int flipCount; atomic_int flipCount;
bool textured;
} mg_tile; } mg_tile;

View File

@ -87,6 +87,7 @@ kernel void ShapeSetup(constant mg_shape* shapeBuffer [[buffer(0)]],
for(int i=0; i<tileCount; i++) for(int i=0; i<tileCount; i++)
{ {
tiles[i].color = shapeBuffer[gid].color; tiles[i].color = shapeBuffer[gid].color;
tiles[i].textured = shapeBuffer[gid].textured;
atomic_store_explicit(&tiles[i].firstElt, -1, memory_order_relaxed); atomic_store_explicit(&tiles[i].firstElt, -1, memory_order_relaxed);
atomic_store_explicit(&tiles[i].eltCount, 0, memory_order_relaxed); atomic_store_explicit(&tiles[i].eltCount, 0, memory_order_relaxed);
atomic_store_explicit(&tiles[i].partial, 0, memory_order_relaxed); atomic_store_explicit(&tiles[i].partial, 0, memory_order_relaxed);
@ -300,15 +301,14 @@ kernel void GatherKernel(const device mg_shape_queue* shapeQueueBuffer [[buffer(
{ {
if(atomic_load_explicit(&tile->flipCount, memory_order_relaxed) & 0x01) if(atomic_load_explicit(&tile->flipCount, 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; int firstEltIndex = *(device int*)&tile->firstElt;
const device mg_tile_elt* elt = &eltBuffer[firstEltIndex]; const device mg_tile_elt* elt = &eltBuffer[firstEltIndex];
count = 0; count = 0;
tileArray[count].kind = mg_cmd_color; tileArray[count] = mg_cmd_color | elt->triangleIndex;
tileArray[count].triangleIndex = elt->triangleIndex;
count++; count++;
continue; continue;
} }
@ -328,8 +328,7 @@ kernel void GatherKernel(const device mg_shape_queue* shapeQueueBuffer [[buffer(
elt = &eltBuffer[eltIndex]; elt = &eltBuffer[eltIndex];
eltIndex = elt->next; eltIndex = elt->next;
tileArray[count].kind = mg_cmd_triangle; tileArray[count] = elt->triangleIndex;
tileArray[count].triangleIndex = elt->triangleIndex;
count++; count++;
} }
} }
@ -420,10 +419,13 @@ kernel void RenderKernel(const device uint* tileCounters [[buffer(0)]],
for(uint tileArrayIndex=0; tileArrayIndex < tileCounter; tileArrayIndex++) for(uint tileArrayIndex=0; tileArrayIndex < tileCounter; tileArrayIndex++)
{ {
const device mg_tile_cmd* cmd = &tileArrayBuffer[RENDERER_TILE_BUFFER_COUNT * tileIndex + tileArrayIndex]; mg_tile_cmd cmd = tileArrayBuffer[RENDERER_TILE_BUFFER_COUNT * tileIndex + tileArrayIndex];
const device mg_triangle_data* triangle = &triangleArray[cmd->triangleIndex];
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; sampleIndex<sampleCount; sampleIndex++) for(int sampleIndex=0; sampleIndex<sampleCount; sampleIndex++)
{ {
@ -432,7 +434,7 @@ kernel void RenderKernel(const device uint* tileCounters [[buffer(0)]],
currentShapeIndex[sampleIndex] = triangle->shapeIndex; currentShapeIndex[sampleIndex] = triangle->shapeIndex;
} }
} }
else if(cmd->kind == mg_cmd_triangle) else
{ {
int2 p0 = triangle->p0; int2 p0 = triangle->p0;
int2 p1 = triangle->p1; int2 p1 = triangle->p1;