[mtl canvas] allow using multiple textures per batch

src/mtl_renderer.h

@@ -104,4 +104,8 @@ typedef struct mg_mtl_screen_tile
 
 } mg_mtl_screen_tile;
 
+enum {
+	MG_MTL_MAX_IMAGES_PER_BATCH = 30
+};
+
 #endif //__MTL_RENDERER_H_

src/mtl_renderer.m

@@ -71,6 +71,8 @@ typedef struct mg_mtl_canvas_backend
 	int maxTileQueueCount;
 	int maxSegmentCount;
 
+	int currentImageIndex;
+
 } mg_mtl_canvas_backend;
 
 typedef struct mg_mtl_image_data
@@ -250,12 +252,8 @@ void mg_mtl_encode_path(mg_mtl_canvas_backend* backend, mg_primitive* primitive,
 		                                simd_make_float3(uvTransform.m[1]/scale, uvTransform.m[4]/scale, 0),
 		                                simd_make_float3(uvTransform.m[2], uvTransform.m[5], 1));
 
-		path->texture = 1;
-	}
-	else
-	{
-		path->texture = 0;
 	}
+	path->texture = backend->currentImageIndex;
 
 	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;
@@ -938,7 +936,7 @@ void mg_mtl_grow_buffer_if_needed(mg_mtl_canvas_backend* backend, id<MTLBuffer>*
 
 void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
                          mg_mtl_surface* surface,
-                         mg_image_data* image,
+                         mg_image* images,
                          int tileSize,
                          int nTilesX,
                          int nTilesY,
@@ -1090,10 +1088,16 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
 
 		[rasterEncoder setTexture:backend->outTexture atIndex:0];
 
+		for(int i=0; i<MG_MTL_MAX_IMAGES_PER_BATCH; i++)
+		{
+			if(images[i].h)
+			{
+				mg_mtl_image_data* image = (mg_mtl_image_data*)mg_image_data_from_handle(images[i]);
 				if(image)
 				{
-			mg_mtl_image_data* mtlImage = (mg_mtl_image_data*)image;
+					[rasterEncoder setTexture: image->texture atIndex: 1+i];
-			[rasterEncoder setTexture: mtlImage->texture atIndex: 1];
+				}
+			}
 		}
 
 		MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1);
@@ -1229,8 +1233,8 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
 
 	//NOTE: encode and render batches
 	vec2 currentPos = {0};
-
+	mg_image images[MG_MTL_MAX_IMAGES_PER_BATCH] = {0};
-	mg_image currentImage = mg_image_nil();
+	int imageCount = 0;
 
 	for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++)
 	{
@@ -1238,20 +1242,42 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
 
 		if(primitive->attributes.image.h != 0)
 		{
-			if(primitiveIndex && (primitive->attributes.image.h != currentImage.h))
+			backend->currentImageIndex = -1;
+			for(int i=0; i<imageCount; i++)
+			{
+				if(images[i].h == primitive->attributes.image.h)
+				{
+					backend->currentImageIndex = i;
+				}
+			}
+			if(backend->currentImageIndex <= 0)
+			{
+				if(imageCount<MG_MTL_MAX_IMAGES_PER_BATCH)
+				{
+					images[imageCount] = primitive->attributes.image;
+					backend->currentImageIndex = imageCount;
+					imageCount++;
+				}
+				else
 				{
-				mg_image_data* imageData = mg_image_data_from_handle(currentImage);
-
 					mg_mtl_render_batch(backend,
 					                    surface,
-				                    imageData,
+					                    images,
 					                    tileSize,
 					                    nTilesX,
 					                    nTilesY,
 					                    viewportSize,
 					                    scale);
+
+					images[0] = primitive->attributes.image;
+					backend->currentImageIndex = 0;
+					imageCount = 1;
 				}
-			currentImage = primitive->attributes.image;
+			}
+		}
+		else
+		{
+			backend->currentImageIndex = -1;
 		}
 
 		if(primitive->path.count)
@@ -1302,11 +1328,9 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface,
 		}
 	}
 
-	mg_image_data* imageData = mg_image_data_from_handle(currentImage);
-
 	mg_mtl_render_batch(backend,
 	                    surface,
-	                    imageData,
+	                    images,
 	                    tileSize,
 	                    nTilesX,
 	                    nTilesY,

src/mtl_renderer.metal

@@ -1450,7 +1450,7 @@ kernel void mtl_merge(constant int* pathCount [[buffer(0)]],
 					{
 						pathOp->kind = MG_MTL_OP_FILL;
 
-						if(pathBuffer[pathIndex].color.a == 1)
+						if(pathBuffer[pathIndex].color.a == 1 && pathBuffer[pathIndex].texture < 0)
 						{
 							screenTilesBuffer[tileIndex].first = pathOpIndex;
 						}
@@ -1513,7 +1513,7 @@ kernel void mtl_raster(const device mg_mtl_screen_tile* screenTilesBuffer [[buff
                        device char* logBuffer [[buffer(7)]],
                        device atomic_int* logOffsetBuffer [[buffer(8)]],
                        texture2d<float, access::write> outTexture [[texture(0)]],
-                       texture2d<float> srcTexture [[texture(1)]],
+                       array<texture2d<float>, MG_MTL_MAX_IMAGES_PER_BATCH> srcTextures [[texture(1)]],
                        uint2  threadGroupCoord [[threadgroup_position_in_grid]],
                        uint2 localCoord [[thread_position_in_threadgroup]])
 {
@@ -1610,7 +1610,8 @@ kernel void mtl_raster(const device mg_mtl_screen_tile* screenTilesBuffer [[buff
 			float4 nextColor = pathBuffer[pathIndex].color;
 			nextColor.rgb *= nextColor.a;
 
-			if(pathBuffer[pathIndex].texture)
+			int textureIndex = pathBuffer[pathIndex].texture;
+			if(textureIndex >= 0 && textureIndex < MG_MTL_MAX_IMAGES_PER_BATCH)
 			{
 				constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear);
 
@@ -1621,7 +1622,7 @@ kernel void mtl_raster(const device mg_mtl_screen_tile* screenTilesBuffer [[buff
 					float3 ph = float3(sampleCoord.xy, 1);
 					float2 uv = (pathBuffer[pathIndex].uvTransform * ph).xy;
 
-					texColor += srcTexture.sample(smp, uv);
+					texColor += srcTextures[textureIndex].sample(smp, uv);
 				}
 				texColor /= srcSampleCount;
 				texColor.rgb *= texColor.a;