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

[mtl canvas] store uv transform and color in triangle data (don't pass vertex or shape buffer to draw kernel)

This commit is contained in:
Martin Fouilleul 2023-03-14 12:20:39 +01:00
parent c9a5b3d52d
commit 221fcbeb6b
3 changed files with 25 additions and 59 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

53
src/mtl_canvas.m
View File

@ -33,7 +33,6 @@ typedef struct mg_mtl_canvas_backend
// permanent metal resources // permanent metal resources
id<MTLComputePipelineState> tilingPipeline; id<MTLComputePipelineState> tilingPipeline;
id<MTLComputePipelineState> sortingPipeline; id<MTLComputePipelineState> sortingPipeline;
id<MTLComputePipelineState> boxingPipeline;
id<MTLComputePipelineState> computePipeline; id<MTLComputePipelineState> computePipeline;
id<MTLRenderPipelineState> renderPipeline; id<MTLRenderPipelineState> renderPipeline;
@ -47,7 +46,6 @@ typedef struct mg_mtl_canvas_backend
id<MTLBuffer> tileCounters; id<MTLBuffer> tileCounters;
id<MTLBuffer> tileArrayBuffer; id<MTLBuffer> tileArrayBuffer;
id<MTLBuffer> triangleArray; id<MTLBuffer> triangleArray;
id<MTLBuffer> boxArray;
} mg_mtl_canvas_backend; } mg_mtl_canvas_backend;
@ -163,29 +161,6 @@ void mg_mtl_canvas_draw_batch(mg_canvas_backend* interface, mg_image_data* image
[blitEncoder fillBuffer: backend->tileCounters range: NSMakeRange(0, RENDERER_MAX_TILES*sizeof(uint)) value: 0]; [blitEncoder fillBuffer: backend->tileCounters range: NSMakeRange(0, RENDERER_MAX_TILES*sizeof(uint)) value: 0];
[blitEncoder endEncoding]; [blitEncoder endEncoding];
/*
//-----------------------------------------------------------
//NOTE(martin): encode the boxing pass
//-----------------------------------------------------------
id<MTLComputeCommandEncoder> boxEncoder = [surface->commandBuffer computeCommandEncoder];
boxEncoder.label = @"boxing pass";
[boxEncoder setComputePipelineState: backend->boxingPipeline];
[boxEncoder setBuffer: backend->vertexBuffer offset:backend->vertexBufferOffset atIndex: 0];
[boxEncoder setBuffer: backend->indexBuffer offset:backend->indexBufferOffset atIndex: 1];
[boxEncoder setBuffer: backend->shapeBuffer offset:backend->shapeBufferOffset 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);
[boxEncoder dispatchThreads: boxGridSize threadsPerThreadgroup: boxGroupSize];
[boxEncoder endEncoding];
*/
//----------------------------------------------------------- //-----------------------------------------------------------
//NOTE(martin): encode the tiling pass //NOTE(martin): encode the tiling pass
//----------------------------------------------------------- //-----------------------------------------------------------
@ -238,10 +213,9 @@ void mg_mtl_canvas_draw_batch(mg_canvas_backend* interface, mg_image_data* image
id<MTLComputeCommandEncoder> drawEncoder = [surface->commandBuffer computeCommandEncoder]; id<MTLComputeCommandEncoder> drawEncoder = [surface->commandBuffer computeCommandEncoder];
drawEncoder.label = @"drawing pass"; drawEncoder.label = @"drawing pass";
[drawEncoder setComputePipelineState:backend->computePipeline]; [drawEncoder setComputePipelineState:backend->computePipeline];
[drawEncoder setBuffer: backend->shapeBuffer offset:backend->shapeBufferOffset atIndex: 0]; [drawEncoder setBuffer: backend->tileCounters offset:0 atIndex: 0];
[drawEncoder setBuffer: backend->tileCounters offset:0 atIndex: 1]; [drawEncoder setBuffer: backend->tileArrayBuffer offset:0 atIndex: 1];
[drawEncoder setBuffer: backend->tileArrayBuffer offset:0 atIndex: 2]; [drawEncoder setBuffer: backend->triangleArray offset:0 atIndex: 2];
[drawEncoder setBuffer: backend->triangleArray offset:0 atIndex: 3];
[drawEncoder setTexture: backend->outTexture atIndex: 0]; [drawEncoder setTexture: backend->outTexture atIndex: 0];
int useTexture = 0; int useTexture = 0;
@ -252,9 +226,9 @@ void mg_mtl_canvas_draw_batch(mg_canvas_backend* interface, mg_image_data* image
useTexture = 1; useTexture = 1;
} }
[drawEncoder setBytes: &clearColorVec4 length: sizeof(vector_float4) atIndex: 4]; [drawEncoder setBytes: &clearColorVec4 length: sizeof(vector_float4) atIndex: 3];
[drawEncoder setBytes: &useTexture length:sizeof(int) atIndex: 5]; [drawEncoder setBytes: &useTexture length:sizeof(int) atIndex: 4];
[drawEncoder setBytes: &scale length: sizeof(float) atIndex: 6]; [drawEncoder setBytes: &scale length: sizeof(float) atIndex: 5];
//TODO: check that we don't exceed maxTotalThreadsPerThreadgroup //TODO: check that we don't exceed maxTotalThreadsPerThreadgroup
DEBUG_ASSERT(RENDERER_TILE_SIZE*RENDERER_TILE_SIZE <= backend->computePipeline.maxTotalThreadsPerThreadgroup); DEBUG_ASSERT(RENDERER_TILE_SIZE*RENDERER_TILE_SIZE <= backend->computePipeline.maxTotalThreadsPerThreadgroup);
@ -341,7 +315,6 @@ void mg_mtl_canvas_destroy(mg_canvas_backend* interface)
[backend->indexBuffer release]; [backend->indexBuffer release];
[backend->tileArrayBuffer release]; [backend->tileArrayBuffer release];
[backend->triangleArray release]; [backend->triangleArray release];
[backend->boxArray release];
[backend->computePipeline release]; [backend->computePipeline release];
} }
} }
@ -473,9 +446,6 @@ mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface)
backend->triangleArray = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_triangle_data) backend->triangleArray = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_triangle_data)
options: MTLResourceStorageModePrivate]; options: MTLResourceStorageModePrivate];
backend->boxArray = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(vector_float4)
options: MTLResourceStorageModePrivate];
//TODO(martin): retain ? //TODO(martin): retain ?
//----------------------------------------------------------- //-----------------------------------------------------------
//NOTE(martin): create and initialize tile counters //NOTE(martin): create and initialize tile counters
@ -499,7 +469,6 @@ mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface)
} }
id<MTLFunction> tilingFunction = [library newFunctionWithName:@"TileKernel"]; id<MTLFunction> tilingFunction = [library newFunctionWithName:@"TileKernel"];
id<MTLFunction> sortingFunction = [library newFunctionWithName:@"SortKernel"]; id<MTLFunction> sortingFunction = [library newFunctionWithName:@"SortKernel"];
id<MTLFunction> boxingFunction = [library newFunctionWithName:@"BoundingBoxKernel"];
id<MTLFunction> computeFunction = [library newFunctionWithName:@"RenderKernel"]; id<MTLFunction> computeFunction = [library newFunctionWithName:@"RenderKernel"];
id<MTLFunction> vertexFunction = [library newFunctionWithName:@"VertexShader"]; id<MTLFunction> vertexFunction = [library newFunctionWithName:@"VertexShader"];
id<MTLFunction> fragmentFunction = [library newFunctionWithName:@"FragmentShader"]; id<MTLFunction> fragmentFunction = [library newFunctionWithName:@"FragmentShader"];
@ -530,16 +499,6 @@ mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface)
reflection: nil reflection: nil
error: &error]; error: &error];
/*
MTLComputePipelineDescriptor* boxingPipelineDesc = [[MTLComputePipelineDescriptor alloc] init];
boxingPipelineDesc.computeFunction = boxingFunction;
// boxingPipelineDesc.threadGroupSizeIsMultipleOfThreadExecutionWidth = true;
backend->boxingPipeline = [metalSurface->device newComputePipelineStateWithDescriptor: boxingPipelineDesc
options: MTLPipelineOptionNone
reflection: nil
error: &error];
*/
//----------------------------------------------------------- //-----------------------------------------------------------
//NOTE(martin): setup our render pipeline state //NOTE(martin): setup our render pipeline state
//----------------------------------------------------------- //-----------------------------------------------------------

