diff --git a/src/gl_canvas.c b/src/gl_canvas.c index bd82050..e5dd87c 100644 --- a/src/gl_canvas.c +++ b/src/gl_canvas.c @@ -30,7 +30,8 @@ typedef struct mg_gl_path vec4 box; vec4 clip; mg_gl_cmd cmd; - u8 pad[12]; + int textureID; + u8 pad[8]; } mg_gl_path; enum _mg_gl_seg_kind{ @@ -125,6 +126,7 @@ enum { MG_GL_INPUT_BUFFERS_COUNT = 3, MG_GL_TILE_SIZE = 16, MG_GL_MSAA_COUNT = 8, + MG_GL_MAX_IMAGES_PER_BATCH = 8, }; typedef struct mg_gl_mapped_buffer @@ -186,6 +188,7 @@ typedef struct mg_gl_canvas_backend int maxTileQueueCount; int maxSegmentCount; + int currentImageIndex; } mg_gl_canvas_backend; static void mg_update_path_extents(vec4* extents, vec2 p) @@ -362,6 +365,12 @@ void mg_gl_canvas_encode_path(mg_gl_canvas_backend* backend, mg_primitive* primi path->uvTransform[9] = uvTransform.m[5]; path->uvTransform[10] = 1; path->uvTransform[11] = 0; + + path->textureID = backend->currentImageIndex; + } + else + { + path->textureID = -1; } int nTilesX = ((path->box.z - path->box.x)*scale - 1) / MG_GL_TILE_SIZE + 1; @@ -1043,7 +1052,7 @@ void mg_gl_grow_buffer_if_needed(GLuint buffer, i32 wantedSize, const char* name void mg_gl_render_batch(mg_gl_canvas_backend* backend, mg_wgl_surface* surface, - mg_image_data* image, + mg_image* images, int tileSize, int nTilesX, int nTilesY, @@ -1220,18 +1229,7 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend, glUniform1i(0, tileSize); glUniform1f(1, scale); glUniform1i(2, pathCount); - - // if there's an image, don't cull solid tiles - if(image) - { - glUniform1i(3, 0); - } - else - { - glUniform1i(3, 1); - } - - glUniform1i(4, backend->pathBatchStart); + glUniform1i(3, backend->pathBatchStart); glDispatchCompute(nTilesX, nTilesY, 1); glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT); @@ -1269,20 +1267,21 @@ void mg_gl_render_batch(mg_gl_canvas_backend* backend, glBindImageTexture(0, backend->outTexture, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_RGBA8); - if(image) + for(int i=0; itexture); - glUniform1ui(2, 1); - } - else - { - glUniform1ui(2, 0); + if(images[i].h) + { + mg_gl_image* image = (mg_gl_image*)mg_image_data_from_handle(images[i]); + if(image) + { + glActiveTexture(GL_TEXTURE1+i); + glBindTexture(GL_TEXTURE_2D, image->texture); + } + } } - glUniform1i(3, backend->pathBatchStart); - glUniform1ui(4, maxWorkGroupCount); + glUniform1i(2, backend->pathBatchStart); + glUniform1ui(3, maxWorkGroupCount); glBindBuffer(GL_DISPATCH_INDIRECT_BUFFER, backend->rasterDispatchBuffer); glDispatchComputeIndirect(0); @@ -1397,28 +1396,53 @@ void mg_gl_canvas_render(mg_canvas_backend* interface, //NOTE: encode and render batches vec2 currentPos = {0}; - mg_image currentImage = mg_image_nil(); - + mg_image images[MG_GL_MAX_IMAGES_PER_BATCH] = {0}; + int imageCount = 0; backend->eltCount = 0; for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++) { mg_primitive* primitive = &primitives[primitiveIndex]; - if(primitiveIndex && (primitive->attributes.image.h != currentImage.h)) + if(primitive->attributes.image.h != 0) { - mg_image_data* imageData = mg_image_data_from_handle(currentImage); + backend->currentImageIndex = -1; + for(int i=0; iattributes.image.h) + { + backend->currentImageIndex = i; + } + } + if(backend->currentImageIndex <= 0) + { + if(imageCountattributes.image; + backend->currentImageIndex = imageCount; + imageCount++; + } + else + { + mg_gl_render_batch(backend, + surface, + images, + tileSize, + nTilesX, + nTilesY, + viewportSize, + scale); - mg_gl_render_batch(backend, - surface, - imageData, - tileSize, - nTilesX, - nTilesY, - viewportSize, - scale); + images[0] = primitive->attributes.image; + backend->currentImageIndex = 0; + imageCount = 1; + } + } + } + else + { + backend->currentImageIndex = -1; } - currentImage = primitive->attributes.image; if(primitive->path.count) { @@ -1470,10 +1494,9 @@ void mg_gl_canvas_render(mg_canvas_backend* interface, } } - mg_image_data* imageData = mg_image_data_from_handle(currentImage); mg_gl_render_batch(backend, surface, - imageData, + images, tileSize, nTilesX, nTilesY, diff --git a/src/glsl_shaders/common.glsl b/src/glsl_shaders/common.glsl index 3dd72c0..49582f8 100644 --- a/src/glsl_shaders/common.glsl +++ b/src/glsl_shaders/common.glsl @@ -34,6 +34,7 @@ struct mg_gl_path vec4 box; vec4 clip; int cmd; + int textureID; }; struct mg_gl_path_elt diff --git a/src/glsl_shaders/merge.glsl b/src/glsl_shaders/merge.glsl index 3a8ffa9..d0a5494 100644 --- a/src/glsl_shaders/merge.glsl +++ b/src/glsl_shaders/merge.glsl @@ -43,8 +43,7 @@ layout(binding = 6) coherent restrict buffer screenTilesCountBufferSSBO layout(location = 0) uniform int tileSize; layout(location = 1) uniform float scale; layout(location = 2) uniform int pathCount; -layout(location = 3) uniform int cullSolidTiles; -layout(location = 4) uniform int pathBufferStart; +layout(location = 3) uniform int pathBufferStart; void main() { @@ -130,7 +129,7 @@ void main() tileOpBuffer.elements[pathOpIndex].kind = MG_GL_OP_FILL; if( pathBuffer.elements[pathBufferStart + pathIndex].color.a == 1 - && cullSolidTiles != 0) + && pathBuffer.elements[pathBufferStart + pathIndex].textureID < 0) { screenTilesBuffer.elements[tileIndex].first = pathOpIndex; } diff --git a/src/glsl_shaders/raster.glsl b/src/glsl_shaders/raster.glsl index e7f6188..2b5f32e 100644 --- a/src/glsl_shaders/raster.glsl +++ b/src/glsl_shaders/raster.glsl @@ -32,12 +32,19 @@ layout(binding = 4) restrict readonly buffer screenTilesCountBufferSSBO layout(location = 0) uniform float scale; layout(location = 1) uniform int msaaSampleCount; -layout(location = 2) uniform uint useTexture; -layout(location = 3) uniform int pathBufferStart; -layout(location = 4) uniform uint maxWorkGroupCount; +layout(location = 2) uniform int pathBufferStart; +layout(location = 3) uniform uint maxWorkGroupCount; layout(rgba8, binding = 0) uniform restrict writeonly image2D outTexture; -layout(binding = 1) uniform sampler2D srcTexture; + +layout(binding = 1) uniform sampler2D srcTexture0; +layout(binding = 2) uniform sampler2D srcTexture1; +layout(binding = 3) uniform sampler2D srcTexture2; +layout(binding = 4) uniform sampler2D srcTexture3; +layout(binding = 5) uniform sampler2D srcTexture4; +layout(binding = 6) uniform sampler2D srcTexture5; +layout(binding = 7) uniform sampler2D srcTexture6; +layout(binding = 8) uniform sampler2D srcTexture7; void main() { @@ -146,15 +153,49 @@ void main() vec4 nextColor = pathBuffer.elements[pathBufferStart + pathIndex].color; nextColor.rgb *= nextColor.a; - if(useTexture != 0) + int textureID = pathBuffer.elements[pathBufferStart+pathIndex].textureID; + if(textureID >= 0) { vec4 texColor = vec4(0); + for(int sampleIndex = 0; sampleIndexuvTransform = simd_matrix(simd_make_float3(uvTransform.m[0]/scale, uvTransform.m[3]/scale, 0), simd_make_float3(uvTransform.m[1]/scale, uvTransform.m[4]/scale, 0), simd_make_float3(uvTransform.m[2], uvTransform.m[5], 1)); + } + 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; @@ -932,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, @@ -1084,14 +1088,17 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend, [rasterEncoder setTexture:backend->outTexture atIndex:0]; - int useTexture = 0; - if(image) + for(int i=0; itexture atIndex: 1]; - useTexture = 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]; + } + } } - [rasterEncoder setBytes: &useTexture length:sizeof(int) atIndex: 9]; MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1); MTLSize rasterGroupSize = MTLSizeMake(MG_MTL_TILE_SIZE, MG_MTL_TILE_SIZE, 1); @@ -1226,27 +1233,52 @@ 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++) { mg_primitive* primitive = &primitives[primitiveIndex]; - if(primitiveIndex && (primitive->attributes.image.h != currentImage.h)) + if(primitive->attributes.image.h != 0) { - mg_image_data* imageData = mg_image_data_from_handle(currentImage); + backend->currentImageIndex = -1; + for(int i=0; iattributes.image.h) + { + backend->currentImageIndex = i; + } + } + if(backend->currentImageIndex <= 0) + { + if(imageCountattributes.image; + backend->currentImageIndex = imageCount; + imageCount++; + } + else + { + mg_mtl_render_batch(backend, + surface, + images, + tileSize, + nTilesX, + nTilesY, + viewportSize, + scale); - mg_mtl_render_batch(backend, - surface, - imageData, - tileSize, - nTilesX, - nTilesY, - viewportSize, - scale); + images[0] = primitive->attributes.image; + backend->currentImageIndex = 0; + imageCount = 1; + } + } + } + else + { + backend->currentImageIndex = -1; } - currentImage = primitive->attributes.image; if(primitive->path.count) { @@ -1296,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 1013f85..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; } @@ -1512,9 +1512,8 @@ kernel void mtl_raster(const device mg_mtl_screen_tile* screenTilesBuffer [[buff constant int* sampleCountBuffer [[buffer(6)]], device char* logBuffer [[buffer(7)]], device atomic_int* logOffsetBuffer [[buffer(8)]], - constant int* useTexture [[buffer(9)]], 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]]) { @@ -1611,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(useTexture[0]) + 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); @@ -1622,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;