From b7f5b84123706485b9c76892cad021a545ed3f67 Mon Sep 17 00:00:00 2001 From: Martin Fouilleul Date: Tue, 11 Apr 2023 11:19:29 +0200 Subject: [PATCH] [mtl canvas] remove old metal canvas stuff --- src/mtl_canvas.m | 667 ------------------------------------------- src/mtl_shader.h | 107 ------- src/mtl_shader.metal | 550 ----------------------------------- 3 files changed, 1324 deletions(-) delete mode 100644 src/mtl_canvas.m delete mode 100644 src/mtl_shader.h delete mode 100644 src/mtl_shader.metal diff --git a/src/mtl_canvas.m b/src/mtl_canvas.m deleted file mode 100644 index 0516f2f..0000000 --- a/src/mtl_canvas.m +++ /dev/null @@ -1,667 +0,0 @@ -/************************************************************//** -* -* @file: mtl_canvas.m -* @author: Martin Fouilleul -* @date: 12/07/2020 -* @revision: 24/01/2023 -* -*****************************************************************/ -#import -#import -#include - -#include"graphics_internal.h" -#include"macro_helpers.h" -#include"osx_app.h" - -#include"mtl_shader.h" - -#define LOG_SUBSYSTEM "Graphics" - -static const int MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH = 4<<20; - -static const int MG_MTL_MAX_BUFFER_AVAILABLE = 3; - -typedef struct mg_mtl_canvas_backend -{ - mg_canvas_backend interface; - mg_surface surface; - - u32 vertexBufferOffset; - u32 indexBufferOffset; - u32 shapeBufferOffset; - - mg_color clearColor; - - // permanent metal resources - id shapePipeline; - id trianglePipeline; - id gatherPipeline; - id computePipeline; - id renderPipeline; - - mp_rect viewPort; - - // triple buffering - u32 bufferIndex; - dispatch_semaphore_t bufferSemaphore; - - // textures and buffers - id backbuffer; - id outTexture; - - id shapeBuffer[MG_MTL_MAX_BUFFER_AVAILABLE]; - id vertexBuffer[MG_MTL_MAX_BUFFER_AVAILABLE]; - id indexBuffer[MG_MTL_MAX_BUFFER_AVAILABLE]; - id shapeQueueBuffer; - id triangleArray; - id tilesBuffer; - id tilesOffset; - id eltBuffer; - id eltOffset; - - id tileArrayBuffer; - id tileCounters; - -} mg_mtl_canvas_backend; - -typedef struct mg_mtl_image_data -{ - mg_image_data interface; - id texture; -} mg_mtl_image_data; - - -mg_mtl_surface* mg_mtl_canvas_get_surface(mg_mtl_canvas_backend* canvas) -{ - mg_mtl_surface* res = 0; - mg_surface_data* data = mg_surface_data_from_handle(canvas->surface); - if(data && data->backend == MG_BACKEND_METAL) - { - res = (mg_mtl_surface*)data; - } - return(res); -} - -void mg_mtl_canvas_update_vertex_layout(mg_mtl_canvas_backend* backend) -{ - char* vertexBase = (char*)[backend->vertexBuffer[backend->bufferIndex] contents] + backend->vertexBufferOffset; - char* shapeBase = (char*)[backend->shapeBuffer[backend->bufferIndex] contents] + backend->shapeBufferOffset; - char* indexBase = (char*)[backend->indexBuffer[backend->bufferIndex] contents] + backend->indexBufferOffset; - - //TODO: add maxShapeCount - - backend->interface.vertexLayout = (mg_vertex_layout){ - .maxVertexCount = MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH - backend->vertexBufferOffset/sizeof(mg_vertex), - .maxIndexCount = MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH - backend->indexBufferOffset/sizeof(int), - - .cubicBuffer = vertexBase + offsetof(mg_vertex, cubic), - .cubicStride = sizeof(mg_vertex), - .posBuffer = vertexBase + offsetof(mg_vertex, pos), - .posStride = sizeof(mg_vertex), - .shapeIndexBuffer = vertexBase + offsetof(mg_vertex, shapeIndex), - .shapeIndexStride = sizeof(mg_vertex), - - .colorBuffer = shapeBase + offsetof(mg_shape, color), - .colorStride = sizeof(mg_shape), - - .texturedBuffer = shapeBase + offsetof(mg_shape, textured), - .texturedStride = sizeof(mg_shape), - - .clipBuffer = shapeBase + offsetof(mg_shape, clip), - .clipStride = sizeof(mg_shape), - .uvTransformBuffer = shapeBase + offsetof(mg_shape, uvTransform), - .uvTransformStride = sizeof(mg_shape), - - .indexBuffer = indexBase, - .indexStride = sizeof(int)}; -} - -void mg_mtl_canvas_begin(mg_canvas_backend* interface, mg_color clearColor) -{ - mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; - mg_mtl_surface* surface = mg_mtl_canvas_get_surface(backend); - if(!surface) - { - return; - } - backend->clearColor = clearColor; - - backend->vertexBufferOffset = 0; - backend->indexBufferOffset = 0; - backend->shapeBufferOffset = 0; - - mg_mtl_canvas_update_vertex_layout(backend); - - mg_mtl_surface_acquire_command_buffer(surface); - - @autoreleasepool - { - MTLClearColor mtlClearColor = MTLClearColorMake(clearColor.r, - clearColor.g, - clearColor.b, - clearColor.a); - - MTLRenderPassDescriptor* renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor]; - renderPassDescriptor.colorAttachments[0].texture = backend->backbuffer; - renderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionClear; - renderPassDescriptor.colorAttachments[0].clearColor = mtlClearColor; - renderPassDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore; - - id renderEncoder = [surface->commandBuffer renderCommandEncoderWithDescriptor:renderPassDescriptor]; - renderEncoder.label = @"clear pass"; - [renderEncoder endEncoding]; - } -} - -void mg_mtl_canvas_end(mg_canvas_backend* interface) -{ - mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; - mg_mtl_surface* surface = mg_mtl_canvas_get_surface(backend); - - if(surface && surface->commandBuffer) - { - @autoreleasepool - { - mg_mtl_surface_acquire_drawable(surface); - if(surface->drawable != nil) - { - f32 scale = surface->mtlLayer.contentsScale; - MTLViewport viewport = {backend->viewPort.x * scale, - backend->viewPort.y * scale, - backend->viewPort.w * scale, - backend->viewPort.h * scale, - 0, - 1}; - - MTLRenderPassDescriptor* renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor]; - renderPassDescriptor.colorAttachments[0].texture = surface->drawable.texture; - renderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionLoad; - renderPassDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore; - - id renderEncoder = [surface->commandBuffer renderCommandEncoderWithDescriptor:renderPassDescriptor]; - renderEncoder.label = @"blit pass"; - [renderEncoder setViewport: viewport]; - [renderEncoder setRenderPipelineState: backend->renderPipeline]; - [renderEncoder setFragmentTexture: backend->backbuffer atIndex: 0]; - [renderEncoder drawPrimitives: MTLPrimitiveTypeTriangle - vertexStart: 0 - vertexCount: 3 ]; - [renderEncoder endEncoding]; - } - [surface->commandBuffer addCompletedHandler:^(id commandBuffer) - { - dispatch_semaphore_signal(backend->bufferSemaphore); - } - ]; - - dispatch_semaphore_wait(backend->bufferSemaphore, DISPATCH_TIME_FOREVER); - backend->bufferIndex = (backend->bufferIndex + 1) % MG_MTL_MAX_BUFFER_AVAILABLE; - } - } -} - -void mg_mtl_canvas_draw_batch(mg_canvas_backend* interface, mg_image_data* image, u32 shapeCount, u32 vertexCount, u32 indexCount) -{ - mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; - mg_mtl_surface* surface = mg_mtl_canvas_get_surface(backend); - - if(!surface || (backend->backbuffer == nil)) - { - return; - } - - //TODO: guard against overflowing buffers... - - @autoreleasepool - { - f32 scale = surface->mtlLayer.contentsScale; - vector_uint2 viewportSize = {backend->viewPort.w * scale, backend->viewPort.h * scale}; - u32 nTilesX = (viewportSize.x + RENDERER_TILE_SIZE - 1)/RENDERER_TILE_SIZE; - u32 nTilesY = (viewportSize.y + RENDERER_TILE_SIZE - 1)/RENDERER_TILE_SIZE; - int triangleCount = indexCount/3; - - printf("triangle count: %i, shape count: %i\n", triangleCount, shapeCount); - - //----------------------------------------------------------- - //NOTE(martin): encode the clear arena offset - //----------------------------------------------------------- - id blitEncoder = [surface->commandBuffer blitCommandEncoder]; - blitEncoder.label = @"clear arena"; - [blitEncoder fillBuffer: backend->tilesOffset range: NSMakeRange(0, sizeof(int)) value: 0]; - [blitEncoder fillBuffer: backend->eltOffset range: NSMakeRange(0, sizeof(int)) value: 0]; - [blitEncoder fillBuffer: backend->tileCounters range: NSMakeRange(0, RENDERER_MAX_TILES*sizeof(uint)) value: 0]; - [blitEncoder endEncoding]; - - //----------------------------------------------------------- - //NOTE(martin): encode the shape setup pass - //----------------------------------------------------------- - id shapeEncoder = [surface->commandBuffer computeCommandEncoder]; - shapeEncoder.label = @"shape pass"; - [shapeEncoder setComputePipelineState: backend->shapePipeline]; - [shapeEncoder setBuffer: backend->shapeBuffer[backend->bufferIndex] offset:backend->shapeBufferOffset atIndex: 0]; - [shapeEncoder setBuffer: backend->shapeQueueBuffer offset:0 atIndex: 1]; - [shapeEncoder setBuffer: backend->tilesBuffer offset:0 atIndex: 2]; - [shapeEncoder setBuffer: backend->tilesOffset offset:0 atIndex: 3]; - [shapeEncoder setBytes: &scale length: sizeof(float) atIndex: 4]; - [shapeEncoder setBytes: &viewportSize length: sizeof(vector_uint2) atIndex: 5]; - - MTLSize shapeGroupSize = MTLSizeMake(backend->shapePipeline.maxTotalThreadsPerThreadgroup, 1, 1); - MTLSize shapeGridSize = MTLSizeMake(shapeCount, 1, 1); - - [shapeEncoder dispatchThreads: shapeGridSize threadsPerThreadgroup: shapeGroupSize]; - [shapeEncoder endEncoding]; - - //----------------------------------------------------------- - //NOTE(martin): encode the triangle setup and binning - //----------------------------------------------------------- - id triangleEncoder = [surface->commandBuffer computeCommandEncoder]; - triangleEncoder.label = @"triangle pass"; - [triangleEncoder setComputePipelineState: backend->trianglePipeline]; - [triangleEncoder setBuffer: backend->vertexBuffer[backend->bufferIndex] offset:backend->vertexBufferOffset atIndex: 0]; - [triangleEncoder setBuffer: backend->indexBuffer[backend->bufferIndex] offset:backend->indexBufferOffset atIndex: 1]; - [triangleEncoder setBuffer: backend->shapeBuffer[backend->bufferIndex] offset:backend->shapeBufferOffset atIndex: 2]; - [triangleEncoder setBuffer: backend->triangleArray offset:0 atIndex: 3]; - [triangleEncoder setBuffer: backend->shapeQueueBuffer offset:0 atIndex: 4]; - [triangleEncoder setBuffer: backend->tilesBuffer offset:0 atIndex: 5]; - [triangleEncoder setBuffer: backend->eltBuffer offset:0 atIndex: 6]; - [triangleEncoder setBuffer: backend->eltOffset offset:0 atIndex: 7]; - - [triangleEncoder setBytes: &scale length: sizeof(float) atIndex: 8]; - - MTLSize triangleGroupSize = MTLSizeMake(backend->trianglePipeline.maxTotalThreadsPerThreadgroup, 1, 1); - MTLSize triangleGridSize = MTLSizeMake(triangleCount, 1, 1); - - [triangleEncoder dispatchThreads: triangleGridSize threadsPerThreadgroup: triangleGroupSize]; - [triangleEncoder endEncoding]; - - //----------------------------------------------------------- - //NOTE(martin): encode gathering pass - //----------------------------------------------------------- - id gatherEncoder = [surface->commandBuffer computeCommandEncoder]; - gatherEncoder.label = @"gather pass"; - [gatherEncoder setComputePipelineState: backend->gatherPipeline]; - [gatherEncoder setBuffer: backend->shapeQueueBuffer offset:0 atIndex: 0]; - [gatherEncoder setBuffer: backend->tilesBuffer offset:0 atIndex: 1]; - [gatherEncoder setBuffer: backend->eltBuffer offset:0 atIndex: 2]; - [gatherEncoder setBuffer: backend->tileCounters offset:0 atIndex: 3]; - [gatherEncoder setBuffer: backend->tileArrayBuffer offset:0 atIndex: 4]; - - [gatherEncoder setBytes: &shapeCount length: sizeof(int) atIndex: 5]; - [gatherEncoder setBytes: &viewportSize length: sizeof(vector_uint2) atIndex: 6]; - - MTLSize gatherGroupSize = MTLSizeMake(16, 16, 1); - MTLSize gatherGridSize = MTLSizeMake(nTilesX, nTilesY, 1); - - [gatherEncoder dispatchThreads: gatherGridSize threadsPerThreadgroup: gatherGroupSize]; - [gatherEncoder endEncoding]; - - //----------------------------------------------------------- - //NOTE(martin): encode drawing pass - //----------------------------------------------------------- - id drawEncoder = [surface->commandBuffer computeCommandEncoder]; - drawEncoder.label = @"drawing pass"; - [drawEncoder setComputePipelineState:backend->computePipeline]; - [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; - if(image) - { - mg_mtl_image_data* mtlImage = (mg_mtl_image_data*)image; - [drawEncoder setTexture: mtlImage->texture atIndex: 1]; - useTexture = 1; - } - - [drawEncoder setBytes: &useTexture length:sizeof(int) atIndex: 3]; - [drawEncoder setBytes: &scale length: sizeof(float) atIndex: 4]; - - //TODO: check that we don't exceed maxTotalThreadsPerThreadgroup - DEBUG_ASSERT(RENDERER_TILE_SIZE*RENDERER_TILE_SIZE <= backend->computePipeline.maxTotalThreadsPerThreadgroup); - MTLSize threadGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1); - MTLSize threadGroupSize = MTLSizeMake(RENDERER_TILE_SIZE, RENDERER_TILE_SIZE, 1); - - [drawEncoder dispatchThreads: threadGridSize threadsPerThreadgroup:threadGroupSize]; - [drawEncoder endEncoding]; - - //----------------------------------------------------------- - //NOTE(martin): blit texture to backbuffer - //----------------------------------------------------------- - - MTLViewport viewport = {backend->viewPort.x * scale, - backend->viewPort.y * scale, - backend->viewPort.w * scale, - backend->viewPort.h * scale, - 0, - 1}; - - MTLRenderPassDescriptor* renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor]; - renderPassDescriptor.colorAttachments[0].texture = backend->backbuffer; - renderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionLoad; - renderPassDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore; - - id renderEncoder = [surface->commandBuffer renderCommandEncoderWithDescriptor:renderPassDescriptor]; - renderEncoder.label = @"blit pass"; - [renderEncoder setViewport: viewport]; - [renderEncoder setRenderPipelineState: backend->renderPipeline]; - [renderEncoder setFragmentTexture: backend->outTexture atIndex: 0]; - [renderEncoder drawPrimitives: MTLPrimitiveTypeTriangle - vertexStart: 0 - vertexCount: 3 ]; - [renderEncoder endEncoding]; - } - - backend->vertexBufferOffset += vertexCount * sizeof(mg_vertex); - backend->indexBufferOffset += indexCount * sizeof(int); - backend->shapeBufferOffset += shapeCount * sizeof(mg_shape); - - mg_mtl_canvas_update_vertex_layout(backend); -} - -/* -void mg_mtl_canvas_viewport(mg_canvas_backend* interface, mp_rect viewPort) -{ - mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; - mg_mtl_surface* surface = mg_mtl_canvas_get_surface(backend); - if(!surface) - { - return; - } - - backend->viewPort = viewPort; - - @autoreleasepool - { - f32 scale = surface->mtlLayer.contentsScale; - CGSize drawableSize = (CGSize){.width = viewPort.w * scale, .height = viewPort.h * scale}; - - [backend->outTexture release]; - - MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; - texDesc.textureType = MTLTextureType2D; - texDesc.storageMode = MTLStorageModePrivate; - texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite; - texDesc.pixelFormat = MTLPixelFormatBGRA8Unorm;// MTLPixelFormatBGRA8Unorm_sRGB; - texDesc.width = drawableSize.width; - texDesc.height = drawableSize.height; - - backend->outTexture = [surface->device newTextureWithDescriptor:texDesc]; - } -} -*/ - -void mg_mtl_canvas_destroy(mg_canvas_backend* interface) -{ - mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; - - @autoreleasepool - { - [backend->outTexture release]; - - for(int i=0; i < MG_MTL_MAX_BUFFER_AVAILABLE; i++) - { - [backend->vertexBuffer[i] release]; - [backend->indexBuffer[i] release]; - [backend->shapeBuffer[i] release]; - } - //NOTE: semaphore does not have a destructor? - - [backend->shapeQueueBuffer release]; - [backend->triangleArray release]; - [backend->tilesBuffer release]; - [backend->tilesOffset release]; - [backend->eltBuffer release]; - [backend->eltOffset release]; - - ////////////////////////////////////////// - //TODO release all pipelines - [backend->computePipeline release]; - } -} - -mg_image_data* mg_mtl_canvas_image_create(mg_canvas_backend* interface, vec2 size) -{ - mg_mtl_image_data* image = 0; - mg_mtl_canvas_backend* backend = (mg_mtl_canvas_backend*)interface; - mg_mtl_surface* surface = mg_mtl_canvas_get_surface(backend); - - if(surface) - { - @autoreleasepool{ - - image = malloc_type(mg_mtl_image_data); - if(image) - { - MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; - texDesc.textureType = MTLTextureType2D; - texDesc.storageMode = MTLStorageModeManaged; - texDesc.usage = MTLTextureUsageShaderRead; - texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm; - texDesc.width = size.x; - texDesc.height = size.y; - - image->texture = [surface->device newTextureWithDescriptor:texDesc]; - if(image->texture != nil) - { - [image->texture retain]; - image->interface.size = size; - } - else - { - free(image); - image = 0; - } - } - } - } - return((mg_image_data*)image); -} - -void mg_mtl_canvas_image_destroy(mg_canvas_backend* backendInterface, mg_image_data* imageInterface) -{ - mg_mtl_image_data* image = (mg_mtl_image_data*)imageInterface; - @autoreleasepool - { - [image->texture release]; - free(image); - } -} - -void mg_mtl_canvas_image_upload_region(mg_canvas_backend* backendInterface, mg_image_data* imageInterface, mp_rect region, u8* pixels) -{@autoreleasepool{ - mg_mtl_image_data* image = (mg_mtl_image_data*)imageInterface; - MTLRegion mtlRegion = MTLRegionMake2D(region.x, region.y, region.w, region.h); - [image->texture replaceRegion:mtlRegion - mipmapLevel:0 - withBytes:(void*)pixels - bytesPerRow: 4 * region.w]; -}} - -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.begin = mg_mtl_canvas_begin; - backend->interface.end = mg_mtl_canvas_end; - backend->interface.drawBatch = mg_mtl_canvas_draw_batch; - - backend->interface.imageCreate = mg_mtl_canvas_image_create; - backend->interface.imageDestroy = mg_mtl_canvas_image_destroy; - backend->interface.imageUploadRegion = mg_mtl_canvas_image_upload_region; - - mp_rect frame = mg_surface_get_frame(surface); - backend->viewPort = (mp_rect){0, 0, frame.w, frame.h}; - - @autoreleasepool - { - f32 scale = metalSurface->mtlLayer.contentsScale; - CGSize drawableSize = (CGSize){.width = backend->viewPort.w * scale, .height = backend->viewPort.h * scale}; - - //----------------------------------------------------------- - //NOTE(martin): create our output texture - //----------------------------------------------------------- - MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; - texDesc.textureType = MTLTextureType2D; - texDesc.storageMode = MTLStorageModePrivate; - texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite; - texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm; - texDesc.width = drawableSize.width; - texDesc.height = drawableSize.height; - - backend->outTexture = [metalSurface->device newTextureWithDescriptor:texDesc]; - - texDesc.usage = MTLTextureUsageRenderTarget | MTLTextureUsageShaderRead; - backend->backbuffer = [metalSurface->device newTextureWithDescriptor:texDesc]; - //TODO(martin): retain ? - - //----------------------------------------------------------- - //NOTE(martin): create buffers - //----------------------------------------------------------- - - backend->bufferSemaphore = dispatch_semaphore_create(MG_MTL_MAX_BUFFER_AVAILABLE); - backend->bufferIndex = 0; - - MTLResourceOptions bufferOptions = MTLResourceCPUCacheModeWriteCombined - | MTLResourceStorageModeShared; - - for(int i=0; iindexBuffer[i] = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(int) - options: bufferOptions]; - - backend->vertexBuffer[i] = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_vertex) - options: bufferOptions]; - - backend->shapeBuffer[i] = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_shape) - options: bufferOptions]; - } - - backend->triangleArray = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_triangle_data) - options: MTLResourceStorageModePrivate]; - - backend->shapeQueueBuffer = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_shape_queue) - options: MTLResourceStorageModePrivate]; - - backend->tilesBuffer = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_tile) - options: MTLResourceStorageModePrivate]; - - backend->tilesOffset = [metalSurface->device newBufferWithLength: sizeof(int) - options: MTLResourceStorageModePrivate]; - - backend->eltBuffer = [metalSurface->device newBufferWithLength: MG_MTL_CANVAS_DEFAULT_BUFFER_LENGTH*sizeof(mg_tile_elt) - options: MTLResourceStorageModePrivate]; - - backend->eltOffset = [metalSurface->device newBufferWithLength: sizeof(int) - options: MTLResourceStorageModePrivate]; - - backend->tileArrayBuffer = [metalSurface->device newBufferWithLength: RENDERER_TILE_BUFFER_COUNT*sizeof(mg_tile_cmd)*RENDERER_MAX_TILES - options: MTLResourceStorageModePrivate]; - - backend->tileCounters = [metalSurface->device newBufferWithLength: RENDERER_MAX_TILES*sizeof(uint) - options: MTLResourceStorageModePrivate]; - - //----------------------------------------------------------- - //NOTE(martin): load the library - //----------------------------------------------------------- - - //TODO(martin): filepath magic to find metallib path when not in the working directory - str8 shaderPath = mp_app_get_resource_path(mem_scratch(), "../resources/mtl_shader.metallib"); - NSString* metalFileName = [[NSString alloc] initWithBytes: shaderPath.ptr length:shaderPath.len encoding: NSUTF8StringEncoding]; - NSError* err = 0; - id 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 shapeFunction = [library newFunctionWithName:@"ShapeSetup"]; - id triangleFunction = [library newFunctionWithName:@"TriangleKernel"]; - id gatherFunction = [library newFunctionWithName:@"GatherKernel"]; - id computeFunction = [library newFunctionWithName:@"RenderKernel"]; - id vertexFunction = [library newFunctionWithName:@"VertexShader"]; - id fragmentFunction = [library newFunctionWithName:@"FragmentShader"]; - - //----------------------------------------------------------- - //NOTE(martin): setup our data layout and pipeline state - //----------------------------------------------------------- - NSError* error = NULL; - backend->computePipeline = [metalSurface->device newComputePipelineStateWithFunction: computeFunction - error:&error]; - ASSERT(backend->computePipeline); - - MTLComputePipelineDescriptor* shapePipelineDesc = [[MTLComputePipelineDescriptor alloc] init]; - shapePipelineDesc.computeFunction = shapeFunction; - - backend->shapePipeline = [metalSurface->device newComputePipelineStateWithDescriptor: shapePipelineDesc - options: MTLPipelineOptionNone - reflection: nil - error: &error]; - - MTLComputePipelineDescriptor* trianglePipelineDesc = [[MTLComputePipelineDescriptor alloc] init]; - trianglePipelineDesc.computeFunction = triangleFunction; - - backend->trianglePipeline = [metalSurface->device newComputePipelineStateWithDescriptor: trianglePipelineDesc - options: MTLPipelineOptionNone - reflection: nil - error: &error]; - - MTLComputePipelineDescriptor* gatherPipelineDesc = [[MTLComputePipelineDescriptor alloc] init]; - gatherPipelineDesc.computeFunction = gatherFunction; - - backend->gatherPipeline = [metalSurface->device newComputePipelineStateWithDescriptor: gatherPipelineDesc - options: MTLPipelineOptionNone - reflection: nil - error: &error]; - - //----------------------------------------------------------- - //NOTE(martin): setup our render pipeline state - //----------------------------------------------------------- - // create and initialize the pipeline state descriptor - MTLRenderPipelineDescriptor *pipelineStateDescriptor = [[MTLRenderPipelineDescriptor alloc] init]; - pipelineStateDescriptor.label = @"My simple 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; - - // create render pipeline - backend->renderPipeline = [metalSurface->device newRenderPipelineStateWithDescriptor: pipelineStateDescriptor error:&err]; - if(err != nil) - { - const char* errStr = [[err localizedDescription] UTF8String]; - const char* descStr = [[err localizedFailureReason] UTF8String]; - const char* recovStr = [[err localizedRecoverySuggestion] UTF8String]; - LOG_ERROR("(%li) %s. %s. %s\n", [err code], errStr, descStr, recovStr); - return(0); - } - } - - mg_mtl_canvas_update_vertex_layout(backend); - } - - return((mg_canvas_backend*)backend); -} - - -#undef LOG_SUBSYSTEM diff --git a/src/mtl_shader.h b/src/mtl_shader.h deleted file mode 100644 index 7317800..0000000 --- a/src/mtl_shader.h +++ /dev/null @@ -1,107 +0,0 @@ -/************************************************************//** -* -* @file: mtl_shader.h -* @author: Martin Fouilleul -* @date: 01/08/2022 -* @revision: -* -*****************************************************************/ -#ifndef __MTL_RENDERER_H_ -#define __MTL_RENDERER_H_ - -#include - -#define RENDERER_TILE_SIZE 16 -#define RENDERER_MAX_TILES 65536 -#define RENDERER_TILE_BUFFER_COUNT 8*(1<<10) - -#define RENDERER_DEBUG_TILE_VISITED 0xf00d -#define RENDERER_DEBUG_TILE_BUFFER_OVERFLOW 0xdead - -typedef struct mg_vertex -{ - vector_float4 cubic; // canonical implicit curve space coordinates - vector_float2 pos; // position - int shapeIndex; -} mg_vertex; - -typedef struct mg_shape -{ - vector_float4 color; - vector_float4 clip; - float uvTransform[6]; - bool textured; -} mg_shape; - -typedef struct mg_triangle_data -{ - matrix_float3x3 uvTransform; - vector_float4 color; - - bool full; - vector_float4 cubic0; - vector_float4 cubic1; - vector_float4 cubic2; - - vector_int4 box; - - vector_int2 p0; - vector_int2 p1; - vector_int2 p2; - - int bias0; - int bias1; - int bias2; - - int cw; - - int shapeIndex; - -} mg_triangle_data; - -#ifndef __METAL_VERSION__ -#define device -#else -using namespace metal; -#endif - -#define MG_TILE_CMD_MASK (3<<30) - -typedef enum mg_tile_cmd_kind -{ - mg_cmd_triangle = 0, - mg_cmd_color = 1<<30, - mg_cmd_flip = 2<<30 -} mg_tile_cmd_kind; - -typedef int mg_tile_cmd; - -typedef struct mg_tile_elt -{ - int triangleIndex; - int next; -} mg_tile_elt; - -typedef struct mg_tile -{ - vector_float4 color; - atomic_int firstElt; - atomic_int eltCount; - atomic_int partial; - atomic_int flipCount; - bool textured; - -} mg_tile; - -typedef struct mg_shape_queue -{ - vector_int4 area; - int tiles; -} mg_shape_queue; - -#ifndef __METAL_VERSION__ -#undef device -#endif - - -#endif //__MTL_RENDERER_H_ diff --git a/src/mtl_shader.metal b/src/mtl_shader.metal deleted file mode 100644 index 7df34d3..0000000 --- a/src/mtl_shader.metal +++ /dev/null @@ -1,550 +0,0 @@ - -#include -#include -#include - -#include"mtl_shader.h" - -using namespace metal; - -struct vs_out -{ - float4 pos [[position]]; - float2 uv; -}; - -vertex vs_out VertexShader(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 FragmentShader(vs_out i [[stage_in]], texture2d tex [[texture(0)]]) -{ - constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); - return(tex.sample(smp, i.uv)); -} - -bool is_top_left(float2 a, float2 b) -{ - return( (a.y == b.y && b.x < a.x) - ||(b.y < a.y)); -} - -////////////////////////////////////////////////////////////////////////////// -//TODO: we should do these computations on 64bits, because otherwise -// we might overflow for values > 2048. -// Unfortunately this is costly. -// Another way is to precompute triangle edges (b - a) in full precision -// once to avoid doing it all the time... -////////////////////////////////////////////////////////////////////////////// - -int orient2d(int2 a, int2 b, int2 c) -{ - return((b.x-a.x)*(c.y-a.y) - (b.y-a.y)*(c.x-a.x)); -} - -device uchar* arena_allocate(device uchar* arenaBuffer, - device volatile atomic_uint* arenaOffset, - uint size) -{ - uint index = atomic_fetch_add_explicit(arenaOffset, size, memory_order_relaxed); - return(&arenaBuffer[index]); -} - -//NOTE: shape setup allocates tile queues for each shape - -kernel void ShapeSetup(constant mg_shape* shapeBuffer [[buffer(0)]], - device mg_shape_queue* shapeQueueBuffer [[buffer(1)]], - device mg_tile* tilesBuffer [[buffer(2)]], - device volatile atomic_uint* tilesOffset [[buffer(3)]], - constant float* scaling [[buffer(4)]], - constant int2* viewport [[buffer(5)]], - uint gid [[thread_position_in_grid]]) -{ - int2 tilesMatrixDim = (*viewport - 1) / RENDERER_TILE_SIZE + 1; - float4 box = shapeBuffer[gid].clip * scaling[0]; - - int2 firstTile = int2(box.xy)/RENDERER_TILE_SIZE; - - //WARN: the following can result in a 1x1 tile allocated even for empty boxes. But if we didn't allocate - // any tile queue, the tileQueues pointer for that shape would alias the tileQueues pointer of another - // shape, and we would have to detect that in the tiling and drawing kernels. Instead, just accept some - // waste and keep the other kernels more uniforms for now... - //TODO limit to screen - int2 lastTile = max(firstTile, min(int2(box.zw)/RENDERER_TILE_SIZE, tilesMatrixDim)); - - int nTilesX = lastTile.x - firstTile.x + 1; - int nTilesY = lastTile.y - firstTile.y + 1; - - int tileCount = nTilesX * nTilesY; - - int tilesIndex = atomic_fetch_add_explicit(tilesOffset, tileCount, memory_order_relaxed); - - shapeQueueBuffer[gid].area = int4(firstTile.x, firstTile.y, nTilesX, nTilesY); - shapeQueueBuffer[gid].tiles = tilesIndex; - - device mg_tile* tiles = &tilesBuffer[tilesIndex]; - - for(int i=0; i 0 ? 1 : -1; - - triangleArray[gid].cw = cw; - triangleArray[gid].bias0 = is_top_left(p1, p2) ? -(1-cw)/2 : -(1+cw)/2; - triangleArray[gid].bias1 = is_top_left(p2, p0) ? -(1-cw)/2 : -(1+cw)/2; - triangleArray[gid].bias2 = is_top_left(p0, p1) ? -(1-cw)/2 : -(1+cw)/2; - - bool triangleFull = all( triangleArray[gid].cubic0 == float4(1, 1, 1, 1) - && triangleArray[gid].cubic1 == float4(1, 1, 1, 1) - && triangleArray[gid].cubic2 == float4(1, 1, 1, 1)); - - triangleArray[gid].full = triangleFull; - - int4 coarseBox = int4(fbox)/RENDERER_TILE_SIZE; - - //NOTE: bucket triangle into tiles - device mg_shape_queue* shapeQueue = &shapeQueueBuffer[shapeIndex]; - device mg_tile* tiles = &tilesBuffer[shapeQueue->tiles]; - - int xMin = max(0, coarseBox.x - shapeQueue->area.x); - int yMin = max(0, coarseBox.y - shapeQueue->area.y); - int xMax = min(coarseBox.z - shapeQueue->area.x, shapeQueue->area.z-1); - int yMax = min(coarseBox.w - shapeQueue->area.y, shapeQueue->area.w-1); - - //NOTE(martin): it's important to do the computation with signed int, so that we can have negative xMax/yMax - // otherwise all triangles on the left or below the x/y axis are attributed to tiles on row/column 0. - - int2 edges[3][2] = {{ip0, ip1}, {ip1, ip2}, {ip2, ip0}}; - - for(int y = yMin; y <= yMax; y++) - { - for(int x = xMin ; x <= xMax; x++) - { - int4 tileBox = int4(shapeQueue->area.x + x, - shapeQueue->area.y + y, - shapeQueue->area.x + x + 1, - shapeQueue->area.y + y + 1) * RENDERER_TILE_SIZE*int(subPixelFactor); - - int2 b[4] = {{tileBox.x, tileBox.y}, - {tileBox.z, tileBox.y}, - {tileBox.z, tileBox.w}, - {tileBox.x, tileBox.w}}; - - //TODO: should add the biases here? - - - int sideFromEdge[3][4]; - for(int edgeIndex=0; edgeIndex<3; edgeIndex++) - { - for(int cornerIndex=0; cornerIndex<4; cornerIndex++) - { - sideFromEdge[edgeIndex][cornerIndex] = cw*orient2d(edges[edgeIndex][0], - edges[edgeIndex][1], - b[cornerIndex]); - } - } - - bool allRightFromEdge0 = sideFromEdge[0][0] < 0 - && sideFromEdge[0][1] < 0 - && sideFromEdge[0][2] < 0 - && sideFromEdge[0][3] < 0; - - bool allRightFromEdge1 = sideFromEdge[1][0] < 0 - && sideFromEdge[1][1] < 0 - && sideFromEdge[1][2] < 0 - && sideFromEdge[1][3] < 0; - - bool allRightFromEdge2 = sideFromEdge[2][0] < 0 - && sideFromEdge[2][1] < 0 - && sideFromEdge[2][2] < 0 - && sideFromEdge[2][3] < 0; - - bool allOutside = allRightFromEdge0 || allRightFromEdge1 || allRightFromEdge2; - if(!allOutside) - { - int tileIndex = y*shapeQueue->area.z + x; - device mg_tile* tile = &tiles[tileIndex]; - - int eltIndex = atomic_fetch_add_explicit(eltOffset, 1, memory_order_relaxed); - - device mg_tile_elt* elt = &eltBuffer[eltIndex]; - elt->triangleIndex = gid; - - elt->next = atomic_exchange_explicit(&tile->firstElt, eltIndex, memory_order_relaxed); - - atomic_fetch_add_explicit(&tile->eltCount, 1, memory_order_relaxed); - - bool allLeftFromEdge0 = sideFromEdge[0][0] > 0 - && sideFromEdge[0][1] > 0 - && sideFromEdge[0][2] > 0 - && sideFromEdge[0][3] > 0; - - bool allLeftFromEdge1 = sideFromEdge[1][0] > 0 - && sideFromEdge[1][1] > 0 - && sideFromEdge[1][2] > 0 - && sideFromEdge[1][3] > 0; - - bool allLeftFromEdge2 = sideFromEdge[2][0] > 0 - && sideFromEdge[2][1] > 0 - && sideFromEdge[2][2] > 0 - && sideFromEdge[2][3] > 0; - - if(allLeftFromEdge0 && allLeftFromEdge1 && allLeftFromEdge2 && triangleFull) - { - elt->triangleIndex |= mg_cmd_flip; - atomic_fetch_add_explicit(&tile->flipCount, 1, memory_order_relaxed); - } - else - { - atomic_store_explicit(&tile->partial, 1, memory_order_relaxed); - } - } - } - } -} - -kernel void GatherKernel(const device mg_shape_queue* shapeQueueBuffer [[buffer(0)]], - const device mg_tile* tilesBuffer [[buffer(1)]], - const device mg_tile_elt* eltBuffer [[buffer(2)]], - device int* tileCounters [[buffer(3)]], - device mg_tile_cmd* tileArrayBuffer [[buffer(4)]], - constant int* shapeCount [[buffer(5)]], - constant uint2* viewport [[buffer(6)]], - uint2 gid [[thread_position_in_grid]]) -{ - uint2 tilesMatrixDim = (*viewport - 1) / RENDERER_TILE_SIZE + 1; - int nTilesX = tilesMatrixDim.x; - - int2 tileCoord = int2(gid); - int tileIndex = tileCoord.y * nTilesX + tileCoord.x; - - device mg_tile_cmd* tileArray = &tileArrayBuffer[tileIndex * RENDERER_TILE_BUFFER_COUNT]; - - int count = 0; - for(int shapeIndex = 0; shapeIndex < shapeCount[0]; shapeIndex++) - { - const device mg_shape_queue* shapeQueue = &shapeQueueBuffer[shapeIndex]; - const device mg_tile* tiles = &tilesBuffer[shapeQueue->tiles]; - - // get the tile queue that corresponds to our tile in the shape area - int2 tileQueueCoord = tileCoord - shapeQueue->area.xy; - - if( tileQueueCoord.x >= 0 - && tileQueueCoord.y >= 0 - && tileQueueCoord.x < shapeQueue->area.z - && tileQueueCoord.y < shapeQueue->area.w) - { - int localIndex = tileQueueCoord.y * shapeQueue->area.z + tileQueueCoord.x; - const device mg_tile* tile = &tiles[localIndex]; - - if(atomic_load_explicit(&tile->partial, memory_order_relaxed) == 0) - { - if(atomic_load_explicit(&tile->flipCount, memory_order_relaxed) & 0x01) - { - if(tile->color.a == 1 && !tile->textured) - { - //NOTE: tile is full covered by a solid color, reset counter and push a color command - int firstEltIndex = *(device int*)&tile->firstElt; - const device mg_tile_elt* elt = &eltBuffer[firstEltIndex]; - - count = 0; - tileArray[count] = mg_cmd_color | (elt->triangleIndex & ~MG_TILE_CMD_MASK); - count++; - continue; - } - } - else - { - //NOTE: tile is fully uncovered, skip that shape - continue; - } - } - - int firstEltIndex = *(device int*)&tile->firstElt; - const device mg_tile_elt* elt = 0; - - for(int eltIndex = firstEltIndex; eltIndex >= 0; eltIndex = elt->next) - { - elt = &eltBuffer[eltIndex]; - eltIndex = elt->next; - - tileArray[count] = elt->triangleIndex; - count++; - } - } - } - tileCounters[tileIndex] = count; -} - -kernel void RenderKernel(const device uint* tileCounters [[buffer(0)]], - const device mg_tile_cmd* tileArrayBuffer [[buffer(1)]], - const device mg_triangle_data* triangleArray [[buffer(2)]], - - constant int* useTexture [[buffer(3)]], - constant float* scaling [[buffer(4)]], - - texture2d outTexture [[texture(0)]], - texture2d texAtlas [[texture(1)]], - - uint2 gid [[thread_position_in_grid]], - uint2 tgid [[threadgroup_position_in_grid]], - uint2 threadsPerThreadgroup [[threads_per_threadgroup]], - uint2 gridSize [[threads_per_grid]]) -{ - //TODO: guard against thread group size not equal to tile size? - const int2 pixelCoord = int2(gid); - const uint2 tileCoord = uint2(pixelCoord)/ RENDERER_TILE_SIZE; - const uint2 tilesMatrixDim = (gridSize - 1) / RENDERER_TILE_SIZE + 1; - const uint tileIndex = tileCoord.y * tilesMatrixDim.x + tileCoord.x; - const uint tileCounter = min(tileCounters[tileIndex], (uint)RENDERER_TILE_BUFFER_COUNT); - -#ifdef RENDERER_DEBUG_TILES - //NOTE(martin): color code debug values and show the tile grid - { - float4 fragColor = float4(0); - - if( pixelCoord.x % 16 == 0 - ||pixelCoord.y % 16 == 0) - { - fragColor = float4(0, 0, 0, 1); - } - else if(tileCounters[tileIndex] == 0xffffu) - { - fragColor = float4(1, 0, 1, 1); - } - else if(tileCounter != 0u) - { - fragColor = float4(0, 1, 0, 1); - } - else - { - fragColor = float4(1, 0, 0, 1); - } - outTexture.write(fragColor, gid); - return; - } -#endif - - const int subPixelFactor = 16; - const int2 centerPoint = int2((float2(pixelCoord) + float2(0.5, 0.5)) * subPixelFactor); - - const int sampleCount = 8; - int2 samplePoints[sampleCount] = {centerPoint + int2(1, 3), - centerPoint + int2(-1, -3), - centerPoint + int2(5, -1), - centerPoint + int2(-3, 5), - centerPoint + int2(-5, -5), - centerPoint + int2(-7, 1), - centerPoint + int2(3, -7), - centerPoint + int2(7, 7)}; - - float4 sampleColor[sampleCount]; - float4 currentColor[sampleCount]; - int flipCount[sampleCount]; - - for(int i=0; i= RENDERER_TILE_BUFFER_COUNT) - { - outTexture.write(float4(1, 0, 0, 1), gid); - return; - } - - for(uint tileArrayIndex=0; tileArrayIndex < tileCounter; tileArrayIndex++) - { - mg_tile_cmd cmd = tileArrayBuffer[RENDERER_TILE_BUFFER_COUNT * tileIndex + tileArrayIndex]; - - int cmdKind = cmd & MG_TILE_CMD_MASK; - int triangleIndex = cmd & ~(MG_TILE_CMD_MASK); - const device mg_triangle_data* triangle = &triangleArray[triangleIndex]; - - if(currentShapeIndex != triangle->shapeIndex) - { - for(int sampleIndex = 0; sampleIndexcolor; - nextColor.rgb *= nextColor.a; - - if(useTexture[0]) - { - int2 samplePoint = samplePoints[sampleIndex]; - float3 sampleFP = float3(float2(samplePoint).xy/(subPixelFactor*2.), 1); - float2 uv = (triangle->uvTransform * sampleFP).xy; - - constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); - float4 texColor = texAtlas.sample(smp, uv); - - texColor.rgb *= texColor.a; - nextColor *= texColor; - } - - currentColor[sampleIndex] = sampleColor[sampleIndex]*(1.-nextColor.a) + nextColor; - flipCount[sampleIndex] = 0; - } - currentShapeIndex = triangle->shapeIndex; - } - - switch(cmdKind) - { - case mg_cmd_color: - { - for(int sampleIndex=0; sampleIndexcolor; - nextColor.rgb *= nextColor.a; - sampleColor[sampleIndex] = nextColor; - flipCount[sampleIndex] = 0; - } - } break; - - case mg_cmd_flip: - { - for(int sampleIndex=0; sampleIndexp0; - int2 p1 = triangle->p1; - int2 p2 = triangle->p2; - - int cw = triangle->cw; - - int bias0 = triangle->bias0; - int bias1 = triangle->bias1; - int bias2 = triangle->bias2; - - float4 cubic0 = triangle->cubic0; - float4 cubic1 = triangle->cubic1; - float4 cubic2 = triangle->cubic2; - - bool fullTriangle = triangle->full; - - int4 clip = triangle->box; - - for(int sampleIndex = 0; sampleIndex < sampleCount; sampleIndex++) - { - int2 samplePoint = samplePoints[sampleIndex]; - - if( samplePoint.x < clip.x - || samplePoint.x > clip.z - || samplePoint.y < clip.y - || samplePoint.y > clip.w) - { - continue; - } - - int w0 = cw*orient2d(p1, p2, samplePoint); - int w1 = cw*orient2d(p2, p0, samplePoint); - int w2 = cw*orient2d(p0, p1, samplePoint); - - if((w0+bias0) >= 0 && (w1+bias1) >= 0 && (w2+bias2) >= 0) - { - float4 cubic = (cubic0*w0 + cubic1*w1 + cubic2*w2)/(w0+w1+w2); - - if( fullTriangle - ||(cubic.w*(cubic.x*cubic.x*cubic.x - cubic.y*cubic.z) <= 0)) - { - flipCount[sampleIndex]++; - } - } - } - } break; - } - } - - float4 pixelColor = float4(0); - for(int sampleIndex = 0; sampleIndex < sampleCount; sampleIndex++) - { - if(flipCount[sampleIndex] & 0x01) - { - sampleColor[sampleIndex] = currentColor[sampleIndex]; - } - pixelColor += sampleColor[sampleIndex]; - } - - outTexture.write(pixelColor/float(sampleCount), gid); -}