[mtl canvas] very first draft of new canvas backend based on edge-counting instead of triangle overlap
This commit is contained in:
parent
c20758f6a2
commit
f6a992e5f4
6
build.sh
6
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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
@ -0,0 +1,112 @@
|
|||
/************************************************************//**
|
||||
*
|
||||
* @file: main.cpp
|
||||
* @author: Martin Fouilleul
|
||||
* @date: 30/07/2022
|
||||
* @revision:
|
||||
*
|
||||
*****************************************************************/
|
||||
#include<stdlib.h>
|
||||
#include<string.h>
|
||||
#include<errno.h>
|
||||
|
||||
#define _USE_MATH_DEFINES //NOTE: necessary for MSVC
|
||||
#include<math.h>
|
||||
|
||||
#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);
|
||||
}
|
||||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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<simd/simd.h>
|
||||
|
||||
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_
|
||||
|
|
@ -0,0 +1,285 @@
|
|||
/************************************************************//**
|
||||
*
|
||||
* @file: mtl_canvas.m
|
||||
* @author: Martin Fouilleul
|
||||
* @date: 12/07/2020
|
||||
* @revision: 24/01/2023
|
||||
*
|
||||
*****************************************************************/
|
||||
#import<Metal/Metal.h>
|
||||
#import<QuartzCore/CAMetalLayer.h>
|
||||
#include<simd/simd.h>
|
||||
|
||||
#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<MTLComputePipelineState> rasterPipeline;
|
||||
id<MTLRenderPipelineState> blitPipeline;
|
||||
|
||||
id<MTLTexture> outTexture;
|
||||
|
||||
id<MTLBuffer> pathBuffer;
|
||||
id<MTLBuffer> 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<MTLComputeCommandEncoder> 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<MTLRenderCommandEncoder> 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<MTLLibrary> 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<MTLFunction> rasterFunction = [library newFunctionWithName:@"mtl_raster"];
|
||||
id<MTLFunction> vertexFunction = [library newFunctionWithName:@"mtl_vertex_shader"];
|
||||
id<MTLFunction> 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
|
||||
|
|
@ -0,0 +1,101 @@
|
|||
|
||||
#include<metal_stdlib>
|
||||
#include<simd/simd.h>
|
||||
#include<metal_simdgroup>
|
||||
|
||||
#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<float, access::write> 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<float> tex [[texture(0)]])
|
||||
{
|
||||
constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear);
|
||||
return(tex.sample(smp, i.uv));
|
||||
}
|
||||
Loading…
Reference in New Issue