From ee3e55dadd6a44a2635112e714cc39d72cd0e6cb Mon Sep 17 00:00:00 2001
From: Martin Fouilleul <martinfouilleul@gmail.com>
Date: Sat, 8 Apr 2023 12:11:38 +0200
Subject: [PATCH] [mtl renderer] fix msaa sample positions

---
 build.sh               |  2 +-
 examples/tiger/main.c  |  2 +-
 src/mtl_renderer.m     |  2 +-
 src/mtl_renderer.metal | 41 +++++++++++++----------------------------
 4 files changed, 16 insertions(+), 31 deletions(-)

diff --git a/build.sh b/build.sh
index 5fcf322..21ac3f2 100755
--- a/build.sh
+++ b/build.sh
@@ -1,6 +1,6 @@
 #!/bin/bash
 
-DEBUG_FLAGS="-g -DDEBUG -DLOG_COMPILE_DEBUG"
+DEBUG_FLAGS="-g -O2 -DDEBUG -DLOG_COMPILE_DEBUG"
 #DEBUG_FLAGS="-O3"
 
 #--------------------------------------------------------------
diff --git a/examples/tiger/main.c b/examples/tiger/main.c
index 753c887..56a90ff 100644
--- a/examples/tiger/main.c
+++ b/examples/tiger/main.c
@@ -200,8 +200,8 @@ int main()
 		if(singlePath)
 		{
 			printf("display single path %i\n", singlePathIndex);
+			printf("viewpos = (%f, %f), zoom = %f\n", startX, startY, zoom);
 		}
-		printf("viewpos = (%f, %f), zoom = %f\n", startX, startY, zoom);
 
 		mg_matrix_pop();
 
diff --git a/src/mtl_renderer.m b/src/mtl_renderer.m
index 2c411a3..1b5dda5 100644
--- a/src/mtl_renderer.m
+++ b/src/mtl_renderer.m
@@ -20,7 +20,7 @@
 
 const int MG_MTL_INPUT_BUFFERS_COUNT = 3,
           MG_MTL_TILE_SIZE = 16,
-          MG_MTL_MSAA_COUNT = 1;
+          MG_MTL_MSAA_COUNT = 8;
 
 typedef struct mg_mtl_canvas_backend
 {
diff --git a/src/mtl_renderer.metal b/src/mtl_renderer.metal
index 2e453ec..aaebf52 100644
--- a/src/mtl_renderer.metal
+++ b/src/mtl_renderer.metal
@@ -1271,22 +1271,6 @@ kernel void mtl_segment_setup(constant int* elementCount [[buffer(0)]],
                               uint eltIndex [[thread_position_in_grid]])
 {
 	const device mg_mtl_path_elt* elt = &elementBuffer[eltIndex];
-
-	//28
-	// 125
-	// 112
-//	if(elt->pathIndex != 96)
-	{
-//		return;
-	}
-
-	/*
-	if(elt->localEltIndex != 21)// && elt->localEltIndex != 22)
-	{
-		return;
-	}
-	*/
-
 	const device mg_mtl_path_queue* pathQueue = &pathQueueBuffer[elt->pathIndex];
 	device mg_mtl_tile_queue* tileQueues = &tileQueueBuffer[pathQueue->tileQueues];
 
@@ -1472,8 +1456,8 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 	                       .offset = logOffsetBuffer,
 	                       .enabled = true};
 */
-	float2 pixelCoord = float2(threadCoord);
-	int2 tileCoord = int2(threadCoord) / tileSize[0];
+	uint2 pixelCoord = threadCoord;
+	int2 tileCoord = int2(pixelCoord) / tileSize[0];
 	int nTilesX = (int(gridSize.x) + tileSize[0] - 1)/tileSize[0];
 	int tileIndex = tileCoord.y * nTilesX + tileCoord.x;
 
@@ -1483,23 +1467,24 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 	const int MG_MTL_MAX_SAMPLE_COUNT = 8;
 	float2 sampleCoords[MG_MTL_MAX_SAMPLE_COUNT];
 	int sampleCount = sampleCountBuffer[0];
+	float2 centerCoord = float2(pixelCoord) + float2(0.5, 0.5);
 
 	if(sampleCount == 8)
 	{
 		sampleCount = 8;
-		sampleCoords[0] = pixelCoord + float2(1, 3)/16;
-		sampleCoords[1] = pixelCoord + float2(-1, -3)/16;
-		sampleCoords[2] = pixelCoord + float2(5, -1)/16;
-		sampleCoords[3] = pixelCoord + float2(-3, 5)/16;
-		sampleCoords[4] = pixelCoord + float2(-5, -5)/16;
-		sampleCoords[5] = pixelCoord + float2(-7, 1)/16;
-		sampleCoords[6] = pixelCoord + float2(3, -7)/16;
-		sampleCoords[7] = pixelCoord + float2(7, 7)/16;
+		sampleCoords[0] = centerCoord + float2(1, 3)/16;
+		sampleCoords[1] = centerCoord + float2(-1, -3)/16;
+		sampleCoords[2] = centerCoord + float2(5, -1)/16;
+		sampleCoords[3] = centerCoord + float2(-3, 5)/16;
+		sampleCoords[4] = centerCoord + float2(-5, -5)/16;
+		sampleCoords[5] = centerCoord + float2(-7, 1)/16;
+		sampleCoords[6] = centerCoord + float2(3, -7)/16;
+		sampleCoords[7] = centerCoord + float2(7, 7)/16;
 	}
 	else
 	{
 		sampleCount = 1;
-		sampleCoords[0] = pixelCoord;
+		sampleCoords[0] = centerCoord;
 	}
 
 	float4 color[MG_MTL_MAX_SAMPLE_COUNT] = {0};
@@ -1584,7 +1569,7 @@ kernel void mtl_raster(const device int* screenTilesBuffer [[buffer(0)]],
 	}
 	//*/
 
-	outTexture.write(pixelColor, uint2(pixelCoord));
+	outTexture.write(pixelColor, pixelCoord);
 }
 
 //------------------------------------------------------------------------------------