diff --git a/src/mtl_canvas.m b/src/mtl_canvas.m index 8a5d820..9552bf9 100644 --- a/src/mtl_canvas.m +++ b/src/mtl_canvas.m @@ -33,7 +33,6 @@ typedef struct mg_mtl_canvas_backend // permanent metal resources id tilingPipeline; id sortingPipeline; - id boxingPipeline; id computePipeline; id renderPipeline; @@ -47,7 +46,6 @@ typedef struct mg_mtl_canvas_backend id tileCounters; id tileArrayBuffer; id triangleArray; - id boxArray; } 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 endEncoding]; - /* - //----------------------------------------------------------- - //NOTE(martin): encode the boxing pass - //----------------------------------------------------------- - id 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 //----------------------------------------------------------- @@ -238,10 +213,9 @@ void mg_mtl_canvas_draw_batch(mg_canvas_backend* interface, mg_image_data* image id drawEncoder = [surface->commandBuffer computeCommandEncoder]; drawEncoder.label = @"drawing pass"; [drawEncoder setComputePipelineState:backend->computePipeline]; - [drawEncoder setBuffer: backend->shapeBuffer offset:backend->shapeBufferOffset atIndex: 0]; - [drawEncoder setBuffer: backend->tileCounters offset:0 atIndex: 1]; - [drawEncoder setBuffer: backend->tileArrayBuffer offset:0 atIndex: 2]; - [drawEncoder setBuffer: backend->triangleArray offset:0 atIndex: 3]; + [drawEncoder setBuffer: backend->tileCounters offset:0 atIndex: 0]; + [drawEncoder setBuffer: backend->tileArrayBuffer offset:0 atIndex: 1]; + [drawEncoder setBuffer: backend->triangleArray offset:0 atIndex: 2]; [drawEncoder setTexture: backend->outTexture atIndex: 0]; int useTexture = 0; @@ -252,9 +226,9 @@ void mg_mtl_canvas_draw_batch(mg_canvas_backend* interface, mg_image_data* image useTexture = 1; } - [drawEncoder setBytes: &clearColorVec4 length: sizeof(vector_float4) atIndex: 4]; - [drawEncoder setBytes: &useTexture length:sizeof(int) atIndex: 5]; - [drawEncoder setBytes: &scale length: sizeof(float) atIndex: 6]; + [drawEncoder setBytes: &clearColorVec4 length: sizeof(vector_float4) atIndex: 3]; + [drawEncoder setBytes: &useTexture length:sizeof(int) atIndex: 4]; + [drawEncoder setBytes: &scale length: sizeof(float) atIndex: 5]; //TODO: check that we don't exceed 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->tileArrayBuffer release]; [backend->triangleArray release]; - [backend->boxArray 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) options: MTLResourceStorageModePrivate]; - backend->boxArray = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(vector_float4) - options: MTLResourceStorageModePrivate]; - //TODO(martin): retain ? //----------------------------------------------------------- //NOTE(martin): create and initialize tile counters @@ -499,7 +469,6 @@ mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface) } id tilingFunction = [library newFunctionWithName:@"TileKernel"]; id sortingFunction = [library newFunctionWithName:@"SortKernel"]; - id boxingFunction = [library newFunctionWithName:@"BoundingBoxKernel"]; id computeFunction = [library newFunctionWithName:@"RenderKernel"]; id vertexFunction = [library newFunctionWithName:@"VertexShader"]; id fragmentFunction = [library newFunctionWithName:@"FragmentShader"]; @@ -530,16 +499,6 @@ mg_canvas_backend* mg_mtl_canvas_create(mg_surface surface) reflection: nil 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 //----------------------------------------------------------- diff --git a/src/mtl_shader.h b/src/mtl_shader.h index b473df2..79c1b33 100644 --- a/src/mtl_shader.h +++ b/src/mtl_shader.h @@ -35,9 +35,13 @@ typedef struct mg_shape typedef struct mg_triangle_data { + matrix_float3x3 uvTransform; + vector_float4 color; + vector_float4 cubic0; vector_float4 cubic1; vector_float4 cubic2; + vector_int4 box; vector_int2 p0; diff --git a/src/mtl_shader.metal b/src/mtl_shader.metal index 23a9f8a..48ebcc6 100644 --- a/src/mtl_shader.metal +++ b/src/mtl_shader.metal @@ -93,6 +93,13 @@ kernel void TileKernel(constant mg_vertex* vertexBuffer [[buffer(0)]], triangleArray[gid].box = int4(fbox * subPixelFactor); 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].cubic1 = vertexBuffer[i1].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)]], - device uint* tileCounters [[buffer(1)]], - const device uint* tileArrayBuffer [[buffer(2)]], - const device mg_triangle_data* triangleArray [[buffer(3)]], +kernel void RenderKernel(const device uint* tileCounters [[buffer(0)]], + const device uint* tileArrayBuffer [[buffer(1)]], + const device mg_triangle_data* triangleArray [[buffer(2)]], - constant float4* clearColor [[buffer(4)]], - constant int* useTexture [[buffer(5)]], - constant float* scaling [[buffer(6)]], + constant float4* clearColor [[buffer(3)]], + constant int* useTexture [[buffer(4)]], + constant float* scaling [[buffer(5)]], texture2d outTexture [[texture(0)]], texture2d texAtlas [[texture(1)]], @@ -277,15 +283,12 @@ kernel void RenderKernel(const device mg_shape* shapeBuffer [[buffer(0)]], float4 cubic2 = triangle->cubic2; int shapeIndex = triangle->shapeIndex; - float4 color = shapeBuffer[shapeIndex].color; + float4 color = triangle->color; color.rgb *= color.a; int4 clip = triangle->box; - const device float* uvTransform2x3 = shapeBuffer[shapeIndex].uvTransform; - matrix_float3x3 uvTransform = {{uvTransform2x3[0], uvTransform2x3[3], 0}, - {uvTransform2x3[1], uvTransform2x3[4], 0}, - {uvTransform2x3[2], uvTransform2x3[5], 1}}; + matrix_float3x3 uvTransform = triangle->uvTransform; for(int sampleIndex = 0; sampleIndex < sampleCount; sampleIndex++) {