Update Files

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/Metal.h

@@ -0,0 +1,5 @@
+#pragma once
+
+#include <kinc/graphics5/graphics.h>
+#include <kinc/image.h>
+#include <kinc/math/matrix.h>

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/Metal.m.h

@@ -0,0 +1,253 @@
+#include "Metal.h"
+
+#include <kinc/color.h>
+#include <kinc/system.h>
+#include <kinc/window.h>
+
+#include <kinc/graphics5/commandlist.h>
+#include <kinc/graphics5/rendertarget.h>
+
+#import <Metal/Metal.h>
+#import <MetalKit/MTKView.h>
+
+id getMetalLayer(void);
+id getMetalDevice(void);
+id getMetalQueue(void);
+
+int renderTargetWidth;
+int renderTargetHeight;
+int newRenderTargetWidth;
+int newRenderTargetHeight;
+
+id<CAMetalDrawable> drawable;
+id<MTLTexture> depthTexture;
+int depthBits;
+int stencilBits;
+
+static kinc_g5_render_target_t fallback_render_target;
+
+id getMetalEncoder(void) {
+	return render_command_encoder;
+}
+
+void kinc_g5_internal_destroy_window(int window) {}
+
+void kinc_g5_internal_destroy(void) {}
+
+extern void kinc_g4_on_g5_internal_resize(int, int, int);
+
+void kinc_internal_resize(int window, int width, int height) {
+	kinc_g4_on_g5_internal_resize(window, width, height);
+}
+
+void kinc_g5_internal_init(void) {}
+
+void kinc_g5_internal_init_window(int window, int depthBufferBits, int stencilBufferBits, bool vsync) {
+	depthBits = depthBufferBits;
+	stencilBits = stencilBufferBits;
+	kinc_g5_render_target_init(&fallback_render_target, 32, 32, KINC_G5_RENDER_TARGET_FORMAT_32BIT, 0, 0);
+}
+
+void kinc_g5_flush(void) {}
+
+void kinc_g5_draw_indexed_vertices_instanced(int instanceCount) {}
+
+void kinc_g5_draw_indexed_vertices_instanced_from_to(int instanceCount, int start, int count) {}
+
+bool kinc_internal_metal_has_depth = false;
+
+bool kinc_internal_current_render_target_has_depth(void) {
+	return kinc_internal_metal_has_depth;
+}
+
+static void start_render_pass(void) {
+	id<MTLTexture> texture = drawable.texture;
+	MTLRenderPassDescriptor *renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor];
+	renderPassDescriptor.colorAttachments[0].texture = texture;
+	renderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionClear;
+	renderPassDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore;
+	renderPassDescriptor.colorAttachments[0].clearColor = MTLClearColorMake(0.0, 0.0, 0.0, 1.0);
+	renderPassDescriptor.depthAttachment.clearDepth = 1;
+	renderPassDescriptor.depthAttachment.loadAction = MTLLoadActionClear;
+	renderPassDescriptor.depthAttachment.storeAction = MTLStoreActionStore;
+	renderPassDescriptor.depthAttachment.texture = depthTexture;
+	renderPassDescriptor.stencilAttachment.clearStencil = 0;
+	renderPassDescriptor.stencilAttachment.loadAction = MTLLoadActionDontCare;
+	renderPassDescriptor.stencilAttachment.storeAction = MTLStoreActionDontCare;
+	renderPassDescriptor.stencilAttachment.texture = depthTexture;
+	
+	render_command_encoder = [command_buffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
+}
+
+static void end_render_pass(void) {
+	[render_command_encoder endEncoding];
+	render_command_encoder = nil;
+}
+
+void kinc_g5_begin(kinc_g5_render_target_t *renderTarget, int window) {
+	CAMetalLayer *metalLayer = getMetalLayer();
+	drawable = [metalLayer nextDrawable];
+
+	if (depthBits > 0 && (depthTexture == nil || depthTexture.width != drawable.texture.width || depthTexture.height != drawable.texture.height)) {
+		MTLTextureDescriptor *descriptor = [MTLTextureDescriptor new];
+		descriptor.textureType = MTLTextureType2D;
+		descriptor.width = drawable.texture.width;
+		descriptor.height = drawable.texture.height;
+		descriptor.depth = 1;
+		descriptor.pixelFormat = MTLPixelFormatDepth32Float_Stencil8;
+		descriptor.arrayLength = 1;
+		descriptor.mipmapLevelCount = 1;
+		descriptor.resourceOptions = MTLResourceStorageModePrivate;
+		descriptor.usage = MTLTextureUsageRenderTarget;
+		id<MTLDevice> device = getMetalDevice();
+		depthTexture = [device newTextureWithDescriptor:descriptor];
+		kinc_internal_metal_has_depth = true;
+	}
+	else {
+		kinc_internal_metal_has_depth = false;
+	}
+
+	id<MTLTexture> texture = drawable.texture;
+	MTLRenderPassDescriptor *renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor];
+	renderPassDescriptor.colorAttachments[0].texture = texture;
+	renderPassDescriptor.colorAttachments[0].loadAction = MTLLoadActionClear;
+	renderPassDescriptor.colorAttachments[0].storeAction = MTLStoreActionStore;
+	renderPassDescriptor.colorAttachments[0].clearColor = MTLClearColorMake(0.0, 0.0, 0.0, 1.0);
+	renderPassDescriptor.depthAttachment.clearDepth = 1;
+	renderPassDescriptor.depthAttachment.loadAction = MTLLoadActionClear;
+	renderPassDescriptor.depthAttachment.storeAction = MTLStoreActionStore;
+	renderPassDescriptor.depthAttachment.texture = depthTexture;
+	renderPassDescriptor.stencilAttachment.clearStencil = 0;
+	renderPassDescriptor.stencilAttachment.loadAction = MTLLoadActionDontCare;
+	renderPassDescriptor.stencilAttachment.storeAction = MTLStoreActionDontCare;
+	renderPassDescriptor.stencilAttachment.texture = depthTexture;
+
+	if (command_buffer != nil && render_command_encoder != nil) {
+		[render_command_encoder endEncoding];
+		[command_buffer commit];
+	}
+
+	id<MTLCommandQueue> commandQueue = getMetalQueue();
+	command_buffer = [commandQueue commandBuffer];
+	render_command_encoder = [command_buffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
+}
+
+void kinc_g5_end(int window) {}
+
+bool kinc_g5_swap_buffers(void) {
+	if (command_buffer != nil && render_command_encoder != nil) {
+		[render_command_encoder endEncoding];
+		[command_buffer presentDrawable:drawable];
+		[command_buffer commit];
+	}
+	drawable = nil;
+	command_buffer = nil;
+	render_command_encoder = nil;
+
+	return true;
+}
+
+bool kinc_window_vsynced(int window) {
+	return true;
+}
+
+void kinc_g5_internal_new_render_pass(kinc_g5_render_target_t **renderTargets, int count, bool wait, unsigned clear_flags, unsigned color, float depth,
+                                      int stencil) {
+	if (command_buffer != nil && render_command_encoder != nil) {
+		[render_command_encoder endEncoding];
+		[command_buffer commit];
+		if (wait) {
+			[command_buffer waitUntilCompleted];
+		}
+	}
+
+	MTLRenderPassDescriptor *renderPassDescriptor = [MTLRenderPassDescriptor renderPassDescriptor];
+	for (int i = 0; i < count; ++i) {
+		if (renderTargets == NULL) {
+			if (drawable == nil) {
+				renderPassDescriptor.colorAttachments[i].texture = (__bridge id<MTLTexture>)fallback_render_target.impl._tex;
+				renderPassDescriptor.depthAttachment.texture = nil;
+				renderPassDescriptor.stencilAttachment.texture = nil;
+				kinc_internal_metal_has_depth = false;
+			}
+			else {
+				renderPassDescriptor.colorAttachments[i].texture = drawable.texture;
+				renderPassDescriptor.depthAttachment.texture = depthTexture;
+				renderPassDescriptor.stencilAttachment.texture = depthTexture;
+				kinc_internal_metal_has_depth = depthTexture != nil;
+			}
+		}
+		else {
+			renderPassDescriptor.colorAttachments[i].texture = (__bridge id<MTLTexture>)renderTargets[i]->impl._tex;
+			renderPassDescriptor.depthAttachment.texture = (__bridge id<MTLTexture>)renderTargets[0]->impl._depthTex;
+			renderPassDescriptor.stencilAttachment.texture = (__bridge id<MTLTexture>)renderTargets[0]->impl._depthTex;
+			kinc_internal_metal_has_depth = renderTargets[0]->impl._depthTex != nil;
+		}
+		if (clear_flags & KINC_G5_CLEAR_COLOR) {
+			float red, green, blue, alpha;
+			kinc_color_components(color, &red, &green, &blue, &alpha);
+			renderPassDescriptor.colorAttachments[i].loadAction = MTLLoadActionClear;
+			renderPassDescriptor.colorAttachments[i].storeAction = MTLStoreActionStore;
+			renderPassDescriptor.colorAttachments[i].clearColor = MTLClearColorMake(red, green, blue, alpha);
+		}
+		else {
+			renderPassDescriptor.colorAttachments[i].loadAction = MTLLoadActionLoad;
+			renderPassDescriptor.colorAttachments[i].storeAction = MTLStoreActionStore;
+			renderPassDescriptor.colorAttachments[i].clearColor = MTLClearColorMake(0.0, 0.0, 0.0, 1.0);
+		}
+	}
+
+	if (clear_flags & KINC_G5_CLEAR_DEPTH) {
+		renderPassDescriptor.depthAttachment.clearDepth = depth;
+		renderPassDescriptor.depthAttachment.loadAction = MTLLoadActionClear;
+		renderPassDescriptor.depthAttachment.storeAction = MTLStoreActionStore;
+	}
+	else {
+		renderPassDescriptor.depthAttachment.clearDepth = 1;
+		renderPassDescriptor.depthAttachment.loadAction = MTLLoadActionLoad;
+		renderPassDescriptor.depthAttachment.storeAction = MTLStoreActionStore;
+	}
+
+	if (clear_flags & KINC_G5_CLEAR_STENCIL) {
+		renderPassDescriptor.stencilAttachment.clearStencil = stencil;
+		renderPassDescriptor.stencilAttachment.loadAction = MTLLoadActionClear;
+		renderPassDescriptor.stencilAttachment.storeAction = MTLStoreActionStore;
+	}
+	else {
+		renderPassDescriptor.stencilAttachment.clearStencil = 0;
+		renderPassDescriptor.stencilAttachment.loadAction = MTLLoadActionDontCare;
+		renderPassDescriptor.stencilAttachment.storeAction = MTLStoreActionDontCare;
+	}
+
+	id<MTLCommandQueue> commandQueue = getMetalQueue();
+	command_buffer = [commandQueue commandBuffer];
+	render_command_encoder = [command_buffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
+}
+
+bool kinc_g5_supports_raytracing(void) {
+	return false;
+}
+
+bool kinc_g5_supports_instanced_rendering(void) {
+	return true;
+}
+
+bool kinc_g5_supports_compute_shaders(void) {
+	return true;
+}
+
+bool kinc_g5_supports_blend_constants(void) {
+	return true;
+}
+
+bool kinc_g5_supports_non_pow2_textures(void) {
+	return true;
+}
+
+bool kinc_g5_render_targets_inverted_y(void) {
+	return false;
+}
+
+int kinc_g5_max_bound_textures(void) {
+	return 16;
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/commandlist.h

@@ -0,0 +1,7 @@
+#pragma once
+
+struct kinc_g5_index_buffer;
+
+typedef struct {
+	struct kinc_g5_index_buffer *current_index_buffer;
+} CommandList5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/commandlist.m.h

@@ -0,0 +1,393 @@
+#include <kinc/graphics5/commandlist.h>
+#include <kinc/graphics5/compute.h>
+#include <kinc/graphics5/constantbuffer.h>
+#include <kinc/graphics5/graphics.h>
+#include <kinc/graphics5/indexbuffer.h>
+#include <kinc/graphics5/pipeline.h>
+#include <kinc/graphics5/vertexbuffer.h>
+#include <kinc/window.h>
+
+#import <Metal/Metal.h>
+#import <MetalKit/MTKView.h>
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+id getMetalDevice(void);
+id getMetalQueue(void);
+id getMetalEncoder(void);
+
+void kinc_g5_internal_new_render_pass(kinc_g5_render_target_t **renderTargets, int count, bool wait, unsigned clear_flags, unsigned color, float depth,
+                                      int stencil);
+void kinc_g5_internal_pipeline_set(kinc_g5_pipeline_t *pipeline);
+
+void kinc_g5_command_list_init(kinc_g5_command_list_t *list) {
+	list->impl.current_index_buffer = NULL;
+}
+
+void kinc_g5_command_list_destroy(kinc_g5_command_list_t *list) {}
+
+static kinc_g5_render_target_t *lastRenderTargets[8] = {NULL, NULL, NULL, NULL, NULL, NULL, NULL, NULL};
+static kinc_g5_pipeline_t *lastPipeline = NULL;
+
+static int formatSize(MTLPixelFormat format) {
+	switch (format) {
+	case MTLPixelFormatRGBA32Float:
+		return 16;
+	case MTLPixelFormatRGBA16Float:
+		return 8;
+	case MTLPixelFormatR16Float:
+		return 2;
+	case MTLPixelFormatR8Unorm:
+		return 1;
+	default:
+		return 4;
+	}
+}
+
+void kinc_g5_command_list_begin(kinc_g5_command_list_t *list) {
+	list->impl.current_index_buffer = NULL;
+	lastRenderTargets[0] = NULL;
+}
+
+void kinc_g5_command_list_end(kinc_g5_command_list_t *list) {}
+
+void kinc_g5_command_list_clear(kinc_g5_command_list_t *list, struct kinc_g5_render_target *renderTarget, unsigned flags, unsigned color, float depth,
+                                int stencil) {
+	if (renderTarget->framebuffer_index >= 0) {
+		kinc_g5_internal_new_render_pass(NULL, 1, false, flags, color, depth, stencil);
+	}
+	else {
+		kinc_g5_internal_new_render_pass(&renderTarget, 1, false, flags, color, depth, stencil);
+	}
+}
+
+void kinc_g5_command_list_render_target_to_framebuffer_barrier(kinc_g5_command_list_t *list, struct kinc_g5_render_target *renderTarget) {}
+
+void kinc_g5_command_list_framebuffer_to_render_target_barrier(kinc_g5_command_list_t *list, struct kinc_g5_render_target *renderTarget) {}
+
+void kinc_g5_command_list_draw_indexed_vertices(kinc_g5_command_list_t *list) {
+	kinc_g5_command_list_draw_indexed_vertices_from_to(list, 0, kinc_g5_index_buffer_count(list->impl.current_index_buffer));
+}
+
+void kinc_g5_command_list_draw_indexed_vertices_from_to(kinc_g5_command_list_t *list, int start, int count) {
+	id<MTLBuffer> indexBuffer = (__bridge id<MTLBuffer>)list->impl.current_index_buffer->impl.metal_buffer;
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	[encoder drawIndexedPrimitives:MTLPrimitiveTypeTriangle
+	                    indexCount:count
+	                     indexType:(list->impl.current_index_buffer->impl.format == KINC_G5_INDEX_BUFFER_FORMAT_16BIT ? MTLIndexTypeUInt16 : MTLIndexTypeUInt32)
+	                   indexBuffer:indexBuffer
+	             indexBufferOffset:start * 4];
+}
+
+void kinc_g5_command_list_draw_indexed_vertices_from_to_from(kinc_g5_command_list_t *list, int start, int count, int vertex_offset) {
+	id<MTLBuffer> indexBuffer = (__bridge id<MTLBuffer>)list->impl.current_index_buffer->impl.metal_buffer;
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	[encoder drawIndexedPrimitives:MTLPrimitiveTypeTriangle
+	                    indexCount:count
+	                     indexType:(list->impl.current_index_buffer->impl.format == KINC_G5_INDEX_BUFFER_FORMAT_16BIT ? MTLIndexTypeUInt16 : MTLIndexTypeUInt32)
+	                   indexBuffer:indexBuffer
+	             indexBufferOffset:start * 4
+	                 instanceCount:1
+	                    baseVertex:vertex_offset
+	                  baseInstance:0];
+}
+
+void kinc_g5_command_list_draw_indexed_vertices_instanced(kinc_g5_command_list_t *list, int instanceCount) {
+	kinc_g5_command_list_draw_indexed_vertices_instanced_from_to(list, instanceCount, 0, kinc_g5_index_buffer_count(list->impl.current_index_buffer));
+}
+
+void kinc_g5_command_list_draw_indexed_vertices_instanced_from_to(kinc_g5_command_list_t *list, int instanceCount, int start, int count) {
+	id<MTLBuffer> indexBuffer = (__bridge id<MTLBuffer>)list->impl.current_index_buffer->impl.metal_buffer;
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	[encoder drawIndexedPrimitives:MTLPrimitiveTypeTriangle
+	                    indexCount:count
+	                     indexType:(list->impl.current_index_buffer->impl.format == KINC_G5_INDEX_BUFFER_FORMAT_16BIT ? MTLIndexTypeUInt16 : MTLIndexTypeUInt32)
+	                   indexBuffer:indexBuffer
+	             indexBufferOffset:start * 4
+	                 instanceCount:instanceCount
+	                    baseVertex:0
+	                  baseInstance:0];
+}
+void kinc_g5_command_list_viewport(kinc_g5_command_list_t *list, int x, int y, int width, int height) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	MTLViewport viewport;
+	viewport.originX = x;
+	viewport.originY = y;
+	viewport.width = width;
+	viewport.height = height;
+	[encoder setViewport:viewport];
+}
+
+void kinc_g5_command_list_scissor(kinc_g5_command_list_t *list, int x, int y, int width, int height) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	MTLScissorRect scissor;
+	scissor.x = x;
+	scissor.y = y;
+	int target_w = -1;
+	int target_h = -1;
+	if (lastRenderTargets[0] != NULL) {
+		target_w = lastRenderTargets[0]->texWidth;
+		target_h = lastRenderTargets[0]->texHeight;
+	}
+	else {
+		target_w = kinc_window_width(0);
+		target_h = kinc_window_height(0);
+	}
+	scissor.width = (x + width <= target_w) ? width : target_w - x;
+	scissor.height = (y + height <= target_h) ? height : target_h - y;
+	[encoder setScissorRect:scissor];
+}
+
+void kinc_g5_command_list_disable_scissor(kinc_g5_command_list_t *list) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	MTLScissorRect scissor;
+	scissor.x = 0;
+	scissor.y = 0;
+	if (lastRenderTargets[0] != NULL) {
+		scissor.width = lastRenderTargets[0]->texWidth;
+		scissor.height = lastRenderTargets[0]->texHeight;
+	}
+	else {
+		scissor.width = kinc_window_width(0);
+		scissor.height = kinc_window_height(0);
+	}
+	[encoder setScissorRect:scissor];
+}
+
+void kinc_g5_command_list_set_pipeline(kinc_g5_command_list_t *list, struct kinc_g5_pipeline *pipeline) {
+	kinc_g5_internal_pipeline_set(pipeline);
+	lastPipeline = pipeline;
+}
+
+void kinc_g5_command_list_set_blend_constant(kinc_g5_command_list_t *list, float r, float g, float b, float a) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	[encoder setBlendColorRed:r green:g blue:b alpha:a];
+}
+
+void kinc_g5_command_list_set_vertex_buffers(kinc_g5_command_list_t *list, struct kinc_g5_vertex_buffer **buffers, int *offsets, int count) {
+	kinc_g5_internal_vertex_buffer_set(buffers[0], offsets[0]);
+}
+
+void kinc_g5_command_list_set_index_buffer(kinc_g5_command_list_t *list, struct kinc_g5_index_buffer *buffer) {
+	list->impl.current_index_buffer = buffer;
+}
+
+extern bool kinc_internal_metal_has_depth;
+
+void kinc_g5_command_list_set_render_targets(kinc_g5_command_list_t *list, struct kinc_g5_render_target **targets, int count) {
+	if (targets[0]->framebuffer_index >= 0) {
+		for (int i = 0; i < 8; ++i)
+			lastRenderTargets[i] = NULL;
+		kinc_g5_internal_new_render_pass(NULL, 1, false, 0, 0, 0.0f, 0);
+	}
+	else {
+		for (int i = 0; i < count; ++i)
+			lastRenderTargets[i] = targets[i];
+		for (int i = count; i < 8; ++i)
+			lastRenderTargets[i] = NULL;
+		kinc_g5_internal_new_render_pass(targets, count, false, 0, 0, 0.0f, 0);
+	}
+}
+
+void kinc_g5_command_list_upload_index_buffer(kinc_g5_command_list_t *list, struct kinc_g5_index_buffer *buffer) {}
+
+void kinc_g5_command_list_upload_vertex_buffer(kinc_g5_command_list_t *list, struct kinc_g5_vertex_buffer *buffer) {}
+
+void kinc_g5_command_list_upload_texture(kinc_g5_command_list_t *list, struct kinc_g5_texture *texture) {}
+
+void kinc_g5_command_list_get_render_target_pixels(kinc_g5_command_list_t *list, kinc_g5_render_target_t *render_target, uint8_t *data) {
+	// Create readback buffer
+	if (render_target->impl._texReadback == NULL) {
+		id<MTLDevice> device = getMetalDevice();
+		MTLTextureDescriptor *descriptor = [MTLTextureDescriptor new];
+		descriptor.textureType = MTLTextureType2D;
+		descriptor.width = render_target->texWidth;
+		descriptor.height = render_target->texHeight;
+		descriptor.depth = 1;
+		descriptor.pixelFormat = [(__bridge id<MTLTexture>)render_target->impl._tex pixelFormat];
+		descriptor.arrayLength = 1;
+		descriptor.mipmapLevelCount = 1;
+		descriptor.usage = MTLTextureUsageUnknown;
+#ifdef KINC_APPLE_SOC
+		descriptor.resourceOptions = MTLResourceStorageModeShared;
+#else
+		descriptor.resourceOptions = MTLResourceStorageModeManaged;
+#endif
+		render_target->impl._texReadback = (__bridge_retained void *)[device newTextureWithDescriptor:descriptor];
+	}
+
+	// Copy render target to readback buffer
+	id<MTLCommandQueue> commandQueue = getMetalQueue();
+	id<MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];
+	id<MTLBlitCommandEncoder> commandEncoder = [commandBuffer blitCommandEncoder];
+	[commandEncoder copyFromTexture:(__bridge id<MTLTexture>)render_target->impl._tex
+	                    sourceSlice:0
+	                    sourceLevel:0
+	                   sourceOrigin:MTLOriginMake(0, 0, 0)
+	                     sourceSize:MTLSizeMake(render_target->texWidth, render_target->texHeight, 1)
+	                      toTexture:(__bridge id<MTLTexture>)render_target->impl._texReadback
+	               destinationSlice:0
+	               destinationLevel:0
+	              destinationOrigin:MTLOriginMake(0, 0, 0)];
+#ifndef KINC_APPLE_SOC
+	[commandEncoder synchronizeResource:(__bridge id<MTLTexture>)render_target->impl._texReadback];
+#endif
+	[commandEncoder endEncoding];
+	[commandBuffer commit];
+	[commandBuffer waitUntilCompleted];
+
+	// Read buffer
+	id<MTLTexture> tex = (__bridge id<MTLTexture>)render_target->impl._texReadback;
+	int formatByteSize = formatSize([(__bridge id<MTLTexture>)render_target->impl._tex pixelFormat]);
+	MTLRegion region = MTLRegionMake2D(0, 0, render_target->texWidth, render_target->texHeight);
+	[tex getBytes:data bytesPerRow:formatByteSize * render_target->texWidth fromRegion:region mipmapLevel:0];
+}
+
+void kinc_g5_command_list_execute(kinc_g5_command_list_t *list) {
+	if (lastRenderTargets[0] == NULL) {
+		kinc_g5_internal_new_render_pass(NULL, 1, false, 0, 0, 0.0f, 0);
+	}
+	else {
+		int count = 1;
+		while (lastRenderTargets[count] != NULL)
+			count++;
+		kinc_g5_internal_new_render_pass(lastRenderTargets, count, false, 0, 0, 0.0f, 0);
+	}
+	if (lastPipeline != NULL)
+		kinc_g5_internal_pipeline_set(lastPipeline);
+}
+
+void kinc_g5_command_list_wait_for_execution_to_finish(kinc_g5_command_list_t *list) {
+	id<MTLCommandQueue> commandQueue = getMetalQueue();
+	id<MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];
+	[commandBuffer commit];
+	[commandBuffer waitUntilCompleted];
+}
+
+void kinc_g5_command_list_set_vertex_constant_buffer(kinc_g5_command_list_t *list, struct kinc_g5_constant_buffer *buffer, int offset, size_t size) {
+	id<MTLBuffer> buf = (__bridge id<MTLBuffer>)buffer->impl._buffer;
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	[encoder setVertexBuffer:buf offset:offset atIndex:1];
+}
+
+void kinc_g5_command_list_set_fragment_constant_buffer(kinc_g5_command_list_t *list, struct kinc_g5_constant_buffer *buffer, int offset, size_t size) {
+	id<MTLBuffer> buf = (__bridge id<MTLBuffer>)buffer->impl._buffer;
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	[encoder setFragmentBuffer:buf offset:offset atIndex:0];
+}
+
+void kinc_g5_command_list_set_compute_constant_buffer(kinc_g5_command_list_t *list, struct kinc_g5_constant_buffer *buffer, int offset, size_t size) {
+	assert(compute_command_encoder != nil);
+	id<MTLBuffer> buf = (__bridge id<MTLBuffer>)buffer->impl._buffer;
+	[compute_command_encoder setBuffer:buf offset:offset atIndex:1];
+}
+
+void kinc_g5_command_list_render_target_to_texture_barrier(kinc_g5_command_list_t *list, struct kinc_g5_render_target *renderTarget) {
+#ifndef KINC_APPLE_SOC
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	[encoder textureBarrier];
+#endif
+}
+
+void kinc_g5_command_list_texture_to_render_target_barrier(kinc_g5_command_list_t *list, struct kinc_g5_render_target *renderTarget) {}
+
+void kinc_g5_command_list_set_texture(kinc_g5_command_list_t *list, kinc_g5_texture_unit_t unit, kinc_g5_texture_t *texture) {
+	id<MTLTexture> tex = (__bridge id<MTLTexture>)texture->impl._tex;
+	if (compute_command_encoder != nil) {
+		[compute_command_encoder setTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_COMPUTE]];
+	}
+	else {
+		if (unit.stages[KINC_G5_SHADER_TYPE_VERTEX] >= 0) {
+			[render_command_encoder setVertexTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_VERTEX]];
+		}
+		if (unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT] >= 0) {
+			[render_command_encoder setFragmentTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT]];
+		}
+	}
+}
+
+void kinc_g5_command_list_set_image_texture(kinc_g5_command_list_t *list, kinc_g5_texture_unit_t unit, kinc_g5_texture_t *texture) {
+	kinc_g5_command_list_set_texture(list, unit, texture);
+}
+
+bool kinc_g5_command_list_init_occlusion_query(kinc_g5_command_list_t *list, unsigned *occlusionQuery) {
+	return false;
+}
+
+void kinc_g5_command_list_delete_occlusion_query(kinc_g5_command_list_t *list, unsigned occlusionQuery) {}
+
+void kinc_g5_command_list_render_occlusion_query(kinc_g5_command_list_t *list, unsigned occlusionQuery, int triangles) {}
+
+bool kinc_g5_command_list_are_query_results_available(kinc_g5_command_list_t *list, unsigned occlusionQuery) {
+	return false;
+}
+
+void kinc_g5_command_list_get_query_result(kinc_g5_command_list_t *list, unsigned occlusionQuery, unsigned *pixelCount) {}
+
+void kinc_g5_command_list_set_render_target_face(kinc_g5_command_list_t *list, kinc_g5_render_target_t *texture, int face) {}
+
+void kinc_g5_command_list_set_texture_from_render_target(kinc_g5_command_list_t *list, kinc_g5_texture_unit_t unit, kinc_g5_render_target_t *target) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	id<MTLTexture> tex = (__bridge id<MTLTexture>)target->impl._tex;
+	if (unit.stages[KINC_G5_SHADER_TYPE_VERTEX] >= 0) {
+		[encoder setVertexTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_VERTEX]];
+	}
+	if (unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT] >= 0) {
+		[encoder setFragmentTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT]];
+	}
+}
+
+void kinc_g5_command_list_set_texture_from_render_target_depth(kinc_g5_command_list_t *list, kinc_g5_texture_unit_t unit, kinc_g5_render_target_t *target) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	id<MTLTexture> depth_tex = (__bridge id<MTLTexture>)target->impl._depthTex;
+	if (unit.stages[KINC_G5_SHADER_TYPE_VERTEX] >= 0) {
+		[encoder setVertexTexture:depth_tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_VERTEX]];
+	}
+	if (unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT] >= 0) {
+		[encoder setFragmentTexture:depth_tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT]];
+	}
+}
+
+void kinc_g5_command_list_set_sampler(kinc_g5_command_list_t *list, kinc_g5_texture_unit_t unit, kinc_g5_sampler_t *sampler) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	id<MTLSamplerState> mtl_sampler = (__bridge id<MTLSamplerState>)sampler->impl.sampler;
+
+	if (unit.stages[KINC_G5_SHADER_TYPE_VERTEX] >= 0) {
+		[encoder setVertexSamplerState:mtl_sampler atIndex:unit.stages[KINC_G5_SHADER_TYPE_VERTEX]];
+	}
+	if (unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT] >= 0) {
+		[encoder setFragmentSamplerState:mtl_sampler atIndex:unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT]];
+	}
+}
+
+void kinc_g5_command_list_set_compute_shader(kinc_g5_command_list_t *list, kinc_g5_compute_shader *shader) {
+	if (compute_command_encoder == nil) {
+		end_render_pass();
+		compute_command_encoder = [command_buffer computeCommandEncoder];
+	}
+	
+	id<MTLComputePipelineState> pipeline = (__bridge id<MTLComputePipelineState>)shader->impl._pipeline;
+	[compute_command_encoder setComputePipelineState:pipeline];
+}
+
+void kinc_g5_command_list_compute(kinc_g5_command_list_t *list, int x, int y, int z) {
+	assert(compute_command_encoder != nil);
+	
+	MTLSize perGrid;
+	perGrid.width = x;
+	perGrid.height = y;
+	perGrid.depth = z;
+	MTLSize perGroup;
+	perGroup.width = 16;
+	perGroup.height = 16;
+	perGroup.depth = 1;
+	[compute_command_encoder dispatchThreadgroups:perGrid threadsPerThreadgroup:perGroup];
+
+	[compute_command_encoder endEncoding];
+	
+	compute_command_encoder = nil;
+
+	start_render_pass();
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/compute.h

@@ -0,0 +1,8 @@
+#pragma once
+
+typedef struct kinc_g5_compute_shader_impl {
+	char name[1024];
+	void *_function;
+	void *_pipeline;
+	void *_reflection;
+} kinc_g5_compute_shader_impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/compute.m.h

@@ -0,0 +1,112 @@
+#include <kinc/graphics5/compute.h>
+#include <kinc/graphics4/texture.h>
+#include <kinc/math/core.h>
+
+#include <Metal/Metal.h>
+
+id getMetalDevice(void);
+id getMetalLibrary(void);
+
+void kinc_g5_compute_shader_init(kinc_g5_compute_shader *shader, void *_data, int length) {
+	shader->impl.name[0] = 0;
+
+	{
+		uint8_t *data = (uint8_t *)_data;
+		if (length > 1 && data[0] == '>') {
+			memcpy(shader->impl.name, data + 1, length - 1);
+			shader->impl.name[length - 1] = 0;
+		}
+		else {
+			for (int i = 3; i < length; ++i) {
+				if (data[i] == '\n') {
+					shader->impl.name[i - 3] = 0;
+					break;
+				}
+				else {
+					shader->impl.name[i - 3] = data[i];
+				}
+			}
+		}
+	}
+
+	char *data = (char *)_data;
+	id<MTLLibrary> library = nil;
+	if (length > 1 && data[0] == '>') {
+		library = getMetalLibrary();
+	}
+	else {
+		id<MTLDevice> device = getMetalDevice();
+		library = [device newLibraryWithSource:[[NSString alloc] initWithBytes:data length:length encoding:NSUTF8StringEncoding] options:nil error:nil];
+	}
+	id<MTLFunction> function = [library newFunctionWithName:[NSString stringWithCString:shader->impl.name encoding:NSUTF8StringEncoding]];
+	assert(function != nil);
+	shader->impl._function = (__bridge_retained void *)function;
+
+	id<MTLDevice> device = getMetalDevice();
+	MTLComputePipelineReflection *reflection = nil;
+	NSError *error = nil;
+	shader->impl._pipeline = (__bridge_retained void *)[device newComputePipelineStateWithFunction:function
+	                                                                                       options:MTLPipelineOptionBufferTypeInfo
+	                                                                                    reflection:&reflection
+	                                                                                         error:&error];
+	if (error != nil)
+		NSLog(@"%@", [error localizedDescription]);
+	assert(shader->impl._pipeline != NULL && !error);
+	shader->impl._reflection = (__bridge_retained void *)reflection;
+}
+
+void kinc_g5_compute_shader_destroy(kinc_g5_compute_shader *shader) {
+	id<MTLFunction> function = (__bridge_transfer id<MTLFunction>)shader->impl._function;
+	function = nil;
+	shader->impl._function = NULL;
+
+	id<MTLComputePipelineState> pipeline = (__bridge_transfer id<MTLComputePipelineState>)shader->impl._pipeline;
+	pipeline = nil;
+	shader->impl._pipeline = NULL;
+
+	MTLComputePipelineReflection *reflection = (__bridge_transfer MTLComputePipelineReflection *)shader->impl._reflection;
+	reflection = nil;
+	shader->impl._reflection = NULL;
+}
+
+kinc_g5_constant_location_t kinc_g5_compute_shader_get_constant_location(kinc_g5_compute_shader *shader, const char *name) {
+	kinc_g5_constant_location_t location;
+	location.impl.vertexOffset = -1;
+	location.impl.fragmentOffset = -1;
+	location.impl.computeOffset = -1;
+
+	MTLComputePipelineReflection *reflection = (__bridge MTLComputePipelineReflection *)shader->impl._reflection;
+
+	for (MTLArgument *arg in reflection.arguments) {
+		if (arg.type == MTLArgumentTypeBuffer && [arg.name isEqualToString:@"uniforms"]) {
+			if ([arg bufferDataType] == MTLDataTypeStruct) {
+				MTLStructType *structObj = [arg bufferStructType];
+				for (MTLStructMember *member in structObj.members) {
+					if (strcmp([[member name] UTF8String], name) == 0) {
+						location.impl.computeOffset = (int)[member offset];
+						break;
+					}
+				}
+			}
+			break;
+		}
+	}
+
+	return location;
+}
+
+kinc_g5_texture_unit_t kinc_g5_compute_shader_get_texture_unit(kinc_g5_compute_shader *shader, const char *name) {
+	kinc_g5_texture_unit_t unit;
+	for (int i = 0; i < KINC_G5_SHADER_TYPE_COUNT; ++i) {
+		unit.stages[i] = -1;
+	}
+
+	MTLComputePipelineReflection *reflection = (__bridge MTLComputePipelineReflection *)shader->impl._reflection;
+	for (MTLArgument *arg in reflection.arguments) {
+		if ([arg type] == MTLArgumentTypeTexture && strcmp([[arg name] UTF8String], name) == 0) {
+			unit.stages[KINC_G5_SHADER_TYPE_COMPUTE] = (int)[arg index];
+		}
+	}
+
+	return unit;
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/constantbuffer.h

@@ -0,0 +1,8 @@
+#pragma once
+
+typedef struct {
+	void *_buffer;
+	int lastStart;
+	int lastCount;
+	int mySize;
+} ConstantBuffer5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/constantbuffer.m.h

@@ -0,0 +1,40 @@
+#include <kinc/graphics5/constantbuffer.h>
+
+#import <Metal/Metal.h>
+
+id getMetalDevice(void);
+
+bool kinc_g5_transposeMat3 = true;
+bool kinc_g5_transposeMat4 = true;
+
+void kinc_g5_constant_buffer_init(kinc_g5_constant_buffer_t *buffer, int size) {
+	buffer->impl.mySize = size;
+	buffer->data = NULL;
+	buffer->impl._buffer = (__bridge_retained void *)[getMetalDevice() newBufferWithLength:size options:MTLResourceOptionCPUCacheModeDefault];
+}
+
+void kinc_g5_constant_buffer_destroy(kinc_g5_constant_buffer_t *buffer) {
+	id<MTLBuffer> buf = (__bridge_transfer id<MTLBuffer>)buffer->impl._buffer;
+	buf = nil;
+	buffer->impl._buffer = NULL;
+}
+
+void kinc_g5_constant_buffer_lock_all(kinc_g5_constant_buffer_t *buffer) {
+	kinc_g5_constant_buffer_lock(buffer, 0, kinc_g5_constant_buffer_size(buffer));
+}
+
+void kinc_g5_constant_buffer_lock(kinc_g5_constant_buffer_t *buffer, int start, int count) {
+	buffer->impl.lastStart = start;
+	buffer->impl.lastCount = count;
+	id<MTLBuffer> buf = (__bridge id<MTLBuffer>)buffer->impl._buffer;
+	uint8_t *data = (uint8_t *)[buf contents];
+	buffer->data = &data[start];
+}
+
+void kinc_g5_constant_buffer_unlock(kinc_g5_constant_buffer_t *buffer) {
+	buffer->data = NULL;
+}
+
+int kinc_g5_constant_buffer_size(kinc_g5_constant_buffer_t *buffer) {
+	return buffer->impl.mySize;
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/graphics.h

@@ -0,0 +1,6 @@
+#pragma once
+
+#include <kinc/backend/graphics5/indexbuffer.h>
+#include <kinc/backend/graphics5/rendertarget.h>
+#include <kinc/backend/graphics5/texture.h>
+#include <kinc/backend/graphics5/vertexbuffer.h>

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/indexbuffer.h

@@ -0,0 +1,10 @@
+#pragma once
+
+typedef struct {
+	void *metal_buffer;
+	int count;
+	bool gpu_memory;
+	int format;
+	int last_start;
+	int last_count;
+} IndexBuffer5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/indexbuffer.m.h

@@ -0,0 +1,73 @@
+#include <kinc/graphics5/graphics.h>
+#include <kinc/graphics5/indexbuffer.h>
+
+#import <Metal/Metal.h>
+
+id getMetalDevice(void);
+
+void kinc_g5_index_buffer_init(kinc_g5_index_buffer_t *buffer, int indexCount, kinc_g5_index_buffer_format_t format, bool gpuMemory) {
+	buffer->impl.count = indexCount;
+	buffer->impl.gpu_memory = gpuMemory;
+	buffer->impl.format = format;
+	buffer->impl.last_start = 0;
+	buffer->impl.last_count = indexCount;
+
+	id<MTLDevice> device = getMetalDevice();
+	MTLResourceOptions options = MTLResourceCPUCacheModeWriteCombined;
+#ifdef KINC_APPLE_SOC
+	options |= MTLResourceStorageModeShared;
+#else
+	if (gpuMemory) {
+		options |= MTLResourceStorageModeManaged;
+	}
+	else {
+		options |= MTLResourceStorageModeShared;
+	}
+#endif
+	buffer->impl.metal_buffer = (__bridge_retained void *)[device
+	    newBufferWithLength:(format == KINC_G5_INDEX_BUFFER_FORMAT_16BIT ? sizeof(uint16_t) * indexCount : sizeof(uint32_t) * indexCount)
+	                options:options];
+}
+
+void kinc_g5_index_buffer_destroy(kinc_g5_index_buffer_t *buffer) {
+	id<MTLBuffer> buf = (__bridge_transfer id<MTLBuffer>)buffer->impl.metal_buffer;
+	buf = nil;
+	buffer->impl.metal_buffer = NULL;
+}
+
+static int kinc_g5_internal_index_buffer_stride(kinc_g5_index_buffer_t *buffer) {
+	return buffer->impl.format == KINC_G5_INDEX_BUFFER_FORMAT_16BIT ? 2 : 4;
+}
+
+void *kinc_g5_index_buffer_lock_all(kinc_g5_index_buffer_t *buffer) {
+	return kinc_g5_index_buffer_lock(buffer, 0, kinc_g5_index_buffer_count(buffer));
+}
+
+void *kinc_g5_index_buffer_lock(kinc_g5_index_buffer_t *buffer, int start, int count) {
+	buffer->impl.last_start = start;
+	buffer->impl.last_count = count;
+
+	id<MTLBuffer> metal_buffer = (__bridge id<MTLBuffer>)buffer->impl.metal_buffer;
+	uint8_t *data = (uint8_t *)[metal_buffer contents];
+	return &data[start * kinc_g5_internal_index_buffer_stride(buffer)];
+}
+
+void kinc_g5_index_buffer_unlock_all(kinc_g5_index_buffer_t *buffer) {
+	kinc_g5_index_buffer_unlock(buffer, buffer->impl.last_count);
+}
+
+void kinc_g5_index_buffer_unlock(kinc_g5_index_buffer_t *buffer, int count) {
+#ifndef KINC_APPLE_SOC
+	if (buffer->impl.gpu_memory) {
+		id<MTLBuffer> metal_buffer = (__bridge id<MTLBuffer>)buffer->impl.metal_buffer;
+		NSRange range;
+		range.location = buffer->impl.last_start * kinc_g5_internal_index_buffer_stride(buffer);
+		range.length = count * kinc_g5_internal_index_buffer_stride(buffer);
+		[metal_buffer didModifyRange:range];
+	}
+#endif
+}
+
+int kinc_g5_index_buffer_count(kinc_g5_index_buffer_t *buffer) {
+	return buffer->impl.count;
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/metalunit.m

@@ -0,0 +1,22 @@
+#import <Metal/Metal.h>
+#import <MetalKit/MTKView.h>
+
+static id<MTLCommandBuffer> command_buffer = nil;
+static id<MTLRenderCommandEncoder> render_command_encoder = nil;
+static id<MTLComputeCommandEncoder> compute_command_encoder = nil;
+
+static void start_render_pass(void);
+static void end_render_pass(void);
+
+#include "Metal.m.h"
+#include "commandlist.m.h"
+#include "compute.m.h"
+#include "constantbuffer.m.h"
+#include "indexbuffer.m.h"
+#include "pipeline.m.h"
+#include "raytrace.m.h"
+#include "rendertarget.m.h"
+#include "sampler.m.h"
+#include "shader.m.h"
+#include "texture.m.h"
+#include "vertexbuffer.m.h"

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/pipeline.h

@@ -0,0 +1,24 @@
+#pragma once
+
+struct kinc_g5_shader;
+
+typedef struct {
+	struct kinc_g5_shader *vertexShader;
+	struct kinc_g5_shader *fragmentShader;
+	void *_pipeline;
+	void *_pipelineDepth;
+	void *_reflection;
+	void *_depthStencil;
+	void *_depthStencilNone;
+	// void _set();
+} PipelineState5Impl;
+
+typedef struct {
+	int a;
+} ComputePipelineState5Impl;
+
+typedef struct {
+	int vertexOffset;
+	int fragmentOffset;
+	int computeOffset;
+} ConstantLocation5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/pipeline.m.h

@@ -0,0 +1,433 @@
+#include <kinc/graphics5/pipeline.h>
+#include <kinc/graphics5/shader.h>
+#include <kinc/log.h>
+
+#import <Metal/Metal.h>
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+id getMetalDevice(void);
+id getMetalEncoder(void);
+
+static MTLBlendFactor convert_blending_factor(kinc_g5_blending_factor_t factor) {
+	switch (factor) {
+	case KINC_G5_BLEND_ONE:
+		return MTLBlendFactorOne;
+	case KINC_G5_BLEND_ZERO:
+		return MTLBlendFactorZero;
+	case KINC_G5_BLEND_SOURCE_ALPHA:
+		return MTLBlendFactorSourceAlpha;
+	case KINC_G5_BLEND_DEST_ALPHA:
+		return MTLBlendFactorDestinationAlpha;
+	case KINC_G5_BLEND_INV_SOURCE_ALPHA:
+		return MTLBlendFactorOneMinusSourceAlpha;
+	case KINC_G5_BLEND_INV_DEST_ALPHA:
+		return MTLBlendFactorOneMinusDestinationAlpha;
+	case KINC_G5_BLEND_SOURCE_COLOR:
+		return MTLBlendFactorSourceColor;
+	case KINC_G5_BLEND_DEST_COLOR:
+		return MTLBlendFactorDestinationColor;
+	case KINC_G5_BLEND_INV_SOURCE_COLOR:
+		return MTLBlendFactorOneMinusSourceColor;
+	case KINC_G5_BLEND_INV_DEST_COLOR:
+		return MTLBlendFactorOneMinusDestinationColor;
+	case KINC_G5_BLEND_CONSTANT:
+		return MTLBlendFactorBlendColor;
+	case KINC_G5_BLEND_INV_CONSTANT:
+		return MTLBlendFactorOneMinusBlendColor;
+	}
+}
+
+static MTLBlendOperation convert_blending_operation(kinc_g5_blending_operation_t op) {
+	switch (op) {
+	case KINC_G5_BLENDOP_ADD:
+		return MTLBlendOperationAdd;
+	case KINC_G5_BLENDOP_SUBTRACT:
+		return MTLBlendOperationSubtract;
+	case KINC_G5_BLENDOP_REVERSE_SUBTRACT:
+		return MTLBlendOperationReverseSubtract;
+	case KINC_G5_BLENDOP_MIN:
+		return MTLBlendOperationMin;
+	case KINC_G5_BLENDOP_MAX:
+		return MTLBlendOperationMax;
+	}
+}
+
+static MTLCompareFunction convert_compare_mode(kinc_g5_compare_mode_t compare) {
+	switch (compare) {
+	case KINC_G5_COMPARE_MODE_ALWAYS:
+		return MTLCompareFunctionAlways;
+	case KINC_G5_COMPARE_MODE_NEVER:
+		return MTLCompareFunctionNever;
+	case KINC_G5_COMPARE_MODE_EQUAL:
+		return MTLCompareFunctionEqual;
+	case KINC_G5_COMPARE_MODE_NOT_EQUAL:
+		return MTLCompareFunctionNotEqual;
+	case KINC_G5_COMPARE_MODE_LESS:
+		return MTLCompareFunctionLess;
+	case KINC_G5_COMPARE_MODE_LESS_EQUAL:
+		return MTLCompareFunctionLessEqual;
+	case KINC_G5_COMPARE_MODE_GREATER:
+		return MTLCompareFunctionGreater;
+	case KINC_G5_COMPARE_MODE_GREATER_EQUAL:
+		return MTLCompareFunctionGreaterEqual;
+	}
+}
+
+static MTLCullMode convert_cull_mode(kinc_g5_cull_mode_t cull) {
+	switch (cull) {
+	case KINC_G5_CULL_MODE_CLOCKWISE:
+		return MTLCullModeFront;
+	case KINC_G5_CULL_MODE_COUNTERCLOCKWISE:
+		return MTLCullModeBack;
+	case KINC_G5_CULL_MODE_NEVER:
+		return MTLCullModeNone;
+	}
+}
+
+static MTLPixelFormat convert_render_target_format(kinc_g5_render_target_format_t format) {
+	switch (format) {
+	case KINC_G5_RENDER_TARGET_FORMAT_128BIT_FLOAT:
+		return MTLPixelFormatRGBA32Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_64BIT_FLOAT:
+		return MTLPixelFormatRGBA16Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_32BIT_RED_FLOAT:
+		return MTLPixelFormatR32Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_16BIT_RED_FLOAT:
+		return MTLPixelFormatR16Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_8BIT_RED:
+		return MTLPixelFormatR8Unorm;
+	case KINC_G5_RENDER_TARGET_FORMAT_32BIT:
+	default:
+		return MTLPixelFormatBGRA8Unorm;
+	}
+}
+
+void kinc_g5_pipeline_init(kinc_g5_pipeline_t *pipeline) {
+	memset(&pipeline->impl, 0, sizeof(pipeline->impl));
+}
+
+void kinc_g5_pipeline_destroy(kinc_g5_pipeline_t *pipeline) {
+	pipeline->impl._reflection = NULL;
+	pipeline->impl._depthStencil = NULL;
+
+	id<MTLRenderPipelineState> pipe = (__bridge_transfer id<MTLRenderPipelineState>)pipeline->impl._pipeline;
+	pipe = nil;
+	pipeline->impl._pipeline = NULL;
+
+	MTLRenderPipelineReflection *reflection = (__bridge_transfer MTLRenderPipelineReflection *)pipeline->impl._reflection;
+	reflection = nil;
+	pipeline->impl._reflection = NULL;
+
+	id<MTLRenderPipelineState> pipeDepth = (__bridge_transfer id<MTLRenderPipelineState>)pipeline->impl._pipelineDepth;
+	pipeDepth = nil;
+	pipeline->impl._pipelineDepth = NULL;
+
+	id<MTLDepthStencilState> depthStencil = (__bridge_transfer id<MTLDepthStencilState>)pipeline->impl._depthStencil;
+	depthStencil = nil;
+	pipeline->impl._depthStencil = NULL;
+
+	id<MTLDepthStencilState> depthStencilNone = (__bridge_transfer id<MTLDepthStencilState>)pipeline->impl._depthStencilNone;
+	depthStencilNone = nil;
+	pipeline->impl._depthStencilNone = NULL;
+}
+
+static int findAttributeIndex(NSArray<MTLVertexAttribute *> *attributes, const char *name) {
+	for (MTLVertexAttribute *attribute in attributes) {
+		if (strcmp(name, [[attribute name] UTF8String]) == 0) {
+			return (int)[attribute attributeIndex];
+		}
+	}
+	return -1;
+}
+
+void kinc_g5_pipeline_compile(kinc_g5_pipeline_t *pipeline) {
+	MTLRenderPipelineDescriptor *renderPipelineDesc = [[MTLRenderPipelineDescriptor alloc] init];
+	renderPipelineDesc.vertexFunction = (__bridge id<MTLFunction>)pipeline->vertexShader->impl.mtlFunction;
+	renderPipelineDesc.fragmentFunction = (__bridge id<MTLFunction>)pipeline->fragmentShader->impl.mtlFunction;
+	for (int i = 0; i < pipeline->colorAttachmentCount; ++i) {
+		renderPipelineDesc.colorAttachments[i].pixelFormat = convert_render_target_format(pipeline->colorAttachment[i]);
+		renderPipelineDesc.colorAttachments[i].blendingEnabled =
+		    pipeline->blend_source != KINC_G5_BLEND_ONE || pipeline->blend_destination != KINC_G5_BLEND_ZERO ||
+		    pipeline->alpha_blend_source != KINC_G5_BLEND_ONE || pipeline->alpha_blend_destination != KINC_G5_BLEND_ZERO;
+		renderPipelineDesc.colorAttachments[i].sourceRGBBlendFactor = convert_blending_factor(pipeline->blend_source);
+		renderPipelineDesc.colorAttachments[i].destinationRGBBlendFactor = convert_blending_factor(pipeline->blend_destination);
+		renderPipelineDesc.colorAttachments[i].rgbBlendOperation = convert_blending_operation(pipeline->blend_operation);
+		renderPipelineDesc.colorAttachments[i].sourceAlphaBlendFactor = convert_blending_factor(pipeline->alpha_blend_source);
+		renderPipelineDesc.colorAttachments[i].destinationAlphaBlendFactor = convert_blending_factor(pipeline->alpha_blend_destination);
+		renderPipelineDesc.colorAttachments[i].alphaBlendOperation = convert_blending_operation(pipeline->alpha_blend_operation);
+		renderPipelineDesc.colorAttachments[i].writeMask =
+		    (pipeline->colorWriteMaskRed[i] ? MTLColorWriteMaskRed : 0) | (pipeline->colorWriteMaskGreen[i] ? MTLColorWriteMaskGreen : 0) |
+		    (pipeline->colorWriteMaskBlue[i] ? MTLColorWriteMaskBlue : 0) | (pipeline->colorWriteMaskAlpha[i] ? MTLColorWriteMaskAlpha : 0);
+	}
+	renderPipelineDesc.depthAttachmentPixelFormat = MTLPixelFormatInvalid;
+	renderPipelineDesc.stencilAttachmentPixelFormat = MTLPixelFormatInvalid;
+
+	float offset = 0;
+	MTLVertexDescriptor *vertexDescriptor = [[MTLVertexDescriptor alloc] init];
+
+	for (int i = 0; i < pipeline->inputLayout[0]->size; ++i) {
+		int index = findAttributeIndex(renderPipelineDesc.vertexFunction.vertexAttributes, pipeline->inputLayout[0]->elements[i].name);
+
+		if (index < 0) {
+			kinc_log(KINC_LOG_LEVEL_WARNING, "Could not find vertex attribute %s\n", pipeline->inputLayout[0]->elements[i].name);
+		}
+
+		if (index >= 0) {
+			vertexDescriptor.attributes[index].bufferIndex = 0;
+			vertexDescriptor.attributes[index].offset = offset;
+		}
+
+		offset += kinc_g4_vertex_data_size(pipeline->inputLayout[0]->elements[i].data);
+		if (index >= 0) {
+			switch (pipeline->inputLayout[0]->elements[i].data) {
+			case KINC_G4_VERTEX_DATA_NONE:
+				assert(false);
+				break;
+			case KINC_G4_VERTEX_DATA_F32_1X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatFloat;
+				break;
+			case KINC_G4_VERTEX_DATA_F32_2X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatFloat2;
+				break;
+			case KINC_G4_VERTEX_DATA_F32_3X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatFloat3;
+				break;
+			case KINC_G4_VERTEX_DATA_F32_4X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatFloat4;
+				break;
+			case KINC_G4_VERTEX_DATA_F32_4X4:
+				assert(false);
+				break;
+			case KINC_G4_VERTEX_DATA_I8_1X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatChar;
+				break;
+			case KINC_G4_VERTEX_DATA_U8_1X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUChar;
+				break;
+			case KINC_G4_VERTEX_DATA_I8_1X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatCharNormalized;
+				break;
+			case KINC_G4_VERTEX_DATA_U8_1X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUCharNormalized;
+				break;
+			case KINC_G4_VERTEX_DATA_I8_2X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatChar2;
+				break;
+			case KINC_G4_VERTEX_DATA_U8_2X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUChar2;
+				break;
+			case KINC_G4_VERTEX_DATA_I8_2X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatChar2Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_U8_2X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUChar2Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_I8_4X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatChar4;
+				break;
+			case KINC_G4_VERTEX_DATA_U8_4X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUChar4;
+				break;
+			case KINC_G4_VERTEX_DATA_I8_4X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatChar4Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_U8_4X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUChar4Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_I16_1X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatShort;
+				break;
+			case KINC_G4_VERTEX_DATA_U16_1X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUShort;
+				break;
+			case KINC_G4_VERTEX_DATA_I16_1X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatShortNormalized;
+				break;
+			case KINC_G4_VERTEX_DATA_U16_1X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUShortNormalized;
+				break;
+			case KINC_G4_VERTEX_DATA_I16_2X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatShort2;
+				break;
+			case KINC_G4_VERTEX_DATA_U16_2X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUShort2;
+				break;
+			case KINC_G4_VERTEX_DATA_I16_2X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatShort2Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_U16_2X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUShort2Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_I16_4X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatShort4;
+				break;
+			case KINC_G4_VERTEX_DATA_U16_4X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUShort4;
+				break;
+			case KINC_G4_VERTEX_DATA_I16_4X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatShort4Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_U16_4X_NORMALIZED:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUShort4Normalized;
+				break;
+			case KINC_G4_VERTEX_DATA_I32_1X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatInt;
+				break;
+			case KINC_G4_VERTEX_DATA_U32_1X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUInt;
+				break;
+			case KINC_G4_VERTEX_DATA_I32_2X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatInt2;
+				break;
+			case KINC_G4_VERTEX_DATA_U32_2X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUInt2;
+				break;
+			case KINC_G4_VERTEX_DATA_I32_3X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatInt3;
+				break;
+			case KINC_G4_VERTEX_DATA_U32_3X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUInt3;
+				break;
+			case KINC_G4_VERTEX_DATA_I32_4X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatInt4;
+				break;
+			case KINC_G4_VERTEX_DATA_U32_4X:
+				vertexDescriptor.attributes[index].format = MTLVertexFormatUInt4;
+				break;
+			default:
+				assert(false);
+				break;
+			}
+		}
+	}
+
+	vertexDescriptor.layouts[0].stride = offset;
+	vertexDescriptor.layouts[0].stepFunction = MTLVertexStepFunctionPerVertex;
+
+	renderPipelineDesc.vertexDescriptor = vertexDescriptor;
+
+	NSError *errors = nil;
+	MTLRenderPipelineReflection *reflection = nil;
+	id<MTLDevice> device = getMetalDevice();
+
+	pipeline->impl._pipeline = (__bridge_retained void *)[device newRenderPipelineStateWithDescriptor:renderPipelineDesc
+	                                                                                          options:MTLPipelineOptionBufferTypeInfo
+	                                                                                       reflection:&reflection
+	                                                                                            error:&errors];
+	if (errors != nil)
+		NSLog(@"%@", [errors localizedDescription]);
+	assert(pipeline->impl._pipeline && !errors);
+
+	renderPipelineDesc.depthAttachmentPixelFormat = MTLPixelFormatDepth32Float_Stencil8;
+	renderPipelineDesc.stencilAttachmentPixelFormat = MTLPixelFormatDepth32Float_Stencil8;
+	pipeline->impl._pipelineDepth = (__bridge_retained void *)[device newRenderPipelineStateWithDescriptor:renderPipelineDesc
+	                                                                                               options:MTLPipelineOptionBufferTypeInfo
+	                                                                                            reflection:&reflection
+	                                                                                                 error:&errors];
+	if (errors != nil)
+		NSLog(@"%@", [errors localizedDescription]);
+	assert(pipeline->impl._pipelineDepth && !errors);
+
+	pipeline->impl._reflection = (__bridge_retained void *)reflection;
+
+	MTLDepthStencilDescriptor *depthStencilDescriptor = [MTLDepthStencilDescriptor new];
+	depthStencilDescriptor.depthCompareFunction = convert_compare_mode(pipeline->depthMode);
+	depthStencilDescriptor.depthWriteEnabled = pipeline->depthWrite;
+	pipeline->impl._depthStencil = (__bridge_retained void *)[device newDepthStencilStateWithDescriptor:depthStencilDescriptor];
+
+	depthStencilDescriptor.depthCompareFunction = MTLCompareFunctionAlways;
+	depthStencilDescriptor.depthWriteEnabled = false;
+	pipeline->impl._depthStencilNone = (__bridge_retained void *)[device newDepthStencilStateWithDescriptor:depthStencilDescriptor];
+}
+
+bool kinc_internal_current_render_target_has_depth(void);
+
+void kinc_g5_internal_pipeline_set(kinc_g5_pipeline_t *pipeline) {
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	if (kinc_internal_current_render_target_has_depth()) {
+		id<MTLRenderPipelineState> pipe = (__bridge id<MTLRenderPipelineState>)pipeline->impl._pipelineDepth;
+		[encoder setRenderPipelineState:pipe];
+		id<MTLDepthStencilState> depthStencil = (__bridge id<MTLDepthStencilState>)pipeline->impl._depthStencil;
+		[encoder setDepthStencilState:depthStencil];
+	}
+	else {
+		id<MTLRenderPipelineState> pipe = (__bridge id<MTLRenderPipelineState>)pipeline->impl._pipeline;
+		[encoder setRenderPipelineState:pipe];
+		id<MTLDepthStencilState> depthStencil = (__bridge id<MTLDepthStencilState>)pipeline->impl._depthStencilNone;
+		[encoder setDepthStencilState:depthStencil];
+	}
+	[encoder setFrontFacingWinding:MTLWindingClockwise];
+	[encoder setCullMode:convert_cull_mode(pipeline->cullMode)];
+}
+
+kinc_g5_constant_location_t kinc_g5_pipeline_get_constant_location(kinc_g5_pipeline_t *pipeline, const char *name) {
+	if (strcmp(name, "bias") == 0) {
+		name = "bias0";
+	}
+
+	kinc_g5_constant_location_t location;
+	location.impl.vertexOffset = -1;
+	location.impl.fragmentOffset = -1;
+	location.impl.computeOffset = -1;
+	
+	MTLRenderPipelineReflection *reflection = (__bridge MTLRenderPipelineReflection *)pipeline->impl._reflection;
+
+	for (MTLArgument *arg in reflection.vertexArguments) {
+		if (arg.type == MTLArgumentTypeBuffer && [arg.name isEqualToString:@"uniforms"]) {
+			if ([arg bufferDataType] == MTLDataTypeStruct) {
+				MTLStructType *structObj = [arg bufferStructType];
+				for (MTLStructMember *member in structObj.members) {
+					if (strcmp([[member name] UTF8String], name) == 0) {
+						location.impl.vertexOffset = (int)[member offset];
+						break;
+					}
+				}
+			}
+			break;
+		}
+	}
+
+	for (MTLArgument *arg in reflection.fragmentArguments) {
+		if ([arg type] == MTLArgumentTypeBuffer && [[arg name] isEqualToString:@"uniforms"]) {
+			if ([arg bufferDataType] == MTLDataTypeStruct) {
+				MTLStructType *structObj = [arg bufferStructType];
+				for (MTLStructMember *member in structObj.members) {
+					if (strcmp([[member name] UTF8String], name) == 0) {
+						location.impl.fragmentOffset = (int)[member offset];
+						break;
+					}
+				}
+			}
+			break;
+		}
+	}
+
+	return location;
+}
+
+kinc_g5_texture_unit_t kinc_g5_pipeline_get_texture_unit(kinc_g5_pipeline_t *pipeline, const char *name) {
+	kinc_g5_texture_unit_t unit = {0};
+	for (int i = 0; i < KINC_G5_SHADER_TYPE_COUNT; ++i) {
+		unit.stages[i] = -1;
+	}
+
+	MTLRenderPipelineReflection *reflection = (__bridge MTLRenderPipelineReflection *)pipeline->impl._reflection;
+	for (MTLArgument *arg in reflection.fragmentArguments) {
+		if ([arg type] == MTLArgumentTypeTexture && strcmp([[arg name] UTF8String], name) == 0) {
+			unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT] = (int)[arg index];
+			break;
+		}
+	}
+
+	for (MTLArgument *arg in reflection.vertexArguments) {
+		if ([arg type] == MTLArgumentTypeTexture && strcmp([[arg name] UTF8String], name) == 0) {
+			unit.stages[KINC_G5_SHADER_TYPE_VERTEX] = (int)[arg index];
+			break;
+		}
+	}
+
+	return unit;
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/raytrace.h

@@ -0,0 +1,9 @@
+#pragma once
+
+typedef struct {
+	void *_raytracingPipeline;
+} kinc_raytrace_pipeline_impl_t;
+
+typedef struct {
+	void *_accelerationStructure;
+} kinc_raytrace_acceleration_structure_impl_t;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/raytrace.m.h

@@ -0,0 +1,174 @@
+#include <kinc/backend/graphics5/raytrace.h>
+#include <kinc/graphics5/commandlist.h>
+#include <kinc/graphics5/constantbuffer.h>
+#include <kinc/graphics5/graphics.h>
+#include <kinc/graphics5/indexbuffer.h>
+#include <kinc/graphics5/raytrace.h>
+#include <kinc/graphics5/vertexbuffer.h>
+
+static kinc_raytrace_acceleration_structure_t *accel;
+static kinc_raytrace_pipeline_t *pipeline;
+static kinc_g5_texture_t *output = NULL;
+static kinc_g5_constant_buffer_t *constant_buf;
+
+id getMetalDevice(void);
+id getMetalQueue(void);
+
+id<MTLComputePipelineState> _raytracing_pipeline;
+NSMutableArray *_primitive_accels;
+id<MTLAccelerationStructure> _instance_accel;
+dispatch_semaphore_t _sem;
+
+void kinc_raytrace_pipeline_init(kinc_raytrace_pipeline_t *pipeline, kinc_g5_command_list_t *command_list, void *ray_shader, int ray_shader_size,
+                                 kinc_g5_constant_buffer_t *constant_buffer) {
+	id<MTLDevice> device = getMetalDevice();
+	constant_buf = constant_buffer;
+
+	NSError *error = nil;
+	id<MTLLibrary> library = [device newLibraryWithSource:[[NSString alloc] initWithBytes:ray_shader length:ray_shader_size encoding:NSUTF8StringEncoding]
+	                                              options:nil
+	                                                error:&error];
+	if (library == nil) {
+		kinc_log(KINC_LOG_LEVEL_ERROR, "%s", error.localizedDescription.UTF8String);
+	}
+
+	MTLComputePipelineDescriptor *descriptor = [[MTLComputePipelineDescriptor alloc] init];
+	descriptor.computeFunction = [library newFunctionWithName:@"raytracingKernel"];
+	descriptor.threadGroupSizeIsMultipleOfThreadExecutionWidth = YES;
+	_raytracing_pipeline = [device newComputePipelineStateWithDescriptor:descriptor options:0 reflection:nil error:&error];
+	_sem = dispatch_semaphore_create(2);
+}
+
+void kinc_raytrace_pipeline_destroy(kinc_raytrace_pipeline_t *pipeline) {}
+
+id<MTLAccelerationStructure> create_acceleration_sctructure(MTLAccelerationStructureDescriptor *descriptor) {
+	id<MTLDevice> device = getMetalDevice();
+	id<MTLCommandQueue> queue = getMetalQueue();
+
+	MTLAccelerationStructureSizes accel_sizes = [device accelerationStructureSizesWithDescriptor:descriptor];
+	id<MTLAccelerationStructure> acceleration_structure = [device newAccelerationStructureWithSize:accel_sizes.accelerationStructureSize];
+
+	id<MTLBuffer> scratch_buffer = [device newBufferWithLength:accel_sizes.buildScratchBufferSize options:MTLResourceStorageModePrivate];
+	id<MTLCommandBuffer> command_buffer = [queue commandBuffer];
+	id<MTLAccelerationStructureCommandEncoder> command_encoder = [command_buffer accelerationStructureCommandEncoder];
+	id<MTLBuffer> compacteds_size_buffer = [device newBufferWithLength:sizeof(uint32_t) options:MTLResourceStorageModeShared];
+
+	[command_encoder buildAccelerationStructure:acceleration_structure descriptor:descriptor scratchBuffer:scratch_buffer scratchBufferOffset:0];
+
+	[command_encoder writeCompactedAccelerationStructureSize:acceleration_structure toBuffer:compacteds_size_buffer offset:0];
+
+	[command_encoder endEncoding];
+	[command_buffer commit];
+	[command_buffer waitUntilCompleted];
+
+	uint32_t compacted_size = *(uint32_t *)compacteds_size_buffer.contents;
+	id<MTLAccelerationStructure> compacted_acceleration_structure = [device newAccelerationStructureWithSize:compacted_size];
+	command_buffer = [queue commandBuffer];
+	command_encoder = [command_buffer accelerationStructureCommandEncoder];
+	[command_encoder copyAndCompactAccelerationStructure:acceleration_structure toAccelerationStructure:compacted_acceleration_structure];
+	[command_encoder endEncoding];
+	[command_buffer commit];
+
+	return compacted_acceleration_structure;
+}
+
+void kinc_raytrace_acceleration_structure_init(kinc_raytrace_acceleration_structure_t *accel, kinc_g5_command_list_t *command_list, kinc_g5_vertex_buffer_t *vb,
+                                               kinc_g5_index_buffer_t *ib) {
+#if !TARGET_OS_IPHONE
+	MTLResourceOptions options = MTLResourceStorageModeManaged;
+#else
+	MTLResourceOptions options = MTLResourceStorageModeShared;
+#endif
+
+	MTLAccelerationStructureTriangleGeometryDescriptor *descriptor = [MTLAccelerationStructureTriangleGeometryDescriptor descriptor];
+	descriptor.indexType = MTLIndexTypeUInt32;
+	descriptor.indexBuffer = (__bridge id<MTLBuffer>)ib->impl.metal_buffer;
+	descriptor.vertexBuffer = (__bridge id<MTLBuffer>)vb->impl.mtlBuffer;
+	descriptor.vertexStride = vb->impl.myStride;
+	descriptor.triangleCount = ib->impl.count / 3;
+
+	MTLPrimitiveAccelerationStructureDescriptor *accel_descriptor = [MTLPrimitiveAccelerationStructureDescriptor descriptor];
+	accel_descriptor.geometryDescriptors = @[ descriptor ];
+	id<MTLAccelerationStructure> acceleration_structure = create_acceleration_sctructure(accel_descriptor);
+	_primitive_accels = [[NSMutableArray alloc] init];
+	[_primitive_accels addObject:acceleration_structure];
+
+	id<MTLDevice> device = getMetalDevice();
+	id<MTLBuffer> instance_buffer = [device newBufferWithLength:sizeof(MTLAccelerationStructureInstanceDescriptor) * 1 options:options];
+
+	MTLAccelerationStructureInstanceDescriptor *instance_descriptors = (MTLAccelerationStructureInstanceDescriptor *)instance_buffer.contents;
+	instance_descriptors[0].accelerationStructureIndex = 0;
+	instance_descriptors[0].options = MTLAccelerationStructureInstanceOptionOpaque;
+	instance_descriptors[0].mask = 1;
+	instance_descriptors[0].transformationMatrix.columns[0] = MTLPackedFloat3Make(1, 0, 0);
+	instance_descriptors[0].transformationMatrix.columns[1] = MTLPackedFloat3Make(0, 1, 0);
+	instance_descriptors[0].transformationMatrix.columns[2] = MTLPackedFloat3Make(0, 0, 1);
+	instance_descriptors[0].transformationMatrix.columns[3] = MTLPackedFloat3Make(0, 0, 0);
+
+#if !TARGET_OS_IPHONE
+	[instance_buffer didModifyRange:NSMakeRange(0, instance_buffer.length)];
+#endif
+
+	MTLInstanceAccelerationStructureDescriptor *inst_accel_descriptor = [MTLInstanceAccelerationStructureDescriptor descriptor];
+	inst_accel_descriptor.instancedAccelerationStructures = _primitive_accels;
+	inst_accel_descriptor.instanceCount = 1;
+	inst_accel_descriptor.instanceDescriptorBuffer = instance_buffer;
+	_instance_accel = create_acceleration_sctructure(inst_accel_descriptor);
+}
+
+void kinc_raytrace_acceleration_structure_destroy(kinc_raytrace_acceleration_structure_t *accel) {}
+
+void kinc_raytrace_set_acceleration_structure(kinc_raytrace_acceleration_structure_t *_accel) {
+	accel = _accel;
+}
+
+void kinc_raytrace_set_pipeline(kinc_raytrace_pipeline_t *_pipeline) {
+	pipeline = _pipeline;
+}
+
+void kinc_raytrace_set_target(kinc_g5_texture_t *_output) {
+	output = _output;
+}
+
+void kinc_raytrace_dispatch_rays(kinc_g5_command_list_t *command_list) {
+	dispatch_semaphore_wait(_sem, DISPATCH_TIME_FOREVER);
+
+	id<MTLCommandQueue> queue = getMetalQueue();
+	id<MTLCommandBuffer> command_buffer = [queue commandBuffer];
+	__block dispatch_semaphore_t sem = _sem;
+	[command_buffer addCompletedHandler:^(id<MTLCommandBuffer> buffer) {
+		dispatch_semaphore_signal(sem);
+	}];
+
+	NSUInteger width = output->texWidth;
+	NSUInteger height = output->texHeight;
+	MTLSize threads_per_threadgroup = MTLSizeMake(8, 8, 1);
+	MTLSize threadgroups = MTLSizeMake((width + threads_per_threadgroup.width - 1) / threads_per_threadgroup.width,
+	                                   (height + threads_per_threadgroup.height - 1) / threads_per_threadgroup.height, 1);
+
+	id<MTLComputeCommandEncoder> compute_encoder = [command_buffer computeCommandEncoder];
+	[compute_encoder setBuffer:(__bridge id<MTLBuffer>)constant_buf->impl._buffer offset:0 atIndex:0];
+	[compute_encoder setAccelerationStructure:_instance_accel atBufferIndex:1];
+	[compute_encoder setTexture:(__bridge id<MTLTexture>)output->impl._tex atIndex:0];
+
+	for (id<MTLAccelerationStructure> primitive_accel in _primitive_accels)
+		[compute_encoder useResource:primitive_accel usage:MTLResourceUsageRead];
+
+	[compute_encoder setComputePipelineState:_raytracing_pipeline];
+	[compute_encoder dispatchThreadgroups:threadgroups threadsPerThreadgroup:threads_per_threadgroup];
+	[compute_encoder endEncoding];
+	[command_buffer commit];
+}
+
+void kinc_raytrace_copy(kinc_g5_command_list_t *command_list, kinc_g5_render_target_t *target, kinc_g5_texture_t *source) {
+	id<MTLCommandQueue> queue = getMetalQueue();
+	id<MTLCommandBuffer> command_buffer = [queue commandBuffer];
+	id<MTLBlitCommandEncoder> command_encoder = [command_buffer blitCommandEncoder];
+	[command_encoder copyFromTexture:(__bridge id<MTLTexture>)source->impl._tex toTexture:(__bridge id<MTLTexture>)target->impl._tex];
+#ifndef KINC_APPLE_SOC
+	[command_encoder synchronizeResource:(__bridge id<MTLTexture>)target->impl._tex];
+#endif
+	[command_encoder endEncoding];
+	[command_buffer commit];
+	[command_buffer waitUntilCompleted];
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/rendertarget.h

@@ -0,0 +1,7 @@
+#pragma once
+
+typedef struct {
+	void *_tex;
+	void *_texReadback;
+	void *_depthTex;
+} RenderTarget5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/rendertarget.m.h

@@ -0,0 +1,164 @@
+#include <kinc/graphics5/rendertarget.h>
+
+#include <kinc/graphics5/graphics.h>
+#include <kinc/graphics5/rendertarget.h>
+
+#import <Metal/Metal.h>
+
+id getMetalDevice(void);
+id getMetalEncoder(void);
+
+static MTLPixelFormat convert_format(kinc_g5_render_target_format_t format) {
+	switch (format) {
+	case KINC_G5_RENDER_TARGET_FORMAT_128BIT_FLOAT:
+		return MTLPixelFormatRGBA32Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_64BIT_FLOAT:
+		return MTLPixelFormatRGBA16Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_32BIT_RED_FLOAT:
+		return MTLPixelFormatR32Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_16BIT_RED_FLOAT:
+		return MTLPixelFormatR16Float;
+	case KINC_G5_RENDER_TARGET_FORMAT_8BIT_RED:
+		return MTLPixelFormatR8Unorm;
+	case KINC_G5_RENDER_TARGET_FORMAT_32BIT:
+	default:
+		return MTLPixelFormatBGRA8Unorm;
+	}
+}
+
+static void render_target_init(kinc_g5_render_target_t *target, int width, int height, kinc_g5_render_target_format_t format, int depthBufferBits,
+                               int stencilBufferBits, int samples_per_pixel, int framebuffer_index) {
+	memset(target, 0, sizeof(kinc_g5_render_target_t));
+
+	target->texWidth = width;
+	target->texHeight = height;
+
+	target->framebuffer_index = framebuffer_index;
+
+	id<MTLDevice> device = getMetalDevice();
+
+	MTLTextureDescriptor *descriptor = [MTLTextureDescriptor new];
+	descriptor.textureType = MTLTextureType2D;
+	descriptor.width = width;
+	descriptor.height = height;
+	descriptor.depth = 1;
+	descriptor.pixelFormat = convert_format(format);
+	descriptor.arrayLength = 1;
+	descriptor.mipmapLevelCount = 1;
+	descriptor.usage = MTLTextureUsageRenderTarget | MTLTextureUsageShaderRead;
+	descriptor.resourceOptions = MTLResourceStorageModePrivate;
+
+	target->impl._tex = (__bridge_retained void *)[device newTextureWithDescriptor:descriptor];
+
+	if (depthBufferBits > 0) {
+		MTLTextureDescriptor *depthDescriptor = [MTLTextureDescriptor new];
+		depthDescriptor.textureType = MTLTextureType2D;
+		depthDescriptor.width = width;
+		depthDescriptor.height = height;
+		depthDescriptor.depth = 1;
+		depthDescriptor.pixelFormat = MTLPixelFormatDepth32Float_Stencil8;
+		depthDescriptor.arrayLength = 1;
+		depthDescriptor.mipmapLevelCount = 1;
+		depthDescriptor.usage = MTLTextureUsageRenderTarget | MTLTextureUsageShaderRead;
+		depthDescriptor.resourceOptions = MTLResourceStorageModePrivate;
+
+		target->impl._depthTex = (__bridge_retained void *)[device newTextureWithDescriptor:depthDescriptor];
+	}
+
+	target->impl._texReadback = NULL;
+}
+
+void kinc_g5_render_target_init_with_multisampling(kinc_g5_render_target_t *target, int width, int height, kinc_g5_render_target_format_t format,
+                                                   int depthBufferBits, int stencilBufferBits, int samples_per_pixel) {
+	render_target_init(target, width, height, format, depthBufferBits, stencilBufferBits, samples_per_pixel, -1);
+}
+
+static int framebuffer_count = 0;
+
+void kinc_g5_render_target_init_framebuffer_with_multisampling(kinc_g5_render_target_t *target, int width, int height, kinc_g5_render_target_format_t format,
+                                                               int depthBufferBits, int stencilBufferBits, int samples_per_pixel) {
+	render_target_init(target, width, height, format, depthBufferBits, stencilBufferBits, samples_per_pixel, framebuffer_count);
+	framebuffer_count += 1;
+}
+
+void kinc_g5_render_target_init_cube_with_multisampling(kinc_g5_render_target_t *target, int cubeMapSize, kinc_g5_render_target_format_t format,
+                                                        int depthBufferBits, int stencilBufferBits, int samples_per_pixel) {
+	target->impl._tex = NULL;
+	target->impl._depthTex = NULL;
+	target->impl._texReadback = NULL;
+}
+
+void kinc_g5_render_target_destroy(kinc_g5_render_target_t *target) {
+	id<MTLTexture> tex = (__bridge_transfer id<MTLTexture>)target->impl._tex;
+	tex = nil;
+	target->impl._tex = NULL;
+
+	id<MTLTexture> depthTex = (__bridge_transfer id<MTLTexture>)target->impl._depthTex;
+	depthTex = nil;
+	target->impl._depthTex = NULL;
+
+	id<MTLTexture> texReadback = (__bridge_transfer id<MTLTexture>)target->impl._texReadback;
+	texReadback = nil;
+	target->impl._texReadback = NULL;
+
+	if (target->framebuffer_index >= 0) {
+		framebuffer_count -= 1;
+	}
+}
+
+#if 0
+void kinc_g5_set_render_target_descriptor(kinc_g5_render_target_t *renderTarget, kinc_g5_texture_descriptor_t descriptor) {
+    MTLSamplerDescriptor* desc = (MTLSamplerDescriptor*) renderTarget->impl._samplerDesc;
+    switch(descriptor.filter_minification) {
+        case KINC_G5_TEXTURE_FILTER_POINT:
+            desc.minFilter = MTLSamplerMinMagFilterNearest;
+            break;
+        default:
+            desc.minFilter = MTLSamplerMinMagFilterLinear;
+    }
+
+    switch(descriptor.filter_magnification) {
+        case KINC_G5_TEXTURE_FILTER_POINT:
+            desc.magFilter = MTLSamplerMinMagFilterNearest;
+            break;
+        default:
+            desc.minFilter = MTLSamplerMinMagFilterLinear;
+    }
+
+    switch(descriptor.addressing_u) {
+        case KINC_G5_TEXTURE_ADDRESSING_REPEAT:
+            desc.sAddressMode = MTLSamplerAddressModeRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_MIRROR:
+            desc.sAddressMode = MTLSamplerAddressModeMirrorRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_CLAMP:
+            desc.sAddressMode = MTLSamplerAddressModeClampToEdge;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_BORDER:
+            desc.sAddressMode = MTLSamplerAddressModeClampToBorderColor;
+            break;
+    }
+
+    switch(descriptor.addressing_v) {
+        case KINC_G5_TEXTURE_ADDRESSING_REPEAT:
+            desc.tAddressMode = MTLSamplerAddressModeRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_MIRROR:
+            desc.tAddressMode = MTLSamplerAddressModeMirrorRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_CLAMP:
+            desc.tAddressMode = MTLSamplerAddressModeClampToEdge;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_BORDER:
+            desc.tAddressMode = MTLSamplerAddressModeClampToBorderColor;
+            break;
+    }
+    id<MTLDevice> device = getMetalDevice();
+    renderTarget->impl._sampler = [device newSamplerStateWithDescriptor:desc];
+}
+#endif
+
+void kinc_g5_render_target_set_depth_stencil_from(kinc_g5_render_target_t *target, kinc_g5_render_target_t *source) {
+	target->impl._depthTex = source->impl._depthTex;
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/sampler.h

@@ -0,0 +1,5 @@
+#pragma once
+
+typedef struct kinc_g5_sampler_impl {
+	void *sampler;
+} kinc_g5_sampler_impl_t;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/sampler.m.h

@@ -0,0 +1,67 @@
+#include <kinc/graphics5/sampler.h>
+
+static MTLSamplerAddressMode convert_addressing(kinc_g5_texture_addressing_t mode) {
+	switch (mode) {
+	case KINC_G5_TEXTURE_ADDRESSING_REPEAT:
+		return MTLSamplerAddressModeRepeat;
+	case KINC_G5_TEXTURE_ADDRESSING_BORDER:
+		return MTLSamplerAddressModeClampToBorderColor;
+	case KINC_G5_TEXTURE_ADDRESSING_CLAMP:
+		return MTLSamplerAddressModeClampToEdge;
+	case KINC_G5_TEXTURE_ADDRESSING_MIRROR:
+		return MTLSamplerAddressModeMirrorRepeat;
+	default:
+		assert(false);
+		return MTLSamplerAddressModeRepeat;
+	}
+}
+
+static MTLSamplerMipFilter convert_mipmap_mode(kinc_g5_mipmap_filter_t filter) {
+	switch (filter) {
+	case KINC_G5_MIPMAP_FILTER_NONE:
+		return MTLSamplerMipFilterNotMipmapped;
+	case KINC_G5_MIPMAP_FILTER_POINT:
+		return MTLSamplerMipFilterNearest;
+	case KINC_G5_MIPMAP_FILTER_LINEAR:
+		return MTLSamplerMipFilterLinear;
+	default:
+		assert(false);
+		return MTLSamplerMipFilterNearest;
+	}
+}
+
+static MTLSamplerMinMagFilter convert_texture_filter(kinc_g5_texture_filter_t filter) {
+	switch (filter) {
+	case KINC_G5_TEXTURE_FILTER_POINT:
+		return MTLSamplerMinMagFilterNearest;
+	case KINC_G5_TEXTURE_FILTER_LINEAR:
+		return MTLSamplerMinMagFilterLinear;
+	case KINC_G5_TEXTURE_FILTER_ANISOTROPIC:
+		return MTLSamplerMinMagFilterLinear; // ?
+	default:
+		assert(false);
+		return MTLSamplerMinMagFilterNearest;
+	}
+}
+
+void kinc_g5_sampler_init(kinc_g5_sampler_t *sampler, const kinc_g5_sampler_options_t *options) {
+	id<MTLDevice> device = getMetalDevice();
+
+	MTLSamplerDescriptor *desc = (MTLSamplerDescriptor *)[[MTLSamplerDescriptor alloc] init];
+	desc.minFilter = convert_texture_filter(options->minification_filter);
+	desc.magFilter = convert_texture_filter(options->magnification_filter);
+	desc.sAddressMode = convert_addressing(options->u_addressing);
+	desc.tAddressMode = convert_addressing(options->v_addressing);
+	desc.mipFilter = convert_mipmap_mode(options->mipmap_filter);
+	desc.maxAnisotropy = options->max_anisotropy;
+	desc.normalizedCoordinates = YES;
+	desc.lodMinClamp = options->lod_min_clamp;
+	desc.lodMaxClamp = options->lod_max_clamp;
+
+	sampler->impl.sampler = (__bridge_retained void *)[device newSamplerStateWithDescriptor:desc];
+}
+
+void kinc_g5_sampler_destroy(kinc_g5_sampler_t *sampler) {
+	id<MTLSamplerState> mtl_sampler = (__bridge_transfer id<MTLSamplerState>)sampler->impl.sampler;
+	mtl_sampler = nil;
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/shader.h

@@ -0,0 +1,6 @@
+#pragma once
+
+typedef struct {
+	char name[1024];
+	void *mtlFunction;
+} Shader5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/shader.m.h

@@ -0,0 +1,61 @@
+#include <kinc/graphics5/graphics.h>
+#include <kinc/graphics5/shader.h>
+#include <kinc/log.h>
+#include <kinc/math/core.h>
+
+#include <Metal/Metal.h>
+
+id getMetalDevice(void);
+id getMetalLibrary(void);
+
+void kinc_g5_shader_destroy(kinc_g5_shader_t *shader) {
+	id<MTLFunction> function = (__bridge_transfer id<MTLFunction>)shader->impl.mtlFunction;
+	function = nil;
+	shader->impl.mtlFunction = NULL;
+}
+
+void kinc_g5_shader_init(kinc_g5_shader_t *shader, const void *source, size_t length, kinc_g5_shader_type_t type) {
+#ifdef KINC_KONG
+	strcpy(shader->impl.name, (const char *)source);
+	shader->impl.mtlFunction = (__bridge_retained void *)[getMetalLibrary() newFunctionWithName:[NSString stringWithCString:shader->impl.name
+	                                                                                                               encoding:NSUTF8StringEncoding]];
+#else
+	shader->impl.name[0] = 0;
+
+	{
+		uint8_t *data = (uint8_t *)source;
+		if (length > 1 && data[0] == '>') {
+			memcpy(shader->impl.name, data + 1, length - 1);
+			shader->impl.name[length - 1] = 0;
+		}
+		else {
+			for (int i = 3; i < length; ++i) {
+				if (data[i] == '\n') {
+					shader->impl.name[i - 3] = 0;
+					break;
+				}
+				else {
+					shader->impl.name[i - 3] = data[i];
+				}
+			}
+		}
+	}
+
+	char *data = (char *)source;
+	id<MTLLibrary> library = nil;
+	if (length > 1 && data[0] == '>') {
+		library = getMetalLibrary();
+	}
+	else {
+		id<MTLDevice> device = getMetalDevice();
+		NSError *error = nil;
+		library = [device newLibraryWithSource:[[NSString alloc] initWithBytes:data length:length encoding:NSUTF8StringEncoding] options:nil error:&error];
+		if (library == nil) {
+			kinc_log(KINC_LOG_LEVEL_ERROR, "%s", error.localizedDescription.UTF8String);
+		}
+	}
+	shader->impl.mtlFunction = (__bridge_retained void *)[library newFunctionWithName:[NSString stringWithCString:shader->impl.name
+	                                                                                                     encoding:NSUTF8StringEncoding]];
+#endif
+	assert(shader->impl.mtlFunction);
+}

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/texture.h

@@ -0,0 +1,14 @@
+#pragma once
+
+#include <kinc/image.h>
+
+typedef struct {
+	int index;
+	bool vertex;
+} TextureUnit5Impl;
+
+typedef struct {
+	void *_tex;
+	void *data;
+	bool has_mipmaps;
+} Texture5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/texture.m.h

@@ -0,0 +1,291 @@
+#include <kinc/graphics5/texture.h>
+
+#include <kinc/graphics5/graphics.h>
+#include <kinc/graphics5/texture.h>
+#include <kinc/image.h>
+#include <kinc/log.h>
+
+#import <Metal/Metal.h>
+
+id getMetalDevice(void);
+
+static MTLPixelFormat convert_image_format(kinc_image_format_t format) {
+	switch (format) {
+	case KINC_IMAGE_FORMAT_RGBA32:
+		return MTLPixelFormatRGBA8Unorm;
+	case KINC_IMAGE_FORMAT_GREY8:
+		return MTLPixelFormatR8Unorm;
+	case KINC_IMAGE_FORMAT_RGB24:
+		return MTLPixelFormatRGBA8Unorm;
+	case KINC_IMAGE_FORMAT_RGBA128:
+		return MTLPixelFormatRGBA32Float;
+	case KINC_IMAGE_FORMAT_RGBA64:
+		return MTLPixelFormatRGBA16Float;
+	case KINC_IMAGE_FORMAT_A32:
+		return MTLPixelFormatR32Float;
+	case KINC_IMAGE_FORMAT_BGRA32:
+		return MTLPixelFormatBGRA8Unorm;
+	case KINC_IMAGE_FORMAT_A16:
+		return MTLPixelFormatR16Float;
+	}
+}
+
+static int formatByteSize(kinc_image_format_t format) {
+	switch (format) {
+	case KINC_IMAGE_FORMAT_RGBA128:
+		return 16;
+	case KINC_IMAGE_FORMAT_RGBA64:
+		return 8;
+	case KINC_IMAGE_FORMAT_RGB24:
+		return 4;
+	case KINC_IMAGE_FORMAT_A32:
+		return 4;
+	case KINC_IMAGE_FORMAT_A16:
+		return 2;
+	case KINC_IMAGE_FORMAT_GREY8:
+		return 1;
+	case KINC_IMAGE_FORMAT_BGRA32:
+	case KINC_IMAGE_FORMAT_RGBA32:
+		return 4;
+	default:
+		assert(false);
+		return 4;
+	}
+}
+
+static void create(kinc_g5_texture_t *texture, int width, int height, int format, bool writable) {
+	texture->impl.has_mipmaps = false;
+	id<MTLDevice> device = getMetalDevice();
+
+	MTLTextureDescriptor *descriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:convert_image_format((kinc_image_format_t)format)
+	                                                                                      width:width
+	                                                                                     height:height
+	                                                                                  mipmapped:NO];
+	descriptor.textureType = MTLTextureType2D;
+	descriptor.width = width;
+	descriptor.height = height;
+	descriptor.depth = 1;
+	descriptor.pixelFormat = convert_image_format((kinc_image_format_t)format);
+	descriptor.arrayLength = 1;
+	descriptor.mipmapLevelCount = 1;
+	// TODO: Make less textures writable
+	if (writable) {
+		descriptor.usage = MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead;
+	}
+
+	texture->impl._tex = (__bridge_retained void *)[device newTextureWithDescriptor:descriptor];
+}
+
+/*void Graphics5::Texture::_init(const char* format, bool readable) {
+    texWidth = width;
+    texHeight = height;
+
+    create(width, height, Image::RGBA32, false);
+    lock();
+    unlock();
+}*/
+
+void kinc_g5_texture_init(kinc_g5_texture_t *texture, int width, int height, kinc_image_format_t format) {
+	// Image(width, height, format, readable);
+	texture->texWidth = width;
+	texture->texHeight = height;
+	texture->format = format;
+	texture->impl.data = malloc(width * height * (format == KINC_IMAGE_FORMAT_GREY8 ? 1 : 4));
+	create(texture, width, height, format, true);
+}
+
+void kinc_g5_texture_init3d(kinc_g5_texture_t *texture, int width, int height, int depth, kinc_image_format_t format) {}
+
+void kinc_g5_texture_init_from_image(kinc_g5_texture_t *texture, struct kinc_image *image) {
+	texture->texWidth = image->width;
+	texture->texHeight = image->height;
+	texture->format = image->format;
+	texture->impl.data = NULL;
+	create(texture, image->width, image->height, image->format, true);
+	id<MTLTexture> tex = (__bridge id<MTLTexture>)texture->impl._tex;
+	[tex replaceRegion:MTLRegionMake2D(0, 0, texture->texWidth, texture->texHeight)
+	       mipmapLevel:0
+	             slice:0
+	         withBytes:image->data
+	       bytesPerRow:kinc_g5_texture_stride(texture)
+	     bytesPerImage:kinc_g5_texture_stride(texture) * texture->texHeight];
+}
+
+void kinc_g5_texture_init_non_sampled_access(kinc_g5_texture_t *texture, int width, int height, kinc_image_format_t format) {
+	texture->texWidth = width;
+	texture->texHeight = height;
+	texture->format = format;
+	texture->impl.data = malloc(width * height * (format == KINC_IMAGE_FORMAT_GREY8 ? 1 : 4));
+	create(texture, width, height, format, true);
+}
+
+void kinc_g5_texture_destroy(kinc_g5_texture_t *texture) {
+	id<MTLTexture> tex = (__bridge_transfer id<MTLTexture>)texture->impl._tex;
+	tex = nil;
+	texture->impl._tex = NULL;
+
+	if (texture->impl.data != NULL) {
+		free(texture->impl.data);
+		texture->impl.data = NULL;
+	}
+}
+
+id getMetalDevice(void);
+id getMetalEncoder(void);
+
+#if 0
+void kinc_g5_internal_set_texture_descriptor(kinc_g5_texture_t *texture, kinc_g5_texture_descriptor_t descriptor) {
+    MTLSamplerDescriptor* desc = (MTLSamplerDescriptor*) texture->impl._samplerDesc;
+    switch(descriptor.filter_minification) {
+        case KINC_G5_TEXTURE_FILTER_POINT:
+            desc.minFilter = MTLSamplerMinMagFilterNearest;
+            break;
+        default:
+            desc.minFilter = MTLSamplerMinMagFilterLinear;
+    }
+
+    switch(descriptor.filter_magnification) {
+        case KINC_G5_TEXTURE_FILTER_POINT:
+            desc.magFilter = MTLSamplerMinMagFilterNearest;
+            break;
+        default:
+            desc.minFilter = MTLSamplerMinMagFilterLinear;
+    }
+
+    switch(descriptor.addressing_u) {
+        case KINC_G5_TEXTURE_ADDRESSING_REPEAT:
+            desc.sAddressMode = MTLSamplerAddressModeRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_MIRROR:
+            desc.sAddressMode = MTLSamplerAddressModeMirrorRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_CLAMP:
+            desc.sAddressMode = MTLSamplerAddressModeClampToEdge;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_BORDER:
+            desc.sAddressMode = MTLSamplerAddressModeClampToBorderColor;
+            break;
+    }
+
+    switch(descriptor.addressing_v) {
+        case KINC_G5_TEXTURE_ADDRESSING_REPEAT:
+            desc.tAddressMode = MTLSamplerAddressModeRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_MIRROR:
+            desc.tAddressMode = MTLSamplerAddressModeMirrorRepeat;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_CLAMP:
+            desc.tAddressMode = MTLSamplerAddressModeClampToEdge;
+            break;
+        case KINC_G5_TEXTURE_ADDRESSING_BORDER:
+            desc.tAddressMode = MTLSamplerAddressModeClampToBorderColor;
+            break;
+    }
+    id<MTLDevice> device = getMetalDevice();
+    texture->impl._sampler = [device newSamplerStateWithDescriptor:desc];
+}
+#endif
+
+int kinc_g5_texture_stride(kinc_g5_texture_t *texture) {
+	switch (texture->format) {
+	case KINC_IMAGE_FORMAT_GREY8:
+		return texture->texWidth;
+	case KINC_IMAGE_FORMAT_RGBA32:
+	case KINC_IMAGE_FORMAT_BGRA32:
+	case KINC_IMAGE_FORMAT_RGB24:
+		return texture->texWidth * 4;
+	case KINC_IMAGE_FORMAT_RGBA64:
+		return texture->texWidth * 8;
+	case KINC_IMAGE_FORMAT_RGBA128:
+		return texture->texWidth * 16;
+	case KINC_IMAGE_FORMAT_A16:
+		return texture->texWidth * 2;
+	case KINC_IMAGE_FORMAT_A32:
+		return texture->texWidth * 4;
+	}
+}
+
+uint8_t *kinc_g5_texture_lock(kinc_g5_texture_t *texture) {
+	return (uint8_t *)texture->impl.data;
+}
+
+void kinc_g5_texture_unlock(kinc_g5_texture_t *tex) {
+	id<MTLTexture> texture = (__bridge id<MTLTexture>)tex->impl._tex;
+	[texture replaceRegion:MTLRegionMake2D(0, 0, tex->texWidth, tex->texHeight)
+	           mipmapLevel:0
+	                 slice:0
+	             withBytes:tex->impl.data
+	           bytesPerRow:kinc_g5_texture_stride(tex)
+	         bytesPerImage:kinc_g5_texture_stride(tex) * tex->texHeight];
+}
+
+void kinc_g5_texture_clear(kinc_g5_texture_t *texture, int x, int y, int z, int width, int height, int depth, unsigned color) {}
+
+void kinc_g5_texture_generate_mipmaps(kinc_g5_texture_t *texture, int levels) {}
+
+void kinc_g5_texture_set_mipmap(kinc_g5_texture_t *texture, kinc_image_t *mipmap, int level) {
+	if (!texture->impl.has_mipmaps) {
+		id<MTLDevice> device = getMetalDevice();
+		MTLTextureDescriptor *descriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:convert_image_format((kinc_image_format_t)texture->format)
+		                                                                                      width:texture->texWidth
+		                                                                                     height:texture->texHeight
+		                                                                                  mipmapped:YES];
+		descriptor.textureType = MTLTextureType2D;
+		descriptor.width = texture->texWidth;
+		descriptor.height = texture->texHeight;
+		descriptor.depth = 1;
+		descriptor.pixelFormat = convert_image_format((kinc_image_format_t)texture->format);
+		descriptor.arrayLength = 1;
+		bool writable = true;
+		if (writable) {
+			descriptor.usage = MTLTextureUsageShaderWrite | MTLTextureUsageShaderRead;
+		}
+		void *mipmaptex = (__bridge_retained void *)[device newTextureWithDescriptor:descriptor];
+
+		id<MTLCommandQueue> commandQueue = getMetalQueue();
+		id<MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];
+		id<MTLBlitCommandEncoder> commandEncoder = [commandBuffer blitCommandEncoder];
+		[commandEncoder copyFromTexture:(__bridge id<MTLTexture>)texture->impl._tex
+		                    sourceSlice:0
+		                    sourceLevel:0
+		                   sourceOrigin:MTLOriginMake(0, 0, 0)
+		                     sourceSize:MTLSizeMake(texture->texWidth, texture->texHeight, 1)
+		                      toTexture:(__bridge id<MTLTexture>)mipmaptex
+		               destinationSlice:0
+		               destinationLevel:0
+		              destinationOrigin:MTLOriginMake(0, 0, 0)];
+#ifndef KINC_APPLE_SOC
+		[commandEncoder synchronizeResource:(__bridge id<MTLTexture>)mipmaptex];
+#endif
+		[commandEncoder endEncoding];
+		[commandBuffer commit];
+		[commandBuffer waitUntilCompleted];
+
+		id<MTLTexture> tex = (__bridge_transfer id<MTLTexture>)texture->impl._tex;
+		tex = nil;
+		texture->impl._tex = mipmaptex;
+
+		texture->impl.has_mipmaps = true;
+	}
+
+	id<MTLTexture> tex = (__bridge id<MTLTexture>)texture->impl._tex;
+	[tex replaceRegion:MTLRegionMake2D(0, 0, mipmap->width, mipmap->height)
+	       mipmapLevel:level
+	         withBytes:mipmap->data
+	       bytesPerRow:mipmap->width * formatByteSize(mipmap->format)];
+}
+
+#include <kinc/graphics4/texture.h>
+
+#if defined(KINC_IOS) || defined(KINC_MACOS)
+void kinc_g4_texture_upload(kinc_g4_texture_t *texture_g4, uint8_t *data, int stride) {
+	kinc_g5_texture_t *tex = &texture_g4->impl._texture;
+	id<MTLTexture> texture = (__bridge id<MTLTexture>)tex->impl._tex;
+	[texture replaceRegion:MTLRegionMake2D(0, 0, tex->texWidth, tex->texHeight)
+	           mipmapLevel:0
+	                 slice:0
+	             withBytes:data
+	           bytesPerRow:stride
+	         bytesPerImage:stride * tex->texHeight];
+}
+#endif

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/vertexbuffer.h

@@ -0,0 +1,14 @@
+#pragma once
+
+#include <kinc/graphics5/vertexstructure.h>
+
+typedef struct {
+	// void unset();
+	int myCount;
+	int myStride;
+	void *mtlBuffer;
+	bool gpuMemory;
+	int lastStart;
+	int lastCount;
+	// static Graphics5::VertexBuffer* current;
+} VertexBuffer5Impl;

Kha/Kinc/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/vertexbuffer.m.h

@@ -0,0 +1,111 @@
+#include <kinc/graphics5/shader.h>
+#include <kinc/graphics5/vertexbuffer.h>
+
+#include <kinc/graphics5/graphics.h>
+#include <kinc/graphics5/indexbuffer.h>
+#include <kinc/graphics5/vertexbuffer.h>
+
+#import <Metal/Metal.h>
+
+id getMetalDevice(void);
+id getMetalEncoder(void);
+
+kinc_g5_vertex_buffer_t *currentVertexBuffer = NULL;
+
+static void vertex_buffer_unset(kinc_g5_vertex_buffer_t *buffer) {
+	if (currentVertexBuffer == buffer)
+		currentVertexBuffer = NULL;
+}
+
+void kinc_g5_vertex_buffer_init(kinc_g5_vertex_buffer_t *buffer, int count, kinc_g5_vertex_structure_t *structure, bool gpuMemory, int instanceDataStepRate) {
+	memset(&buffer->impl, 0, sizeof(buffer->impl));
+	buffer->impl.myCount = count;
+	buffer->impl.gpuMemory = gpuMemory;
+	for (int i = 0; i < structure->size; ++i) {
+		kinc_g5_vertex_element_t element = structure->elements[i];
+		buffer->impl.myStride += kinc_g4_vertex_data_size(element.data);
+	}
+
+	id<MTLDevice> device = getMetalDevice();
+	MTLResourceOptions options = MTLResourceCPUCacheModeWriteCombined;
+#ifdef KINC_APPLE_SOC
+	options |= MTLResourceStorageModeShared;
+#else
+	if (gpuMemory) {
+		options |= MTLResourceStorageModeManaged;
+	}
+	else {
+		options |= MTLResourceStorageModeShared;
+	}
+#endif
+	id<MTLBuffer> buf = [device newBufferWithLength:count * buffer->impl.myStride options:options];
+	buffer->impl.mtlBuffer = (__bridge_retained void *)buf;
+
+	buffer->impl.lastStart = 0;
+	buffer->impl.lastCount = 0;
+}
+
+void kinc_g5_vertex_buffer_destroy(kinc_g5_vertex_buffer_t *buf) {
+	id<MTLBuffer> buffer = (__bridge_transfer id<MTLBuffer>)buf->impl.mtlBuffer;
+	buffer = nil;
+	buf->impl.mtlBuffer = NULL;
+	vertex_buffer_unset(buf);
+}
+
+float *kinc_g5_vertex_buffer_lock_all(kinc_g5_vertex_buffer_t *buf) {
+	buf->impl.lastStart = 0;
+	buf->impl.lastCount = kinc_g5_vertex_buffer_count(buf);
+	id<MTLBuffer> buffer = (__bridge id<MTLBuffer>)buf->impl.mtlBuffer;
+	float *floats = (float *)[buffer contents];
+	return floats;
+}
+
+float *kinc_g5_vertex_buffer_lock(kinc_g5_vertex_buffer_t *buf, int start, int count) {
+	buf->impl.lastStart = start;
+	buf->impl.lastCount = count;
+	id<MTLBuffer> buffer = (__bridge id<MTLBuffer>)buf->impl.mtlBuffer;
+	float *floats = (float *)[buffer contents];
+	return &floats[start * buf->impl.myStride / sizeof(float)];
+}
+
+void kinc_g5_vertex_buffer_unlock_all(kinc_g5_vertex_buffer_t *buf) {
+#ifndef KINC_APPLE_SOC
+	if (buf->impl.gpuMemory) {
+		id<MTLBuffer> buffer = (__bridge id<MTLBuffer>)buf->impl.mtlBuffer;
+		NSRange range;
+		range.location = buf->impl.lastStart * buf->impl.myStride;
+		range.length = buf->impl.lastCount * buf->impl.myStride;
+		[buffer didModifyRange:range];
+	}
+#endif
+}
+
+void kinc_g5_vertex_buffer_unlock(kinc_g5_vertex_buffer_t *buf, int count) {
+#ifndef KINC_APPLE_SOC
+	if (buf->impl.gpuMemory) {
+		id<MTLBuffer> buffer = (__bridge id<MTLBuffer>)buf->impl.mtlBuffer;
+		NSRange range;
+		range.location = buf->impl.lastStart * buf->impl.myStride;
+		range.length = count * buf->impl.myStride;
+		[buffer didModifyRange:range];
+	}
+#endif
+}
+
+int kinc_g5_internal_vertex_buffer_set(kinc_g5_vertex_buffer_t *buf, int offset_) {
+	currentVertexBuffer = buf;
+
+	id<MTLRenderCommandEncoder> encoder = getMetalEncoder();
+	id<MTLBuffer> buffer = (__bridge id<MTLBuffer>)buf->impl.mtlBuffer;
+	[encoder setVertexBuffer:buffer offset:offset_ * buf->impl.myStride atIndex:0];
+
+	return offset_;
+}
+
+int kinc_g5_vertex_buffer_count(kinc_g5_vertex_buffer_t *buffer) {
+	return buffer->impl.myCount;
+}
+
+int kinc_g5_vertex_buffer_stride(kinc_g5_vertex_buffer_t *buffer) {
+	return buffer->impl.myStride;
+}