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

[mtl canvas] multi-sampling

This commit is contained in:
Martin Fouilleul 2023-04-07 18:18:50 +02:00
parent 65b5a4b52a
commit 2b7aadf180
2 changed files with 60 additions and 58 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
build.sh
View File

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

116
src/mtl_renderer.metal
View File

@ -1487,94 +1487,96 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
uint2 threadCoord [[thread_position_in_grid]], uint2 threadCoord [[thread_position_in_grid]],
uint2 gridSize [[threads_per_grid]]) uint2 gridSize [[threads_per_grid]])
{ {
int2 pixelCoord = int2(threadCoord); float2 pixelCoord = float2(threadCoord);
int2 tileCoord = pixelCoord / tileSize[0]; int2 tileCoord = int2(threadCoord) / tileSize[0];
int nTilesX = (int(gridSize.x) + tileSize[0] - 1)/tileSize[0]; int nTilesX = (int(gridSize.x) + tileSize[0] - 1)/tileSize[0];
int tileIndex = tileCoord.y * nTilesX + tileCoord.x; int tileIndex = tileCoord.y * nTilesX + tileCoord.x;
float4 color = float4(0, 0, 0, 0);
int pathIndex = 0; int pathIndex = 0;
int winding = 0;
int opIndex = screenTilesBuffer[tileIndex]; int opIndex = screenTilesBuffer[tileIndex];
const int MG_MTL_SAMPLE_COUNT = 8;
const float2 sampleCoords[MG_MTL_SAMPLE_COUNT] = {pixelCoord + float2(1, 3)/16,
pixelCoord + float2(-1, -3)/16,
pixelCoord + float2(5, -1)/16,
pixelCoord + float2(-3, 5)/16,
pixelCoord + float2(-5, -5)/16,
pixelCoord + float2(-7, 1)/16,
pixelCoord + float2(3, -7)/16,
pixelCoord + float2(7, 7)/16};
float4 color[MG_MTL_SAMPLE_COUNT] = {0};
int winding[MG_MTL_SAMPLE_COUNT] = {0};
while(opIndex != -1) while(opIndex != -1)
{ {
const device mg_mtl_tile_op* op = &tileOpBuffer[opIndex]; const device mg_mtl_tile_op* op = &tileOpBuffer[opIndex];
if(op->kind == MG_MTL_OP_START) if(op->kind == MG_MTL_OP_START)
{ {
bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding & 1)) float4 pathColor = pathBuffer[pathIndex].color;
||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding != 0)); pathColor.rgb *= pathColor.a;
if(filled)
{
float4 pathColor = pathBuffer[pathIndex].color;
pathColor.rgb *= pathColor.a;
color = color*(1-pathColor.a) + pathColor;
}
for(int sampleIndex=0; sampleIndex<MG_MTL_SAMPLE_COUNT; sampleIndex++)
{
bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1))
||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0));
if(filled)
{
color[sampleIndex] = color[sampleIndex]*(1-pathColor.a) + pathColor;
}
winding[sampleIndex] = op->windingOffset;
}
pathIndex = op->index; pathIndex = op->index;
winding = op->windingOffset;
/*
if(winding < 0)
{
color = float4(1, 0, 0, 1);
}
else if(winding > 0)
{
color = float4(0, 1, 0, 1);
}
if(winding && ((pixelCoord.x % 16) == 0))
{
outTexture.write(color, uint2(pixelCoord));
return;
}
if(op->next != -1)
{
color = float4(0, 0.5, 1, 1);
}
*/
} }
else if(op->kind == MG_MTL_OP_SEGMENT) else if(op->kind == MG_MTL_OP_SEGMENT)
{ {
const device mg_mtl_segment* seg = &segmentBuffer[op->index]; const device mg_mtl_segment* seg = &segmentBuffer[op->index];
if( (pixelCoord.y > seg->box.y) for(int sampleIndex=0; sampleIndex<MG_MTL_SAMPLE_COUNT; sampleIndex++)
&&(pixelCoord.y <= seg->box.w)
&&(mtl_side_of_segment(float2(pixelCoord), seg) < 0))
{ {
winding += seg->windingIncrement; float2 sampleCoord = sampleCoords[sampleIndex];
}
if(op->crossRight) if( (sampleCoord.y > seg->box.y)
{ &&(sampleCoord.y <= seg->box.w)
//color = float4(0, 1, 1, 1); &&(mtl_side_of_segment(sampleCoord, seg) < 0))
if( (seg->config == MG_MTL_BR || seg->config == MG_MTL_TL)
&&(pixelCoord.y > seg->box.w))
{ {
winding += seg->windingIncrement; winding[sampleIndex] += seg->windingIncrement;
} }
else if( (seg->config == MG_MTL_BL || seg->config == MG_MTL_TR)
&&(pixelCoord.y > seg->box.y)) if(op->crossRight)
{ {
winding -= seg->windingIncrement; if( (seg->config == MG_MTL_BR || seg->config == MG_MTL_TL)
&&(sampleCoord.y > seg->box.w))
{
winding[sampleIndex] += seg->windingIncrement;
}
else if( (seg->config == MG_MTL_BL || seg->config == MG_MTL_TR)
&&(sampleCoord.y > seg->box.y))
{
winding[sampleIndex] -= seg->windingIncrement;
}
} }
} }
} }
opIndex = op->next; opIndex = op->next;
} }
bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding & 1)) float4 pixelColor = float4(0, 0, 0, 0);
||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding != 0)); float4 pathColor = pathBuffer[pathIndex].color;
if(filled) pathColor.rgb *= pathColor.a;
for(int sampleIndex=0; sampleIndex<MG_MTL_SAMPLE_COUNT; sampleIndex++)
{ {
float4 pathColor = pathBuffer[pathIndex].color; bool filled = (pathBuffer[pathIndex].cmd == MG_MTL_FILL && (winding[sampleIndex] & 1))
pathColor.rgb *= pathColor.a; ||(pathBuffer[pathIndex].cmd == MG_MTL_STROKE && (winding[sampleIndex] != 0));
color = color*(1-pathColor.a) + pathColor; if(filled)
{
color[sampleIndex] = color[sampleIndex]*(1-pathColor.a) + pathColor;
}
pixelColor += color[sampleIndex];
} }
pixelColor /= MG_MTL_SAMPLE_COUNT;
/* /*
if( (pixelCoord.x % tileSize[0] == 0) if( (pixelCoord.x % tileSize[0] == 0)
@ -1585,7 +1587,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
} }
*/ */
outTexture.write(color, uint2(pixelCoord)); outTexture.write(pixelColor, uint2(pixelCoord));
} }
//------------------------------------------------------------------------------------ //------------------------------------------------------------------------------------