From f6a992e5f4d4b55cd69c684c06e9a0c90caa05c4 Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Tue, 28 Mar 2023 13:09:48 +0200 Subject: [PATCH] [mtl canvas] very first draft of new canvas backend based on edge-counting instead of triangle overlap --- build.sh | 6 +- examples/polygon/build.sh | 11 ++ examples/polygon/main.c | 112 +++++++++++++++ src/graphics.c | 88 +----------- src/graphics_internal.h | 91 ++++++++++++ src/milepost.m | 3 +- src/mtl_renderer.h | 53 +++++++ src/mtl_renderer.m | 285 ++++++++++++++++++++++++++++++++++++++ src/mtl_renderer.metal | 101 ++++++++++++++ 9 files changed, 665 insertions(+), 85 deletions(-) create mode 100755 examples/polygon/build.sh create mode 100644 examples/polygon/main.c create mode 100644 src/mtl_renderer.h create mode 100644 src/mtl_renderer.m create mode 100644 src/mtl_renderer.metal diff --git a/build.sh b/build.sh index d658aba..3ac7a17 100755 --- a/build.sh +++ b/build.sh @@ -1,6 +1,6 @@ #!/bin/bash -DEBUG_FLAGS="-g -O2 -DDEBUG -DLOG_COMPILE_DEBUG" +DEBUG_FLAGS="-g -DDEBUG -DLOG_COMPILE_DEBUG" #DEBUG_FLAGS="-O3" #-------------------------------------------------------------- @@ -55,8 +55,8 @@ fi if [ $target = 'lib' ] ; then # compile metal shader - xcrun -sdk macosx metal $shaderFlagParam -c -o $BINDIR/mtl_shader.air $SRCDIR/mtl_shader.metal - xcrun -sdk macosx metallib -o $RESDIR/mtl_shader.metallib $BINDIR/mtl_shader.air + xcrun -sdk macosx metal $shaderFlagParam -c -o $BINDIR/mtl_renderer.air $SRCDIR/mtl_renderer.metal + xcrun -sdk macosx metallib -o $RESDIR/mtl_renderer.metallib $BINDIR/mtl_renderer.air # compile milepost. We use one compilation unit for all C code, and one compilation # unit for all ObjectiveC code diff --git a/examples/polygon/build.sh b/examples/polygon/build.sh new file mode 100755 index 0000000..f9ccf43 --- /dev/null +++ b/examples/polygon/build.sh @@ -0,0 +1,11 @@ +#!/bin/bash + +BINDIR=../../bin +RESDIR=../../resources +SRCDIR=../../src + +INCLUDES="-I$SRCDIR -I$SRCDIR/util -I$SRCDIR/platform -I$SRCDIR/app" +LIBS="-L$BINDIR -lmilepost" +FLAGS="-mmacos-version-min=10.15.4 -DDEBUG -DLOG_COMPILE_DEBUG" + +clang -g $FLAGS $LIBS $INCLUDES -o $BINDIR/example_polygon main.c diff --git a/examples/polygon/main.c b/examples/polygon/main.c new file mode 100644 index 0000000..cedefea --- /dev/null +++ b/examples/polygon/main.c @@ -0,0 +1,112 @@ +/************************************************************//** +* +* @file: main.cpp +* @author: Martin Fouilleul +* @date: 30/07/2022 +* @revision: +* +*****************************************************************/ +#include +#include +#include + +#define _USE_MATH_DEFINES //NOTE: necessary for MSVC +#include + +#include"milepost.h" + +#define LOG_SUBSYSTEM "Main" + +int main() +{ + LogLevel(LOG_LEVEL_WARNING); + + mp_init(); + mp_clock_init(); //TODO put that in mp_init()? + + mp_rect windowRect = {.x = 100, .y = 100, .w = 810, .h = 610}; + mp_window window = mp_window_create(windowRect, "test", 0); + + mp_rect contentRect = mp_window_get_content_rect(window); + + //NOTE: create surface + mg_surface surface = mg_surface_create_for_window(window, MG_BACKEND_DEFAULT); + mg_surface_swap_interval(surface, 0); + + //TODO: create canvas + mg_canvas canvas = mg_canvas_create(surface); + + if(mg_canvas_is_nil(canvas)) + { + printf("Error: couldn't create canvas\n"); + return(-1); + } + + // start app + mp_window_bring_to_front(window); + mp_window_focus(window); + + f64 frameTime = 0; + + while(!mp_should_quit()) + { + f64 startTime = mp_get_time(MP_CLOCK_MONOTONIC); + + mp_pump_events(0); + mp_event event = {0}; + while(mp_next_event(&event)) + { + switch(event.type) + { + case MP_EVENT_WINDOW_CLOSE: + { + mp_request_quit(); + } break; + + default: + break; + } + } + + mg_surface_prepare(surface); + + // background + mg_set_color_rgba(0, 1, 1, 1); + mg_clear(); + + mg_move_to(100, 100); + mg_line_to(150, 150); + mg_line_to(100, 200); + mg_line_to(50, 150); + mg_close_path(); + mg_set_color_rgba(1, 0, 0, 1); + mg_fill(); + + mg_move_to(200, 100); + mg_line_to(400, 100); + mg_line_to(400, 200); + mg_line_to(200, 200); + mg_close_path(); + mg_set_color_rgba(0, 1, 0, 1); + mg_fill(); + + + printf("Milepost vector graphics test program (frame time = %fs, fps = %f)...\n", + frameTime, + 1./frameTime); + + mg_flush(); + mg_surface_present(surface); + + mem_arena_clear(mem_scratch()); + frameTime = mp_get_time(MP_CLOCK_MONOTONIC) - startTime; + } + + mg_canvas_destroy(canvas); + mg_surface_destroy(surface); + mp_window_destroy(window); + + mp_terminate(); + + return(0); +} diff --git a/src/graphics.c b/src/graphics.c index 1f405a4..ad9bfbf 100644 --- a/src/graphics.c +++ b/src/graphics.c @@ -21,87 +21,6 @@ #define LOG_SUBSYSTEM "Graphics" -//------------------------------------------------------------------------ -// canvas structs -//------------------------------------------------------------------------ -typedef enum { MG_PATH_MOVE, - MG_PATH_LINE, - MG_PATH_QUADRATIC, - MG_PATH_CUBIC } mg_path_elt_type; - -typedef struct mg_path_elt -{ - mg_path_elt_type type; - vec2 p[3]; - -} mg_path_elt; - -typedef struct mg_path_descriptor -{ - u32 startIndex; - u32 count; - vec2 startPoint; - -} mg_path_descriptor; - -typedef struct mg_attributes -{ - f32 width; - f32 tolerance; - mg_color color; - mg_joint_type joint; - f32 maxJointExcursion; - mg_cap_type cap; - - mg_font font; - f32 fontSize; - - mg_image image; - mp_rect srcRegion; - - mg_mat2x3 transform; - mp_rect clip; - -} mg_attributes; - -typedef struct mg_rounded_rect -{ - f32 x; - f32 y; - f32 w; - f32 h; - f32 r; -} mg_rounded_rect; - -typedef enum { MG_CMD_FILL, - MG_CMD_STROKE, - MG_CMD_RECT_FILL, - MG_CMD_RECT_STROKE, - MG_CMD_ROUND_RECT_FILL, - MG_CMD_ROUND_RECT_STROKE, - MG_CMD_ELLIPSE_FILL, - MG_CMD_ELLIPSE_STROKE, - MG_CMD_JUMP, - MG_CMD_CLIP_PUSH, - MG_CMD_CLIP_POP, - } mg_primitive_cmd; - -typedef struct mg_primitive -{ - mg_primitive_cmd cmd; - mg_attributes attributes; - - union - { - mg_path_descriptor path; - mp_rect rect; - mg_rounded_rect roundedRect; - utf32 codePoint; - u32 jump; - }; - -} mg_primitive; - typedef struct mg_glyph_map_entry { unicode_range range; @@ -3055,6 +2974,13 @@ void mg_flush_commands(int primitiveCount, mg_primitive* primitives, mg_path_elt return; } + + if(canvas->backend && canvas->backend->render) + { + canvas->backend->render(canvas->backend, primitiveCount, primitives, canvas->path.startIndex + canvas->path.count, pathElements); + return; + } + u32 nextIndex = 0; mg_reset_shape_index(canvas); diff --git a/src/graphics_internal.h b/src/graphics_internal.h index 90fdfd4..d38e8da 100644 --- a/src/graphics_internal.h +++ b/src/graphics_internal.h @@ -123,6 +123,93 @@ typedef void (*mg_canvas_backend_image_upload_region_proc)(mg_canvas_backend* ba mp_rect region, u8* pixels); + +//------------------------------------------------------------------------ +// canvas structs +//------------------------------------------------------------------------ +typedef enum { MG_PATH_MOVE, + MG_PATH_LINE, + MG_PATH_QUADRATIC, + MG_PATH_CUBIC } mg_path_elt_type; + +typedef struct mg_path_elt +{ + mg_path_elt_type type; + vec2 p[3]; + +} mg_path_elt; + +typedef struct mg_path_descriptor +{ + u32 startIndex; + u32 count; + vec2 startPoint; + +} mg_path_descriptor; + +typedef struct mg_attributes +{ + f32 width; + f32 tolerance; + mg_color color; + mg_joint_type joint; + f32 maxJointExcursion; + mg_cap_type cap; + + mg_font font; + f32 fontSize; + + mg_image image; + mp_rect srcRegion; + + mg_mat2x3 transform; + mp_rect clip; + +} mg_attributes; + +typedef struct mg_rounded_rect +{ + f32 x; + f32 y; + f32 w; + f32 h; + f32 r; +} mg_rounded_rect; + +typedef enum { MG_CMD_FILL, + MG_CMD_STROKE, + MG_CMD_RECT_FILL, + MG_CMD_RECT_STROKE, + MG_CMD_ROUND_RECT_FILL, + MG_CMD_ROUND_RECT_STROKE, + MG_CMD_ELLIPSE_FILL, + MG_CMD_ELLIPSE_STROKE, + MG_CMD_JUMP, + MG_CMD_CLIP_PUSH, + MG_CMD_CLIP_POP, + } mg_primitive_cmd; + +typedef struct mg_primitive +{ + mg_primitive_cmd cmd; + mg_attributes attributes; + + union + { + mg_path_descriptor path; + mp_rect rect; + mg_rounded_rect roundedRect; + u32 jump; + }; + +} mg_primitive; + +typedef void (*mg_canvas_backend_render_proc)(mg_canvas_backend* backend, + u32 primitiveCount, + mg_primitive* primitives, + u32 eltCount, + mg_path_elt* pathElements); + typedef struct mg_canvas_backend { mg_vertex_layout vertexLayout; @@ -135,6 +222,10 @@ typedef struct mg_canvas_backend mg_canvas_backend_image_create_proc imageCreate; mg_canvas_backend_image_destroy_proc imageDestroy; mg_canvas_backend_image_upload_region_proc imageUploadRegion; + + + mg_canvas_backend_render_proc render; + } mg_canvas_backend; #ifdef __cplusplus diff --git a/src/milepost.m b/src/milepost.m index a064fe3..df7dd7a 100644 --- a/src/milepost.m +++ b/src/milepost.m @@ -12,7 +12,8 @@ #if MG_COMPILE_BACKEND_METAL #include"mtl_surface.m" - #include"mtl_canvas.m" +// #include"mtl_canvas.m" + #include"mtl_renderer.m" #endif #if MG_COMPILE_BACKEND_GLES diff --git a/src/mtl_renderer.h b/src/mtl_renderer.h new file mode 100644 index 0000000..75109de --- /dev/null +++ b/src/mtl_renderer.h @@ -0,0 +1,53 @@ +/************************************************************//** +* +* @file: mtl_renderer.h +* @author: Martin Fouilleul +* @date: 01/08/2022 +* @revision: +* +*****************************************************************/ +#ifndef __MTL_RENDERER_H_ +#define __MTL_RENDERER_H_ + +#include + +typedef enum { + MG_MTL_LINE, +} mg_mtl_seg_kind; + +typedef enum { + MG_MTL_BL, // curve on bottom left + MG_MTL_BR, // curve on bottom right + MG_MTL_TL, // curve on top left + MG_MTL_TR // curve on top right +} mg_mtl_seg_config; + +typedef enum { + MG_MTL_FILL, +} mg_mtl_cmd; + +typedef struct mg_mtl_path_elt +{ + int pathIndex; + mg_mtl_seg_kind kind; + vector_float2 p[4]; +} mg_mtl_path_elt; + +typedef struct mg_mtl_segment +{ + int pathIndex; + mg_mtl_seg_config config; //TODO pack these + int windingIncrement; + vector_float4 box; +} mg_mtl_segment; + +typedef struct mg_mtl_path +{ + mg_mtl_cmd cmd; + vector_float4 color; + vector_float4 box; + +} mg_mtl_path; + + +#endif //__MTL_RENDERER_H_ diff --git a/src/mtl_renderer.m b/src/mtl_renderer.m new file mode 100644 index 0000000..809c19e --- /dev/null +++ b/src/mtl_renderer.m @@ -0,0 +1,285 @@ +/************************************************************//** +* +* @file: mtl_canvas.m +* @author: Martin Fouilleul +* @date: 12/07/2020 +* @revision: 24/01/2023 +* +*****************************************************************/ +#import +#import +#include + +#include"graphics_internal.h" +#include"macro_helpers.h" +#include"osx_app.h" + +#include"mtl_renderer.h" + +#define LOG_SUBSYSTEM "Graphics" + +typedef struct mg_mtl_canvas_backend +{ + mg_canvas_backend interface; + mg_surface surface; + + id rasterPipeline; + id blitPipeline; + + id outTexture; + + id pathBuffer; + id segmentBuffer; + +} mg_mtl_canvas_backend; + + +static void mg_update_path_extents(vec4* extents, vec2 p) +{ + extents->x = minimum(extents->x, p.x); + extents->y = minimum(extents->y, p.y); + extents->z = maximum(extents->z, p.x); + extents->w = maximum(extents->w, p.y); +} + +void mg_mtl_canvas_render(mg_canvas_backend* interface, + u32 primitiveCount, + mg_primitive* primitives, + u32 eltCount, + mg_path_elt* pathElements) +{ + mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; + + //TODO: update rolling buffers + mg_mtl_segment* segmentBufferData = (mg_mtl_segment*)[backend->segmentBuffer contents]; + mg_mtl_path* pathBufferData = (mg_mtl_path*)[backend->pathBuffer contents]; + + //NOTE: fill renderer input buffers + int segCount = 0; + int pathCount = 0; + vec2 currentPos = {0}; + + for(int primitiveIndex = 0; primitiveIndex < primitiveCount; primitiveIndex++) + { + mg_primitive* primitive = &primitives[primitiveIndex]; + if(primitive->cmd == MG_CMD_FILL && primitive->path.count) + { + vec4 pathExtents = (vec4){FLT_MAX, FLT_MAX, -FLT_MAX, -FLT_MAX}; + + for(int eltIndex = 0; + (eltIndex < primitive->path.count) && (primitive->path.startIndex + eltIndex < eltCount); + eltIndex++) + { + mg_path_elt* elt = &pathElements[primitive->path.startIndex + eltIndex]; + if(elt->type == MG_PATH_MOVE) + { + currentPos = elt->p[0]; + } + else if(elt->type == MG_PATH_LINE) + { + //NOTE: transform and push path elt + update primitive bounding box + vec2 p0 = mg_mat2x3_mul(primitive->attributes.transform, currentPos); + vec2 p3 = mg_mat2x3_mul(primitive->attributes.transform, elt->p[0]); + currentPos = elt->p[0]; + + if(p0.y != p3.y) + { + mg_mtl_segment* seg = &segmentBufferData[segCount]; + segCount++; + + seg->pathIndex = primitiveIndex; + seg->box = (vector_float4){minimum(p0.x, p3.x), + minimum(p0.y, p3.y), + maximum(p0.x, p3.x), + maximum(p0.y, p3.y)}; + + if( (p3.x > p0.x && p3.y < p0.y) + ||(p3.x <= p0.x && p3.y > p0.y)) + { + seg->config = MG_MTL_TR; + } + else if( (p3.x > p0.x && p3.y > p0.y) + ||(p3.x <= p0.x && p3.y < p0.y)) + { + seg->config = MG_MTL_BR; + } + + seg->windingIncrement = (p3.y > p0.y)? 1 : -1; + } + } + } + + //NOTE: push path + mg_mtl_path* path = &pathBufferData[pathCount]; + pathCount++; + + path->cmd = (mg_mtl_cmd)primitive->cmd; + path->box = (vector_float4){maximum(primitive->attributes.clip.x, pathExtents.x), + maximum(primitive->attributes.clip.y, pathExtents.y), + minimum(primitive->attributes.clip.x + primitive->attributes.clip.w, pathExtents.z), + minimum(primitive->attributes.clip.y + primitive->attributes.clip.h, pathExtents.w)}; + + path->color = (vector_float4){primitive->attributes.color.r, + primitive->attributes.color.g, + primitive->attributes.color.b, + primitive->attributes.color.a}; + + //TODO: compute uv transform + } + } + + mg_mtl_surface* surface = (mg_mtl_surface*)mg_surface_data_from_handle(backend->surface); + ASSERT(surface && surface->interface.backend == MG_BACKEND_METAL); + + mp_rect frame = mg_surface_get_frame(backend->surface); + f32 scale = surface->mtlLayer.contentsScale; + vec2 viewportSize = {frame.w * scale, frame.h * scale}; + + //NOTE: encode GPU commands + @autoreleasepool + { + mg_mtl_surface_acquire_command_buffer(surface); + + //NOTE: raster pass + id rasterEncoder = [surface->commandBuffer computeCommandEncoder]; + rasterEncoder.label = @"raster pass"; + [rasterEncoder setComputePipelineState: backend->rasterPipeline]; + + [rasterEncoder setBytes:&pathCount length:sizeof(int) atIndex:0]; + [rasterEncoder setBuffer:backend->pathBuffer offset:0 atIndex:1]; + [rasterEncoder setBytes:&segCount length:sizeof(int) atIndex:2]; + [rasterEncoder setBuffer:backend->segmentBuffer offset:0 atIndex:3]; + + [rasterEncoder setTexture:backend->outTexture atIndex:0]; + + MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1); + MTLSize rasterGroupSize = MTLSizeMake(16, 16, 1); + [rasterEncoder dispatchThreads: rasterGridSize threadsPerThreadgroup: rasterGroupSize]; + + [rasterEncoder endEncoding]; + + //NOTE: blit pass + mg_mtl_surface_acquire_drawable(surface); + if(surface->drawable != nil) + { + MTLViewport viewport = {0, 0, viewportSize.x, viewportSize.y, 0, 1}; + + //TODO: clear here? + MTLRenderPassDescriptor* renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor]; + renderPassDescriptor.colorAttachments[0].texture = surface->drawable.texture; + renderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionLoad; + renderPassDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore; + + id renderEncoder = [surface->commandBuffer renderCommandEncoderWithDescriptor:renderPassDescriptor]; + renderEncoder.label = @"blit pass"; + [renderEncoder setViewport: viewport]; + [renderEncoder setRenderPipelineState: backend->blitPipeline]; + [renderEncoder setFragmentTexture: backend->outTexture atIndex: 0]; + [renderEncoder drawPrimitives: MTLPrimitiveTypeTriangle + vertexStart: 0 + vertexCount: 3 ]; + [renderEncoder endEncoding]; + } + } +} + +void mg_mtl_canvas_destroy(mg_canvas_backend* interface) +{ + mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; + + @autoreleasepool + { + [backend->pathBuffer release]; + [backend->segmentBuffer release]; + } + + free(backend); +} + +const u32 MG_MTL_PATH_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_path), + MG_MTL_SEGMENT_BUFFER_SIZE = (4<<20)*sizeof(mg_mtl_segment); + +mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface) +{ + mg_mtl_canvas_backend* backend = 0; + + mg_surface_data* surfaceData = mg_surface_data_from_handle(surface); + if(surfaceData && surfaceData->backend == MG_BACKEND_METAL) + { + mg_mtl_surface* metalSurface = (mg_mtl_surface*)surfaceData; + + backend = malloc_type(mg_mtl_canvas_backend); + memset(backend, 0, sizeof(mg_mtl_canvas_backend)); + + backend->surface = surface; + + //NOTE(martin): setup interface functions + backend->interface.destroy = mg_mtl_canvas_destroy; + backend->interface.render = mg_mtl_canvas_render; + + @autoreleasepool{ + //NOTE: load metal library + str8 shaderPath = mp_app_get_resource_path(mem_scratch(), "../resources/mtl_renderer.metallib"); + NSString* metalFileName = [[NSString alloc] initWithBytes: shaderPath.ptr length:shaderPath.len encoding: NSUTF8StringEncoding]; + NSError* err = 0; + id library = [metalSurface->device newLibraryWithFile: metalFileName error:&err]; + if(err != nil) + { + const char* errStr = [[err localizedDescription] UTF8String]; + LOG_ERROR("error : %s\n", errStr); + return(0); + } + id rasterFunction = [library newFunctionWithName:@"mtl_raster"]; + id vertexFunction = [library newFunctionWithName:@"mtl_vertex_shader"]; + id fragmentFunction = [library newFunctionWithName:@"mtl_fragment_shader"]; + + //NOTE: create pipelines + NSError* error = NULL; + backend->rasterPipeline = [metalSurface->device newComputePipelineStateWithFunction: rasterFunction + error:&error]; + + MTLRenderPipelineDescriptor *pipelineStateDescriptor = [[MTLRenderPipelineDescriptor alloc] init]; + pipelineStateDescriptor.label = @"blit pipeline"; + pipelineStateDescriptor.vertexFunction = vertexFunction; + pipelineStateDescriptor.fragmentFunction = fragmentFunction; + pipelineStateDescriptor.colorAttachments[0].pixelFormat = metalSurface->mtlLayer.pixelFormat; + pipelineStateDescriptor.colorAttachments[0].blendingEnabled = YES; + pipelineStateDescriptor.colorAttachments[0].rgbBlendOperation = MTLBlendOperationAdd; + pipelineStateDescriptor.colorAttachments[0].sourceRGBBlendFactor = MTLBlendFactorOne; + pipelineStateDescriptor.colorAttachments[0].destinationRGBBlendFactor = MTLBlendFactorOneMinusSourceAlpha; + pipelineStateDescriptor.colorAttachments[0].alphaBlendOperation = MTLBlendOperationAdd; + pipelineStateDescriptor.colorAttachments[0].sourceAlphaBlendFactor = MTLBlendFactorOne; + pipelineStateDescriptor.colorAttachments[0].destinationAlphaBlendFactor = MTLBlendFactorOneMinusSourceAlpha; + + backend->blitPipeline = [metalSurface->device newRenderPipelineStateWithDescriptor: pipelineStateDescriptor error:&err]; + + //NOTE: create textures + mp_rect frame = mg_surface_get_frame(surface); + f32 scale = metalSurface->mtlLayer.contentsScale; + + MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; + texDesc.textureType = MTLTextureType2D; + texDesc.storageMode = MTLStorageModePrivate; + texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite; + texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm; + texDesc.width = frame.w * scale; + texDesc.height = frame.h * scale; + + backend->outTexture = [metalSurface->device newTextureWithDescriptor:texDesc]; + + //NOTE: create buffers + MTLResourceOptions bufferOptions = MTLResourceCPUCacheModeWriteCombined + | MTLResourceStorageModeShared; + + backend->pathBuffer = [metalSurface->device newBufferWithLength: MG_MTL_PATH_BUFFER_SIZE + options: bufferOptions]; + + backend->segmentBuffer = [metalSurface->device newBufferWithLength: MG_MTL_SEGMENT_BUFFER_SIZE + options: bufferOptions]; + } + + } + return((mg_canvas_backend*)backend); +} + +#undef LOG_SUBSYSTEM diff --git a/src/mtl_renderer.metal b/src/mtl_renderer.metal new file mode 100644 index 0000000..f46b273 --- /dev/null +++ b/src/mtl_renderer.metal @@ -0,0 +1,101 @@ + +#include +#include +#include + +#include"mtl_renderer.h" + +using namespace metal; + +kernel void mtl_raster(constant int* pathCount [[buffer(0)]], + const device mg_mtl_path* pathBuffer [[buffer(1)]], + constant int* segCount [[buffer(2)]], + const device mg_mtl_segment* segmentBuffer [[buffer(3)]], + texture2d outTexture [[texture(0)]], + uint2 threadCoord [[thread_position_in_grid]]) +{ + int2 pixelCoord = int2(threadCoord); + + float4 color = float4(0, 0, 0, 0); + int currentPath = 0; + int winding = 0; + + for(int segIndex = 0; segIndex < segCount[0]; segIndex++) + { + const device mg_mtl_segment* seg = &segmentBuffer[segIndex]; + + if(seg->pathIndex != currentPath) + { + //depending on winding number, update color + if(winding & 1) + { + float4 pathColor = pathBuffer[currentPath].color; + pathColor.rgb *= pathColor.a; + color = color*(1-pathColor.a) + pathColor; + } + currentPath = seg->pathIndex; + winding = 0; + } + + if(pixelCoord.y >= seg->box.y && pixelCoord.y < seg->box.w) + { + if(pixelCoord.x < seg->box.x) + { + winding += seg->windingIncrement; + } + else if(pixelCoord.x < seg->box.z) + { + /*TODO: if pixel is on opposite size of diagonal as curve on the right, increment + otherwise if not on same size of diagonal as curve, do implicit test + */ + float alpha = (seg->box.w - seg->box.y)/(seg->box.z - seg->box.x); + float ofs = seg->box.w - seg->box.y; + float dx = pixelCoord.x - seg->box.x; + float dy = pixelCoord.y - seg->box.y; + + if( (seg->config == MG_MTL_BR && dy > alpha*dx) + ||(seg->config == MG_MTL_TR && dy < ofs - alpha*dx)) + { + winding += seg->windingIncrement; + } + else if( !(seg->config == MG_MTL_TL && dy < alpha*dx) + && !(seg->config == MG_MTL_BL && dy > ofs - alpha*dx)) + { + //Need implicit test, but for lines, we only have config BR or TR, so the test is always negative for now + } + } + } + } + + if(winding & 1) + { + float4 pathColor = pathBuffer[currentPath].color; + pathColor.rgb *= pathColor.a; + color = color*(1-pathColor.a) + pathColor; + } + + outTexture.write(color, uint2(pixelCoord)); +} + +//------------------------------------------------------------------------------------ +// Blit shader +//------------------------------------------------------------------------------------ +struct vs_out +{ + float4 pos [[position]]; + float2 uv; +}; + +vertex vs_out mtl_vertex_shader(ushort vid [[vertex_id]]) +{ + vs_out out; + out.uv = float2((vid << 1) & 2, vid & 2); + out.pos = float4(out.uv * float2(2, -2) + float2(-1, 1), 0, 1); + return(out); +} + +fragment float4 mtl_fragment_shader(vs_out i [[stage_in]], texture2d tex [[texture(0)]]) +{ + constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); + return(tex.sample(smp, i.uv)); +}