From afaa95fe077c8e38e76e9a1dc2d73df6e3699c7a Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Wed, 8 Feb 2023 10:42:21 +0100 Subject: [PATCH] applied canvas changes to metal backend --- examples/canvas/build.sh | 2 +- examples/canvas/main.c | 170 ++++++++++++++++++++++++------------- examples/win_canvas/main.c | 1 - src/graphics.c | 6 +- src/metal_canvas.m | 70 ++++++++++----- src/metal_shader.h | 12 ++- src/metal_shader.metal | 38 +++++---- src/metal_surface.h | 2 + 8 files changed, 200 insertions(+), 101 deletions(-) diff --git a/examples/canvas/build.sh b/examples/canvas/build.sh index 65a3ee3..684662b 100755 --- a/examples/canvas/build.sh +++ b/examples/canvas/build.sh @@ -6,6 +6,6 @@ SRCDIR=../../src INCLUDES="-I$SRCDIR -I$SRCDIR/util -I$SRCDIR/platform -I$SRCDIR/app" LIBS="-L$BINDIR -lmilepost -framework Carbon -framework Cocoa -framework Metal -framework QuartzCore" -FLAGS="-mmacos-version-min=10.15.4 -DDEBUG -DLOG_COMPILE_DEBUG -Wl,-dead_strip" +FLAGS="-mmacos-version-min=10.15.4 -DDEBUG -DLOG_COMPILE_DEBUG -Wl,-dead_strip -DMG_IMPLEMENTS_BACKEND_METAL" clang -g $FLAGS $LIBS $INCLUDES -o $BINDIR/example_canvas main.c diff --git a/examples/canvas/main.c b/examples/canvas/main.c index 532c364..0829682 100644 --- a/examples/canvas/main.c +++ b/examples/canvas/main.c @@ -16,11 +16,44 @@ #define LOG_SUBSYSTEM "Main" +mg_font create_font() +{ + //NOTE(martin): create font + str8 fontPath = mp_app_get_resource_path(mem_scratch(), "../resources/OpenSansLatinSubset.ttf"); + char* fontPathCString = str8_to_cstring(mem_scratch(), fontPath); + + FILE* fontFile = fopen(fontPathCString, "r"); + if(!fontFile) + { + LOG_ERROR("Could not load font file '%s'\n", fontPathCString); + return(mg_font_nil()); + } + unsigned char* fontData = 0; + fseek(fontFile, 0, SEEK_END); + u32 fontDataSize = ftell(fontFile); + rewind(fontFile); + fontData = (unsigned char*)malloc(fontDataSize); + fread(fontData, 1, fontDataSize, fontFile); + fclose(fontFile); + + unicode_range ranges[5] = {UNICODE_RANGE_BASIC_LATIN, + UNICODE_RANGE_C1_CONTROLS_AND_LATIN_1_SUPPLEMENT, + UNICODE_RANGE_LATIN_EXTENDED_A, + UNICODE_RANGE_LATIN_EXTENDED_B, + UNICODE_RANGE_SPECIALS}; + + mg_font font = mg_font_create_from_memory(fontDataSize, fontData, 5, ranges); + free(fontData); + + return(font); +} + int main() { LogLevel(LOG_LEVEL_DEBUG); mp_init(); + mp_clock_init(); mp_rect rect = {.x = 100, .y = 100, .w = 800, .h = 600}; mp_window window = mp_window_create(rect, "test", 0); @@ -29,20 +62,27 @@ int main() #if defined(OS_MACOS) mg_surface surface = mg_metal_surface_create_for_window(window); #elif defined(OS_WIN64) - mg_surface surface = mg_gles_surface_create_for_window(window); + mg_surface surface = mg_gl_surface_create_for_window(window); #else #error "unsupported OS" #endif //TODO: create canvas mg_canvas canvas = mg_canvas_create(surface); + mg_font font = create_font(); // start app mp_window_bring_to_front(window); mp_window_focus(window); + f32 x = 400, y = 300; + f32 dx = 5, dy = 5; + 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)) @@ -63,58 +103,29 @@ int main() event.frame.rect.h); } break; - case MP_EVENT_WINDOW_MOVE: - { - printf("moved, rect = {%f, %f, %f, %f}\n", - event.frame.rect.x, - event.frame.rect.y, - event.frame.rect.w, - event.frame.rect.h); - } break; - - case MP_EVENT_MOUSE_MOVE: - { - printf("mouse moved, pos = {%f, %f}, delta = {%f, %f}\n", - event.move.x, - event.move.y, - event.move.deltaX, - event.move.deltaY); - } break; - - case MP_EVENT_MOUSE_WHEEL: - { - printf("mouse wheel, delta = {%f, %f}\n", - event.move.deltaX, - event.move.deltaY); - } break; - - case MP_EVENT_MOUSE_ENTER: - { - printf("mouse enter\n"); - } break; - - case MP_EVENT_MOUSE_LEAVE: - { - printf("mouse leave\n"); - } break; - - case MP_EVENT_MOUSE_BUTTON: - { - printf("mouse button %i: %i\n", - event.key.code, - event.key.action == MP_KEY_PRESS ? 1 : 0); - } break; - case MP_EVENT_KEYBOARD_KEY: { - printf("key %i: %s\n", - event.key.code, - event.key.action == MP_KEY_PRESS ? "press" : (event.key.action == MP_KEY_RELEASE ? "release" : "repeat")); - } break; - - case MP_EVENT_KEYBOARD_CHAR: - { - printf("entered char %s\n", event.character.sequence); + if(event.key.action == MP_KEY_PRESS || event.key.action == MP_KEY_REPEAT) + { + /* + if(event.key.code == MP_KEY_LEFT) + { + dx-=5.1; + } + else if(event.key.code == MP_KEY_RIGHT) + { + dx+=5.1; + } + else if(event.key.code == MP_KEY_UP) + { + dy+=5.1; + } + else if(event.key.code == MP_KEY_DOWN) + { + dy-=5.1; + } + */ + } } break; default: @@ -122,29 +133,72 @@ int main() } } + if(x-200 < 0) + { + dx = 5; + } + if(x+200 > 800) + { + dx = -5; + } + if(y-200 < 0) + { + dy = 5; + } + if(y+200 > 550) + { + dy = -5; + } + x += dx; + y += dy; + mg_surface_prepare(surface); + // background - mg_set_color_rgba(1, 0, 1, 1); + mg_set_color_rgba(0, 1, 1, 1); mg_clear(); // head mg_set_color_rgba(1, 1, 0, 1); - mg_circle_fill(400, 300, 200); + mg_circle_fill(x, y, 200); // smile - mg_set_color_rgba(1, 0, 0, 1); + f32 frown = frameTime > 0.033 ? 100 : 0; + + mg_set_color_rgba(0, 0, 0, 1); mg_set_width(20); - mg_move_to(300, 200); - mg_cubic_to(350, 150, 450, 150, 500, 200); + mg_move_to(x-100, y-100); + mg_cubic_to(x-50, y-150+frown, x+50, y-150+frown, x+100, y-100); mg_stroke(); // eyes - mg_ellipse_fill(330, 350, 30, 50); - mg_ellipse_fill(470, 350, 30, 50); + mg_ellipse_fill(x-70, y+50, 30, 50); + mg_ellipse_fill(x+70, y+50, 30, 50); + + // text + mg_set_color_rgba(0, 0, 1, 1); + mg_set_font(font); + mg_set_font_size(12); + mg_move_to(50, 50); + + str8 text = str8_pushf(mem_scratch(), + "Milepost vector graphics test program (frame time = %fs, fps = %f)...", + frameTime, + 1./frameTime); + mg_text_outlines(text); + 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; + } mp_terminate(); diff --git a/examples/win_canvas/main.c b/examples/win_canvas/main.c index cb3209b..8b1dfee 100644 --- a/examples/win_canvas/main.c +++ b/examples/win_canvas/main.c @@ -80,7 +80,6 @@ int main() f32 x = 400, y = 300; f32 dx = 5, dy = 5; - f64 frameTime = 0; while(!mp_should_quit()) diff --git a/src/graphics.c b/src/graphics.c index 7e77d23..5b3230a 100644 --- a/src/graphics.c +++ b/src/graphics.c @@ -2605,15 +2605,19 @@ mg_canvas mg_canvas_create(mg_surface surface) break; #endif +/* #ifdef MG_IMPLEMENTS_BACKEND_GLES case MG_BACKEND_GLES: backend = mg_gles_canvas_create(surface); break; #endif - +*/ + #ifdef MG_IMPLEMENTS_BACKEND_GL case MG_BACKEND_GL: backend = mg_gl_canvas_create(surface); break; + #endif + default: break; } diff --git a/src/metal_canvas.m b/src/metal_canvas.m index faee470..8cf82da 100644 --- a/src/metal_canvas.m +++ b/src/metal_canvas.m @@ -25,6 +25,8 @@ typedef struct mg_metal_canvas_backend mg_canvas_backend interface; mg_surface surface; + mg_color clearColor; + // permanent metal resources id tilingPipeline; id sortingPipeline; @@ -37,6 +39,7 @@ typedef struct mg_metal_canvas_backend // textures and buffers id outTexture; id atlasTexture; + id shapeBuffer; id vertexBuffer; id indexBuffer; id tileCounters; @@ -57,7 +60,20 @@ mg_metal_surface* mg_metal_canvas_get_surface(mg_metal_canvas_backend* canvas) return(res); } -void mg_metal_canvas_draw_buffers(mg_canvas_backend* interface, u32 vertexCount, u32 indexCount, mg_color clearColor) +void mg_metal_canvas_begin(mg_canvas_backend* interface) +{} + +void mg_metal_canvas_end(mg_canvas_backend* interface) +{} + +void mg_metal_canvas_clear(mg_canvas_backend* interface, mg_color clearColor) +{ + //TODO + mg_metal_canvas_backend* backend = (mg_metal_canvas_backend*)interface; + backend->clearColor = clearColor; +} + +void mg_metal_canvas_draw_batch(mg_canvas_backend* interface, u32 vertexCount, u32 indexCount, u32 shapeCount) { mg_metal_canvas_backend* backend = (mg_metal_canvas_backend*)interface; mg_metal_surface* surface = mg_metal_canvas_get_surface(backend); @@ -92,9 +108,10 @@ void mg_metal_canvas_draw_buffers(mg_canvas_backend* interface, u32 vertexCount, [boxEncoder setComputePipelineState: backend->boxingPipeline]; [boxEncoder setBuffer: backend->vertexBuffer offset:0 atIndex: 0]; [boxEncoder setBuffer: backend->indexBuffer offset:0 atIndex: 1]; - [boxEncoder setBuffer: backend->triangleArray offset:0 atIndex: 2]; - [boxEncoder setBuffer: backend->boxArray offset:0 atIndex: 3]; - [boxEncoder setBytes: &scale length: sizeof(float) atIndex: 4]; + [boxEncoder setBuffer: backend->shapeBuffer offset:0 atIndex: 2]; + [boxEncoder setBuffer: backend->triangleArray offset:0 atIndex: 3]; + [boxEncoder setBuffer: backend->boxArray offset:0 atIndex: 4]; + [boxEncoder setBytes: &scale length: sizeof(float) atIndex: 5]; MTLSize boxGroupSize = MTLSizeMake(backend->boxingPipeline.maxTotalThreadsPerThreadgroup, 1, 1); MTLSize boxGridSize = MTLSizeMake(indexCount/3, 1, 1); @@ -137,20 +154,21 @@ void mg_metal_canvas_draw_buffers(mg_canvas_backend* interface, u32 vertexCount, [sortEncoder endEncoding]; //----------------------------------------------------------- - //NOTE(martin): create compute encoder and encode commands + //NOTE(martin): encode drawing pass //----------------------------------------------------------- - vector_float4 clearColorVec4 = {clearColor.r, clearColor.g, clearColor.b, clearColor.a}; + vector_float4 clearColorVec4 = {backend->clearColor.r, backend->clearColor.g, backend->clearColor.b, backend->clearColor.a}; id encoder = [surface->commandBuffer computeCommandEncoder]; [encoder setComputePipelineState:backend->computePipeline]; [encoder setTexture: backend->outTexture atIndex: 0]; [encoder setTexture: backend->atlasTexture atIndex: 1]; [encoder setBuffer: backend->vertexBuffer offset:0 atIndex: 0]; - [encoder setBuffer: backend->tileCounters offset:0 atIndex: 1]; - [encoder setBuffer: backend->tilesArray offset:0 atIndex: 2]; - [encoder setBuffer: backend->triangleArray offset:0 atIndex: 3]; - [encoder setBuffer: backend->boxArray offset:0 atIndex: 4]; - [encoder setBytes: &clearColorVec4 length: sizeof(vector_float4) atIndex: 5]; + [encoder setBuffer: backend->shapeBuffer offset:0 atIndex: 1]; + [encoder setBuffer: backend->tileCounters offset:0 atIndex: 2]; + [encoder setBuffer: backend->tilesArray offset:0 atIndex: 3]; + [encoder setBuffer: backend->triangleArray offset:0 atIndex: 4]; + [encoder setBuffer: backend->boxArray offset:0 atIndex: 5]; + [encoder setBytes: &clearColorVec4 length: sizeof(vector_float4) atIndex: 6]; //TODO: check that we don't exceed maxTotalThreadsPerThreadgroup DEBUG_ASSERT(RENDERER_TILE_SIZE*RENDERER_TILE_SIZE <= backend->computePipeline.maxTotalThreadsPerThreadgroup); @@ -221,23 +239,27 @@ void mg_metal_canvas_viewport(mg_canvas_backend* interface, mp_rect viewPort) void mg_metal_canvas_update_vertex_layout(mg_metal_canvas_backend* backend) { char* vertexBase = (char*)[backend->vertexBuffer contents]; + char* shapeBase = (char*)[backend->shapeBuffer contents]; + char* indexBase = (char*)[backend->indexBuffer contents]; backend->interface.vertexLayout = (mg_vertex_layout){ .maxVertexCount = MG_METAL_CANVAS_DEFAULT_BUFFER_LENGTH, .maxIndexCount = MG_METAL_CANVAS_DEFAULT_BUFFER_LENGTH, - .posBuffer = vertexBase + offsetof(mg_vertex, pos), - .posStride = sizeof(mg_vertex), .cubicBuffer = vertexBase + offsetof(mg_vertex, cubic), .cubicStride = sizeof(mg_vertex), - .uvBuffer = vertexBase + offsetof(mg_vertex, uv), - .uvStride = sizeof(mg_vertex), - .colorBuffer = vertexBase + offsetof(mg_vertex, color), - .colorStride = sizeof(mg_vertex), + .posBuffer = vertexBase + offsetof(mg_vertex, pos), + .posStride = sizeof(mg_vertex), .zIndexBuffer = vertexBase + offsetof(mg_vertex, zIndex), .zIndexStride = sizeof(mg_vertex), - .clipBuffer = vertexBase + offsetof(mg_vertex, clip), - .clipStride = sizeof(mg_vertex), - .indexBuffer = [backend->indexBuffer contents], + + .colorBuffer = shapeBase + offsetof(mg_shape, color), + .colorStride = sizeof(mg_shape), + .clipBuffer = shapeBase + offsetof(mg_shape, clip), + .clipStride = sizeof(mg_shape), + .uvBuffer = shapeBase + offsetof(mg_shape, uv), + .uvStride = sizeof(mg_shape), + + .indexBuffer = indexBase, .indexStride = sizeof(int)}; } @@ -283,7 +305,10 @@ mg_canvas_backend* mg_metal_canvas_create(mg_surface surface) //NOTE(martin): setup interface functions backend->interface.destroy = mg_metal_canvas_destroy; - backend->interface.drawBuffers = mg_metal_canvas_draw_buffers; + backend->interface.begin = mg_metal_canvas_begin; + backend->interface.end = mg_metal_canvas_end; + backend->interface.clear = mg_metal_canvas_clear; + backend->interface.drawBatch = mg_metal_canvas_draw_batch; backend->interface.atlasUpload = mg_metal_canvas_atlas_upload; mp_rect frame = mg_surface_get_frame(surface); @@ -333,6 +358,9 @@ mg_canvas_backend* mg_metal_canvas_create(mg_surface surface) backend->vertexBuffer = [metalSurface->device newBufferWithLength: MG_METAL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_vertex) options: bufferOptions]; + backend->shapeBuffer = [metalSurface->device newBufferWithLength: MG_METAL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_shape) + options: bufferOptions]; + backend->tilesArray = [metalSurface->device newBufferWithLength: RENDERER_TILE_BUFFER_SIZE*sizeof(int)*RENDERER_MAX_TILES options: MTLResourceStorageModePrivate]; diff --git a/src/metal_shader.h b/src/metal_shader.h index 67c4a27..a30825b 100644 --- a/src/metal_shader.h +++ b/src/metal_shader.h @@ -20,14 +20,18 @@ typedef struct mg_vertex { - vector_float2 pos; // position vector_float4 cubic; // canonical implicit curve space coordinates - vector_float2 uv; // texture coordinates + vector_float2 pos; // position + int zIndex; +} mg_vertex; + +typedef struct mg_shape +{ vector_float4 color; vector_float4 clip; - int zIndex; + vector_float2 uv; // texture coordinates -} mg_vertex; +} mg_shape; typedef struct mg_triangle_data { diff --git a/src/metal_shader.metal b/src/metal_shader.metal index 8a71d4e..a9ac950 100644 --- a/src/metal_shader.metal +++ b/src/metal_shader.metal @@ -35,9 +35,10 @@ bool is_top_left(float2 a, float2 b) kernel void BoundingBoxKernel(constant mg_vertex* vertexBuffer [[buffer(0)]], constant uint* indexBuffer [[buffer(1)]], - device mg_triangle_data* triangleArray [[buffer(2)]], - device float4* boxArray [[buffer(3)]], - constant float* contentsScaling [[buffer(4)]], + constant mg_shape* shapeBuffer [[buffer(2)]], + device mg_triangle_data* triangleArray [[buffer(3)]], + device float4* boxArray [[buffer(4)]], + constant float* contentsScaling [[buffer(5)]], uint gid [[thread_position_in_grid]]) { uint triangleIndex = gid; @@ -56,7 +57,9 @@ kernel void BoundingBoxKernel(constant mg_vertex* vertexBuffer [[buffer(0)]], float2 boxMax = max(max(p0, p1), p2); //NOTE(martin): clip bounding box against clip rect - vector_float4 clip = contentsScaling[0]*vertexBuffer[indexBuffer[vertexIndex]].clip; + int shapeIndex = vertexBuffer[i0].zIndex; + + vector_float4 clip = contentsScaling[0]*shapeBuffer[shapeIndex].clip; float2 clipMin(clip.x, clip.y); float2 clipMax(clip.x + clip.z-1, clip.y + clip.w-1); @@ -83,7 +86,7 @@ kernel void BoundingBoxKernel(constant mg_vertex* vertexBuffer [[buffer(0)]], //NOTE(martin): fill triangle data boxArray[triangleIndex] = float4(boxMin.x, boxMin.y, boxMax.x, boxMax.y); - triangleArray[triangleIndex].zIndex = vertexBuffer[i0].zIndex; + triangleArray[triangleIndex].zIndex = shapeIndex; triangleArray[triangleIndex].i0 = i0; triangleArray[triangleIndex].i1 = i1; triangleArray[triangleIndex].i2 = i2; @@ -174,11 +177,12 @@ float orient2d(float2 a, float2 b, float2 c) kernel void RenderKernel(texture2d outTexture [[texture(0)]], texture2d texAtlas [[texture(1)]], const device mg_vertex* vertexBuffer [[buffer(0)]], - device uint* tileCounters [[buffer(1)]], - const device uint* tilesArray [[buffer(2)]], - const device mg_triangle_data* triangleArray [[buffer(3)]], - const device float4* boxArray [[buffer(4)]], - constant vector_float4* clearColor [[buffer(5)]], + const device mg_shape* shapeBuffer [[buffer(1)]], + device uint* tileCounters [[buffer(2)]], + const device uint* tilesArray [[buffer(3)]], + const device mg_triangle_data* triangleArray [[buffer(4)]], + const device float4* boxArray [[buffer(5)]], + constant vector_float4* clearColor [[buffer(6)]], uint2 gid [[thread_position_in_grid]], uint2 tgid [[threadgroup_position_in_grid]], uint2 threadsPerThreadgroup [[threads_per_threadgroup]], @@ -274,12 +278,16 @@ kernel void RenderKernel(texture2d outTexture [[texture(0) float4 cubic1 = v1->cubic; float4 cubic2 = v2->cubic; - float2 uv0 = v0->uv; - float2 uv1 = v1->uv; - float2 uv2 = v2->uv; - int zIndex = v0->zIndex; - float4 color = v0->color; + float4 color = shapeBuffer[zIndex].color; + + ///////////////////////////////////////////////////////////////////////// + //TODO: dummy uv while we figure out image handling. + float2 uv0 = shapeBuffer[zIndex].uv; + float2 uv1 = uv0; + float2 uv2 = uv0; + ///////////////////////////////////////////////////////////////////////// + for(int i=0; i<6; i++) { diff --git a/src/metal_surface.h b/src/metal_surface.h index 430b5ba..dd06e70 100644 --- a/src/metal_surface.h +++ b/src/metal_surface.h @@ -9,6 +9,8 @@ #ifndef __METAL_SURFACE_H_ #define __METAL_SURFACE_H_ +#include"graphics.h" + #ifdef __OBJC__ #import #endif