hmn
/
orca
5
Fork 2
Code Issues 21 Pull Requests 1 Packages Projects Releases Wiki Activity

applied canvas changes to metal backend

This commit is contained in:
Martin Fouilleul 2023-02-08 10:42:21 +01:00
parent 2419ab7889
commit afaa95fe07
8 changed files with 200 additions and 101 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
examples/canvas/build.sh
View File

@ -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

170
examples/canvas/main.c
View File

@ -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();

1
examples/win_canvas/main.c
View File

@ -80,7 +80,6 @@ int main()
f32 x = 400, y = 300;
f32 dx = 5, dy = 5;
f64 frameTime = 0;
while(!mp_should_quit())

6
src/graphics.c
View File

@ -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;
}

70
src/metal_canvas.m
View File

@ -25,6 +25,8 @@ typedef struct mg_metal_canvas_backend
mg_canvas_backend interface;
mg_surface surface;
mg_color clearColor;
// permanent metal resources
id<MTLComputePipelineState> tilingPipeline;
id<MTLComputePipelineState> sortingPipeline;
@ -37,6 +39,7 @@ typedef struct mg_metal_canvas_backend
// textures and buffers
id<MTLTexture> outTexture;
id<MTLTexture> atlasTexture;
id<MTLBuffer> shapeBuffer;
id<MTLBuffer> vertexBuffer;
id<MTLBuffer> indexBuffer;
id<MTLBuffer> 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<MTLComputeCommandEncoder> 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];

12
src/metal_shader.h
View File

@ -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
{

38
src/metal_shader.metal
View File

@ -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<float, access::write> outTexture [[texture(0)]],
texture2d<float> 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<float, access::write> 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++)
{

2
src/metal_surface.h
View File

@ -9,6 +9,8 @@
#ifndef __METAL_SURFACE_H_
#define __METAL_SURFACE_H_
#include"graphics.h"
#ifdef __OBJC__
#import<Metal/Metal.h>
#endif