4
src/mtl_shader.h
View File

@ -35,9 +35,13 @@ typedef struct mg_shape
typedef struct mg_triangle_data typedef struct mg_triangle_data
{ {
matrix_float3x3 uvTransform;
vector_float4 color;
vector_float4 cubic0; vector_float4 cubic0;
vector_float4 cubic1; vector_float4 cubic1;
vector_float4 cubic2; vector_float4 cubic2;
vector_int4 box; vector_int4 box;
vector_int2 p0; vector_int2 p0;

27
src/mtl_shader.metal
View File

@ -93,6 +93,13 @@ kernel void TileKernel(constant mg_vertex* vertexBuffer [[buffer(0)]],
triangleArray[gid].box = int4(fbox * subPixelFactor); triangleArray[gid].box = int4(fbox * subPixelFactor);
triangleArray[gid].shapeIndex = shapeIndex; triangleArray[gid].shapeIndex = shapeIndex;
triangleArray[gid].color = shapeBuffer[shapeIndex].color;
constant float* uvTransform2x3 = shapeBuffer[shapeIndex].uvTransform;
triangleArray[gid].uvTransform = (matrix_float3x3){{uvTransform2x3[0], uvTransform2x3[3], 0},
{uvTransform2x3[1], uvTransform2x3[4], 0},
{uvTransform2x3[2], uvTransform2x3[5], 1}};
triangleArray[gid].cubic0 = vertexBuffer[i0].cubic; triangleArray[gid].cubic0 = vertexBuffer[i0].cubic;
triangleArray[gid].cubic1 = vertexBuffer[i1].cubic; triangleArray[gid].cubic1 = vertexBuffer[i1].cubic;
triangleArray[gid].cubic2 = vertexBuffer[i2].cubic; triangleArray[gid].cubic2 = vertexBuffer[i2].cubic;
@ -182,14 +189,13 @@ int is_clockwise(int2 p0, int2 p1, int2 p2)
} }
kernel void RenderKernel(const device mg_shape* shapeBuffer [[buffer(0)]], kernel void RenderKernel(const device uint* tileCounters [[buffer(0)]],
device uint* tileCounters [[buffer(1)]], const device uint* tileArrayBuffer [[buffer(1)]],
const device uint* tileArrayBuffer [[buffer(2)]], const device mg_triangle_data* triangleArray [[buffer(2)]],
const device mg_triangle_data* triangleArray [[buffer(3)]],
constant float4* clearColor [[buffer(4)]], constant float4* clearColor [[buffer(3)]],
constant int* useTexture [[buffer(5)]], constant int* useTexture [[buffer(4)]],
constant float* scaling [[buffer(6)]], constant float* scaling [[buffer(5)]],
texture2d<float, access::write> outTexture [[texture(0)]], texture2d<float, access::write> outTexture [[texture(0)]],
texture2d<float> texAtlas [[texture(1)]], texture2d<float> texAtlas [[texture(1)]],
@ -277,15 +283,12 @@ kernel void RenderKernel(const device mg_shape* shapeBuffer [[buffer(0)]],
float4 cubic2 = triangle->cubic2; float4 cubic2 = triangle->cubic2;
int shapeIndex = triangle->shapeIndex; int shapeIndex = triangle->shapeIndex;
float4 color = shapeBuffer[shapeIndex].color; float4 color = triangle->color;
color.rgb *= color.a; color.rgb *= color.a;
int4 clip = triangle->box; int4 clip = triangle->box;
const device float* uvTransform2x3 = shapeBuffer[shapeIndex].uvTransform; matrix_float3x3 uvTransform = triangle->uvTransform;
matrix_float3x3 uvTransform = {{uvTransform2x3[0], uvTransform2x3[3], 0},
{uvTransform2x3[1], uvTransform2x3[4], 0},
{uvTransform2x3[2], uvTransform2x3[5], 1}};
for(int sampleIndex = 0; sampleIndex < sampleCount; sampleIndex++) for(int sampleIndex = 0; sampleIndex < sampleCount; sampleIndex++)
{ {