diff --git a/src/mtl_renderer.h b/src/mtl_renderer.h index 94e8910..287ed4c 100644 --- a/src/mtl_renderer.h +++ b/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_ diff --git a/src/mtl_renderer.m b/src/mtl_renderer.m index efa956f..61de9db 100644 --- a/src/mtl_renderer.m +++ b/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* 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]; - if(image) + for(int i=0; itexture atIndex: 1]; + if(images[i].h) + { + mg_mtl_image_data* image = (mg_mtl_image_data*)mg_image_data_from_handle(images[i]); + if(image) + { + [rasterEncoder setTexture: image->texture atIndex: 1+i]; + } + } } 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 currentImage = mg_image_nil(); + mg_image images[MG_MTL_MAX_IMAGES_PER_BATCH] = {0}; + 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; iattributes.image.h) + { + backend->currentImageIndex = i; + } } - currentImage = primitive->attributes.image; + if(backend->currentImageIndex <= 0) + { + if(imageCountattributes.image; + backend->currentImageIndex = imageCount; + imageCount++; + } + else + { + mg_mtl_render_batch(backend, + surface, + images, + tileSize, + nTilesX, + nTilesY, + viewportSize, + scale); + + images[0] = primitive->attributes.image; + backend->currentImageIndex = 0; + imageCount = 1; + } + } + } + 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, diff --git a/src/mtl_renderer.metal b/src/mtl_renderer.metal index d6d39e2..55c9179 100644 --- a/src/mtl_renderer.metal +++ b/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 outTexture [[texture(0)]], - texture2d srcTexture [[texture(1)]], + array, 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;