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

[osx, canvas] use 32 pixels-wide tiles

This commit is contained in:
Martin Fouilleul 2023-07-11 14:27:15 +02:00
parent 9e8a0f5f69
commit fc07a0d018
2 changed files with 13 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
src/mtl_renderer.m
View File

@ -17,7 +17,7 @@
#include"mtl_renderer.h" #include"mtl_renderer.h"
const int MG_MTL_INPUT_BUFFERS_COUNT = 3, const int MG_MTL_INPUT_BUFFERS_COUNT = 3,
MG_MTL_TILE_SIZE = 16, MG_MTL_TILE_SIZE = 32,
MG_MTL_MSAA_COUNT = 8; MG_MTL_MSAA_COUNT = 8;
typedef struct mg_mtl_canvas_backend typedef struct mg_mtl_canvas_backend
@ -914,7 +914,7 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
[mergeEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:11]; [mergeEncoder setBuffer:backend->logOffsetBuffer[backend->bufferIndex] offset:0 atIndex:11];
MTLSize mergeGridSize = MTLSizeMake(nTilesX, nTilesY, 1); MTLSize mergeGridSize = MTLSizeMake(nTilesX, nTilesY, 1);
MTLSize mergeGroupSize = MTLSizeMake(16, 16, 1); MTLSize mergeGroupSize = MTLSizeMake(MG_MTL_TILE_SIZE, MG_MTL_TILE_SIZE, 1);
[mergeEncoder dispatchThreads: mergeGridSize threadsPerThreadgroup: mergeGroupSize]; [mergeEncoder dispatchThreads: mergeGridSize threadsPerThreadgroup: mergeGroupSize];
[mergeEncoder endEncoding]; [mergeEncoder endEncoding];
@ -946,7 +946,7 @@ void mg_mtl_render_batch(mg_mtl_canvas_backend* backend,
[rasterEncoder setBytes: &useTexture length:sizeof(int) atIndex: 9]; [rasterEncoder setBytes: &useTexture length:sizeof(int) atIndex: 9];
MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1); MTLSize rasterGridSize = MTLSizeMake(viewportSize.x, viewportSize.y, 1);
MTLSize rasterGroupSize = MTLSizeMake(16, 16, 1); MTLSize rasterGroupSize = MTLSizeMake(MG_MTL_TILE_SIZE, MG_MTL_TILE_SIZE, 1);
// [rasterEncoder dispatchThreads: rasterGridSize threadsPerThreadgroup: rasterGroupSize]; // [rasterEncoder dispatchThreads: rasterGridSize threadsPerThreadgroup: rasterGroupSize];
[rasterEncoder dispatchThreadgroupsWithIndirectBuffer: backend->rasterDispatchBuffer [rasterEncoder dispatchThreadgroupsWithIndirectBuffer: backend->rasterDispatchBuffer
@ -999,7 +999,7 @@ void mg_mtl_canvas_resize(mg_mtl_canvas_backend* backend, vec2 size)
MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init];
texDesc.textureType = MTLTextureType2D; texDesc.textureType = MTLTextureType2D;
texDesc.storageMode = MTLStorageModePrivate; texDesc.storageMode = MTLStorageModePrivate;
texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite; texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite | MTLTextureUsageRenderTarget;
texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm; texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm;
texDesc.width = size.x; texDesc.width = size.x;
texDesc.height = size.y; texDesc.height = size.y;
@ -1407,7 +1407,7 @@ mg_canvas_backend* mtl_canvas_backend_create(mg_mtl_surface* surface)
MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init]; MTLTextureDescriptor* texDesc = [[MTLTextureDescriptor alloc] init];
texDesc.textureType = MTLTextureType2D; texDesc.textureType = MTLTextureType2D;
texDesc.storageMode = MTLStorageModePrivate; texDesc.storageMode = MTLStorageModePrivate;
texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite; texDesc.usage = MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite | MTLTextureUsageRenderTarget;
texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm; texDesc.pixelFormat = MTLPixelFormatRGBA8Unorm;
texDesc.width = backend->frameSize.x; texDesc.width = backend->frameSize.x;
texDesc.height = backend->frameSize.y; texDesc.height = backend->frameSize.y;

9
src/mtl_renderer.metal
View File

@ -1611,6 +1611,13 @@ kernel void mtl_raster(const device mg_mtl_screen_tile* screenTilesBuffer [[buff
} }
opIndex = op->next; opIndex = op->next;
} }
/*
if((pixelCoord.x % tileSize[0] == 0) || (pixelCoord.y % tileSize[0] == 0))
{
color = float4(0, 0, 0, 1);
}
//*/
outTexture.write(color, pixelCoord); outTexture.write(color, pixelCoord);
} }
@ -1631,7 +1638,7 @@ vertex vs_out mtl_vertex_shader(ushort vid [[vertex_id]])
return(out); return(out);
} }
fragment float4 mtl_fragment_shader(vs_out i [[stage_in]], texture2d<float> tex [[texture(0)]]) fragment float4 mtl_fragment_shader(vs_out i [[stage_in]], texture2d<float, access::sample> tex [[texture(0)]])
{ {
constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear); constexpr sampler smp(mip_filter::nearest, mag_filter::linear, min_filter::linear);
return(tex.sample(smp, i.uv)); return(tex.sample(smp, i.uv));