[mtl canvas] remove old metal canvas stuff
This commit is contained in:
parent
03b5802529
commit
b7f5b84123
667
src/mtl_canvas.m
667
src/mtl_canvas.m
|
|
@ -1,667 +0,0 @@
|
||||||
/************************************************************//**
|
|
||||||
*
|
|
||||||
* @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_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<MTLComputePipelineState> shapePipeline;
|
|
||||||
id<MTLComputePipelineState> trianglePipeline;
|
|
||||||
id<MTLComputePipelineState> gatherPipeline;
|
|
||||||
id<MTLComputePipelineState> computePipeline;
|
|
||||||
id<MTLRenderPipelineState> renderPipeline;
|
|
||||||
|
|
||||||
mp_rect viewPort;
|
|
||||||
|
|
||||||
// triple buffering
|
|
||||||
u32 bufferIndex;
|
|
||||||
dispatch_semaphore_t bufferSemaphore;
|
|
||||||
|
|
||||||
// textures and buffers
|
|
||||||
id<MTLTexture> backbuffer;
|
|
||||||
id<MTLTexture> outTexture;
|
|
||||||
|
|
||||||
id<MTLBuffer> shapeBuffer[MG_MTL_MAX_BUFFER_AVAILABLE];
|
|
||||||
id<MTLBuffer> vertexBuffer[MG_MTL_MAX_BUFFER_AVAILABLE];
|
|
||||||
id<MTLBuffer> indexBuffer[MG_MTL_MAX_BUFFER_AVAILABLE];
|
|
||||||
id<MTLBuffer> shapeQueueBuffer;
|
|
||||||
id<MTLBuffer> triangleArray;
|
|
||||||
id<MTLBuffer> tilesBuffer;
|
|
||||||
id<MTLBuffer> tilesOffset;
|
|
||||||
id<MTLBuffer> eltBuffer;
|
|
||||||
id<MTLBuffer> eltOffset;
|
|
||||||
|
|
||||||
id<MTLBuffer> tileArrayBuffer;
|
|
||||||
id<MTLBuffer> tileCounters;
|
|
||||||
|
|
||||||
} mg_mtl_canvas_backend;
|
|
||||||
|
|
||||||
typedef struct mg_mtl_image_data
|
|
||||||
{
|
|
||||||
mg_image_data interface;
|
|
||||||
id<MTLTexture> 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<MTLRenderCommandEncoder> 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<MTLRenderCommandEncoder> 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<MTLCommandBuffer> 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<MTLBlitCommandEncoder> 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<MTLComputeCommandEncoder> 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<MTLComputeCommandEncoder> 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<MTLComputeCommandEncoder> 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<MTLComputeCommandEncoder> 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<MTLRenderCommandEncoder> 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; i<MG_MTL_MAX_BUFFER_AVAILABLE; i++)
|
|
||||||
{
|
|
||||||
backend->indexBuffer[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<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> shapeFunction = [library newFunctionWithName:@"ShapeSetup"];
|
|
||||||
id<MTLFunction> triangleFunction = [library newFunctionWithName:@"TriangleKernel"];
|
|
||||||
id<MTLFunction> gatherFunction = [library newFunctionWithName:@"GatherKernel"];
|
|
||||||
id<MTLFunction> computeFunction = [library newFunctionWithName:@"RenderKernel"];
|
|
||||||
id<MTLFunction> vertexFunction = [library newFunctionWithName:@"VertexShader"];
|
|
||||||
id<MTLFunction> 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
|
|
||||||
107
src/mtl_shader.h
107
src/mtl_shader.h
|
|
@ -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<simd/simd.h>
|
|
||||||
|
|
||||||
#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_
|
|
||||||
|
|
@ -1,550 +0,0 @@
|
||||||
|
|
||||||
#include<metal_stdlib>
|
|
||||||
#include<simd/simd.h>
|
|
||||||
#include<metal_simdgroup>
|
|
||||||
|
|
||||||
#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<float> 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<tileCount; i++)
|
|
||||||
{
|
|
||||||
tiles[i].color = shapeBuffer[gid].color;
|
|
||||||
tiles[i].textured = shapeBuffer[gid].textured;
|
|
||||||
atomic_store_explicit(&tiles[i].firstElt, -1, memory_order_relaxed);
|
|
||||||
atomic_store_explicit(&tiles[i].eltCount, 0, memory_order_relaxed);
|
|
||||||
atomic_store_explicit(&tiles[i].partial, 0, memory_order_relaxed);
|
|
||||||
atomic_store_explicit(&tiles[i].flipCount, 0, memory_order_relaxed);
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
//NOTE: setup triangle data and bucket triangle into tile queues
|
|
||||||
|
|
||||||
kernel void TriangleKernel(constant mg_vertex* vertexBuffer [[buffer(0)]],
|
|
||||||
constant uint* indexBuffer [[buffer(1)]],
|
|
||||||
constant mg_shape* shapeBuffer [[buffer(2)]],
|
|
||||||
device mg_triangle_data* triangleArray [[buffer(3)]],
|
|
||||||
device mg_shape_queue* shapeQueueBuffer [[buffer(4)]],
|
|
||||||
device mg_tile* tilesBuffer [[buffer(5)]],
|
|
||||||
device mg_tile_elt* eltBuffer [[buffer(6)]],
|
|
||||||
device volatile atomic_uint* eltOffset [[buffer(7)]],
|
|
||||||
constant float* scaling [[buffer(8)]],
|
|
||||||
uint gid [[thread_position_in_grid]])
|
|
||||||
{
|
|
||||||
//NOTE: triangle setup
|
|
||||||
uint triangleIndex = gid * 3;
|
|
||||||
|
|
||||||
uint i0 = indexBuffer[triangleIndex];
|
|
||||||
uint i1 = indexBuffer[triangleIndex+1];
|
|
||||||
uint i2 = indexBuffer[triangleIndex+2];
|
|
||||||
|
|
||||||
float2 p0 = vertexBuffer[i0].pos * scaling[0];
|
|
||||||
float2 p1 = vertexBuffer[i1].pos * scaling[0];
|
|
||||||
float2 p2 = vertexBuffer[i2].pos * scaling[0];
|
|
||||||
|
|
||||||
int shapeIndex = vertexBuffer[i0].shapeIndex;
|
|
||||||
|
|
||||||
//NOTE(martin): compute triangle bounding box and clip it
|
|
||||||
float4 clip = shapeBuffer[shapeIndex].clip * scaling[0];
|
|
||||||
float4 fbox = float4(min(min(p0, p1), p2), max(max(p0, p1), p2));
|
|
||||||
fbox = float4(max(fbox.xy, clip.xy), min(fbox.zw, clip.zw));
|
|
||||||
|
|
||||||
//NOTE(martin): fill triangle data
|
|
||||||
const float subPixelFactor = 16;
|
|
||||||
|
|
||||||
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;
|
|
||||||
|
|
||||||
int2 ip0 = int2(p0 * subPixelFactor);
|
|
||||||
int2 ip1 = int2(p1 * subPixelFactor);
|
|
||||||
int2 ip2 = int2(p2 * subPixelFactor);
|
|
||||||
|
|
||||||
triangleArray[gid].p0 = ip0;
|
|
||||||
triangleArray[gid].p1 = ip1;
|
|
||||||
triangleArray[gid].p2 = ip2;
|
|
||||||
|
|
||||||
int cw = orient2d(ip0, ip1, ip2) > 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<float, access::write> outTexture [[texture(0)]],
|
|
||||||
texture2d<float> 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<sampleCount; i++)
|
|
||||||
{
|
|
||||||
flipCount[i] = 0;
|
|
||||||
sampleColor[i] = float4(0, 0, 0, 0);
|
|
||||||
currentColor[i] = float4(0, 0, 0, 0);
|
|
||||||
}
|
|
||||||
|
|
||||||
int currentShapeIndex = -1;
|
|
||||||
|
|
||||||
if(tileCounter >= 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; sampleIndex<sampleCount; sampleIndex++)
|
|
||||||
{
|
|
||||||
if(flipCount[sampleIndex] & 0x01)
|
|
||||||
{
|
|
||||||
sampleColor[sampleIndex] = currentColor[sampleIndex];
|
|
||||||
}
|
|
||||||
|
|
||||||
float4 nextColor = triangle->color;
|
|
||||||
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; sampleIndex<sampleCount; sampleIndex++)
|
|
||||||
{
|
|
||||||
float4 nextColor = triangle->color;
|
|
||||||
nextColor.rgb *= nextColor.a;
|
|
||||||
sampleColor[sampleIndex] = nextColor;
|
|
||||||
flipCount[sampleIndex] = 0;
|
|
||||||
}
|
|
||||||
} break;
|
|
||||||
|
|
||||||
case mg_cmd_flip:
|
|
||||||
{
|
|
||||||
for(int sampleIndex=0; sampleIndex<sampleCount; sampleIndex++)
|
|
||||||
{
|
|
||||||
flipCount[sampleIndex]++;
|
|
||||||
}
|
|
||||||
} break;
|
|
||||||
|
|
||||||
case mg_cmd_triangle:
|
|
||||||
{
|
|
||||||
int2 p0 = triangle->p0;
|
|
||||||
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);
|
|
||||||
}
|
|
||||||
Loading…
Reference in New Issue