From 0c8e698b6e412a8ef98bb348567b4d9511614bb1 Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Thu, 27 Jul 2023 16:10:30 +0200 Subject: [PATCH 1/4] [mtl canvas] don't trigger a new batch when setting the current image to nil (avoids creating a lot of batches when we just alternate between the same image and no image, eg painting multiple instances of the same image with a stroked border) --- src/mtl_renderer.h | 1 + src/mtl_renderer.m | 34 ++++++++++++++++++++-------------- src/mtl_renderer.metal | 3 +-- 3 files changed, 22 insertions(+), 16 deletions(-) diff --git a/src/mtl_renderer.h b/src/mtl_renderer.h index 788e1df..94e8910 100644 --- a/src/mtl_renderer.h +++ b/src/mtl_renderer.h @@ -23,6 +23,7 @@ typedef struct mg_mtl_path vector_float4 color; vector_float4 box; vector_float4 clip; + int texture; } mg_mtl_path; typedef enum { diff --git a/src/mtl_renderer.m b/src/mtl_renderer.m index 8cd5811..efa956f 100644 --- a/src/mtl_renderer.m +++ b/src/mtl_renderer.m @@ -249,6 +249,12 @@ void mg_mtl_encode_path(mg_mtl_canvas_backend* backend, mg_primitive* primitive, path->uvTransform = 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 = 1; + } + else + { + path->texture = 0; } int nTilesX = ((path->box.z - path->box.x)*scale - 1) / MG_MTL_TILE_SIZE + 1; @@ -1084,14 +1090,11 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend, [rasterEncoder setTexture:backend->outTexture atIndex:0]; - int useTexture = 0; if(image) { mg_mtl_image_data* mtlImage = (mg_mtl_image_data*)image; [rasterEncoder setTexture: mtlImage->texture atIndex: 1]; - useTexture = 1; } - [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); @@ -1233,20 +1236,23 @@ void mg_mtl_canvas_render(mg_canvas_backend* interface, { 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); + if(primitiveIndex && (primitive->attributes.image.h != currentImage.h)) + { + mg_image_data* imageData = mg_image_data_from_handle(currentImage); - mg_mtl_render_batch(backend, - surface, - imageData, - tileSize, - nTilesX, - nTilesY, - viewportSize, - scale); + mg_mtl_render_batch(backend, + surface, + imageData, + tileSize, + nTilesX, + nTilesY, + viewportSize, + scale); + } + currentImage = primitive->attributes.image; } - currentImage = primitive->attributes.image; if(primitive->path.count) { diff --git a/src/mtl_renderer.metal b/src/mtl_renderer.metal index 1013f85..d6d39e2 100644 --- a/src/mtl_renderer.metal +++ b/src/mtl_renderer.metal @@ -1512,7 +1512,6 @@ 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)]], uint2 threadGroupCoord [[threadgroup_position_in_grid]], @@ -1611,7 +1610,7 @@ kernel void mtl_raster(const device mg_mtl_screen_tile* screenTilesBuffer [[buff float4 nextColor = pathBuffer[pathIndex].color; nextColor.rgb *= nextColor.a; - if(useTexture[0]) + if(pathBuffer[pathIndex].texture) { constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); From 203adfd4706a8b2f03a1815ef19c6d9ca51c6365 Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Thu, 27 Jul 2023 16:46:10 +0200 Subject: [PATCH 2/4] [mtl canvas] allow using multiple textures per batch --- src/mtl_renderer.h | 4 +++ src/mtl_renderer.m | 76 +++++++++++++++++++++++++++--------------- src/mtl_renderer.metal | 9 ++--- 3 files changed, 59 insertions(+), 30 deletions(-) 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; From c983c39f44d810f472668492ba33d12540ebf627 Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Thu, 27 Jul 2023 18:11:34 +0200 Subject: [PATCH 3/4] [canvas] removed old begin/end/renderBatch function pointers --- src/gl_canvas.c | 2 +- src/graphics_surface.h | 11 ----------- 2 files changed, 1 insertion(+), 12 deletions(-) diff --git a/src/gl_canvas.c b/src/gl_canvas.c index bd82050..09b5a07 100644 --- a/src/gl_canvas.c +++ b/src/gl_canvas.c @@ -1661,10 +1661,10 @@ mg_canvas_backend* gl_canvas_backend_create(mg_wgl_surface* surface) //NOTE(martin): setup interface functions backend->interface.destroy = mg_gl_canvas_destroy; - backend->interface.render = mg_gl_canvas_render; backend->interface.imageCreate = mg_gl_canvas_image_create; backend->interface.imageDestroy = mg_gl_canvas_image_destroy; backend->interface.imageUploadRegion = mg_gl_canvas_image_upload_region; + backend->interface.render = mg_gl_canvas_render; surface->interface.prepare((mg_surface_data*)surface); diff --git a/src/graphics_surface.h b/src/graphics_surface.h index 0dca19b..d494d9a 100644 --- a/src/graphics_surface.h +++ b/src/graphics_surface.h @@ -78,14 +78,6 @@ typedef struct mg_image_data } mg_image_data; typedef void (*mg_canvas_backend_destroy_proc)(mg_canvas_backend* backend); -typedef void (*mg_canvas_backend_begin_proc)(mg_canvas_backend* backend, mg_color clearColor); -typedef void (*mg_canvas_backend_end_proc)(mg_canvas_backend* backend); -typedef void (*mg_canvas_backend_draw_batch_proc)(mg_canvas_backend* backend, - mg_image_data* imageData, - u32 vertexCount, - u32 shapeCount, - u32 indexCount); - typedef mg_image_data* (*mg_canvas_backend_image_create_proc)(mg_canvas_backend* backend, vec2 size); typedef void (*mg_canvas_backend_image_destroy_proc)(mg_canvas_backend* backend, mg_image_data* image); @@ -104,9 +96,6 @@ typedef void (*mg_canvas_backend_render_proc)(mg_canvas_backend* backend, typedef struct mg_canvas_backend { mg_canvas_backend_destroy_proc destroy; - mg_canvas_backend_begin_proc begin; - mg_canvas_backend_end_proc end; - mg_canvas_backend_draw_batch_proc drawBatch; mg_canvas_backend_image_create_proc imageCreate; mg_canvas_backend_image_destroy_proc imageDestroy; From b05caf984c1c5359565fd7496f378c4c59bbd85f Mon Sep 17 00:00:00 2001 From: martinfouilleul Date: Thu, 27 Jul 2023 20:07:07 +0200 Subject: [PATCH 4/4] [gl canvas] allow multiple textures per batch --- src/gl_canvas.c | 105 +++++++++++++++++++++-------------- src/glsl_shaders/common.glsl | 1 + src/glsl_shaders/merge.glsl | 5 +- src/glsl_shaders/raster.glsl | 53 ++++++++++++++++-- 4 files changed, 114 insertions(+), 50 deletions(-) diff --git a/src/gl_canvas.c b/src/gl_canvas.c index 09b5a07..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, @@ -1661,10 +1684,10 @@ mg_canvas_backend* gl_canvas_backend_create(mg_wgl_surface* surface) //NOTE(martin): setup interface functions backend->interface.destroy = mg_gl_canvas_destroy; + backend->interface.render = mg_gl_canvas_render; backend->interface.imageCreate = mg_gl_canvas_image_create; backend->interface.imageDestroy = mg_gl_canvas_image_destroy; backend->interface.imageUploadRegion = mg_gl_canvas_image_upload_region; - backend->interface.render = mg_gl_canvas_render; surface->interface.prepare((mg_surface_data*)surface); 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; sampleIndex