[Metal] Make compute things compile again

Kode · Jan 17, 2024 · 6e50435 · 6e50435
1 parent 97b08b8
commit 6e50435
Show file tree

Hide file tree

Showing 8 changed files with 121 additions and 101 deletions.
diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/Metal.m.h b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/Metal.m.h
@@ -20,16 +20,14 @@ int newRenderTargetWidth;
 int newRenderTargetHeight;
 
 id<CAMetalDrawable> drawable;
-id<MTLCommandBuffer> commandBuffer;
-id<MTLRenderCommandEncoder> commandEncoder;
 id<MTLTexture> depthTexture;
 int depthBits;
 int stencilBits;
 
 static kinc_g5_render_target_t fallback_render_target;
 
 id getMetalEncoder(void) {
-	return commandEncoder;
+	return render_command_encoder;
 }
 
 void kinc_g5_internal_destroy_window(int window) {}
@@ -62,6 +60,29 @@ 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];
+}
+
 void kinc_g5_begin(kinc_g5_render_target_t *renderTarget, int window) {
 	CAMetalLayer *metalLayer = getMetalLayer();
 	drawable = [metalLayer nextDrawable];
@@ -100,27 +121,27 @@ void kinc_g5_begin(kinc_g5_render_target_t *renderTarget, int window) {
 	renderPassDescriptor.stencilAttachment.storeAction = MTLStoreActionDontCare;
 	renderPassDescriptor.stencilAttachment.texture = depthTexture;
 
-	if (commandBuffer != nil && commandEncoder != nil) {
-		[commandEncoder endEncoding];
-		[commandBuffer commit];
+	if (command_buffer != nil && render_command_encoder != nil) {
+		[render_command_encoder endEncoding];
+		[command_buffer commit];
 	}
 
 	id<MTLCommandQueue> commandQueue = getMetalQueue();
-	commandBuffer = [commandQueue commandBuffer];
-	commandEncoder = [commandBuffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
+	command_buffer = [commandQueue commandBuffer];
+	render_command_encoder = [command_buffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
 }
 
 void kinc_g5_end(int window) {}
 
 bool kinc_g5_swap_buffers(void) {
-	if (commandBuffer != nil && commandEncoder != nil) {
-		[commandEncoder endEncoding];
-		[commandBuffer presentDrawable:drawable];
-		[commandBuffer commit];
+	if (command_buffer != nil && render_command_encoder != nil) {
+		[render_command_encoder endEncoding];
+		[command_buffer presentDrawable:drawable];
+		[command_buffer commit];
 	}
 	drawable = nil;
-	commandBuffer = nil;
-	commandEncoder = nil;
+	command_buffer = nil;
+	render_command_encoder = nil;
 
 	return true;
 }
@@ -131,11 +152,11 @@ bool kinc_window_vsynced(int window) {
 
 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 (commandBuffer != nil && commandEncoder != nil) {
-		[commandEncoder endEncoding];
-		[commandBuffer commit];
+	if (command_buffer != nil && render_command_encoder != nil) {
+		[render_command_encoder endEncoding];
+		[command_buffer commit];
 		if (wait) {
-			[commandBuffer waitUntilCompleted];
+			[command_buffer waitUntilCompleted];
 		}
 	}
 
@@ -198,8 +219,8 @@ void kinc_g5_internal_new_render_pass(kinc_g5_render_target_t **renderTargets, i
 	}
 
 	id<MTLCommandQueue> commandQueue = getMetalQueue();
-	commandBuffer = [commandQueue commandBuffer];
-	commandEncoder = [commandBuffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
+	command_buffer = [commandQueue commandBuffer];
+	render_command_encoder = [command_buffer renderCommandEncoderWithDescriptor:renderPassDescriptor];
 }
 
 bool kinc_g5_supports_raytracing(void) {

diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/commandlist.m.h b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/commandlist.m.h
@@ -1,4 +1,5 @@
 #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>
@@ -277,6 +278,12 @@ void kinc_g5_command_list_set_fragment_constant_buffer(kinc_g5_command_list_t *l
 	[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:0];
+}
+
 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();
@@ -287,13 +294,17 @@ void kinc_g5_command_list_render_target_to_texture_barrier(kinc_g5_command_list_
 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<MTLRenderCommandEncoder> encoder = getMetalEncoder();
 	id<MTLTexture> tex = (__bridge id<MTLTexture>)texture->impl._tex;
-	if (unit.stages[KINC_G5_SHADER_TYPE_VERTEX] >= 0) {
-		[encoder setVertexTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_VERTEX]];
+	if (compute_command_encoder != nil) {
+		[compute_command_encoder setTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_COMPUTE]];
 	}
-	if (unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT] >= 0) {
-		[encoder setFragmentTexture:tex atIndex:unit.stages[KINC_G5_SHADER_TYPE_FRAGMENT]];
+	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]];
+		}
 	}
 }
 
@@ -348,3 +359,31 @@ void kinc_g5_command_list_set_sampler(kinc_g5_command_list_t *list, kinc_g5_text
 		[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];
+
+	start_render_pass();
+}
diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/compute.m.h b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/compute.m.h
@@ -7,40 +7,6 @@
 id getMetalDevice(void);
 id getMetalLibrary(void);
 
-static id<MTLCommandQueue> commandQueue;
-static id<MTLCommandBuffer> commandBuffer;
-static id<MTLComputeCommandEncoder> commandEncoder;
-static id<MTLBuffer> buffer;
-
-void initMetalCompute(id<MTLDevice> device, id<MTLCommandQueue> queue) {
-	commandQueue = queue;
-	commandBuffer = [commandQueue commandBuffer];
-	commandEncoder = [commandBuffer computeCommandEncoder];
-	buffer = [device newBufferWithLength:constantsSize options:MTLResourceOptionCPUCacheModeDefault];
-	constantsMemory = (uint8_t *)[buffer contents];
-}
-
-void shutdownMetalCompute(void) {
-	[commandEncoder endEncoding];
-	commandEncoder = nil;
-	commandBuffer = nil;
-	commandQueue = nil;
-}
-
-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;
-}
-
 void kinc_g5_compute_shader_init(kinc_g5_compute_shader *shader, void *_data, int length) {
 	shader->impl.name[0] = 0;
 
@@ -89,9 +55,25 @@ void kinc_g5_compute_shader_init(kinc_g5_compute_shader *shader, void *_data, in
 	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._offset = -1;
+	location.impl.vertexOffset = -1;
+	location.impl.fragmentOffset = -1;
+	location.impl.computeOffset = -1;
 
 	MTLComputePipelineReflection *reflection = (__bridge MTLComputePipelineReflection *)shader->impl._reflection;
 
@@ -101,7 +83,7 @@ kinc_g5_constant_location_t kinc_g5_compute_shader_get_constant_location(kinc_g5
 				MTLStructType *structObj = [arg bufferStructType];
 				for (MTLStructMember *member in structObj.members) {
 					if (strcmp([[member name] UTF8String], name) == 0) {
-						location.impl._offset = (int)[member offset];
+						location.impl.computeOffset = (int)[member offset];
 						break;
 					}
 				}
@@ -115,40 +97,16 @@ kinc_g5_constant_location_t kinc_g5_compute_shader_get_constant_location(kinc_g5
 
 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;
-	unit.impl._index = -1;
+	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.impl._index = (int)[arg index];
+			unit.stages[KINC_G5_SHADER_TYPE_COMPUTE] = (int)[arg index];
 		}
 	}
 
 	return unit;
 }
-
-void kinc_compute_set_shader(kinc_compute_shader_t *shader) {
-	id<MTLComputePipelineState> pipeline = (__bridge id<MTLComputePipelineState>)shader->impl._pipeline;
-	[commandEncoder setComputePipelineState:pipeline];
-}
-
-void kinc_compute(int x, int y, int z) {
-	[commandEncoder setBuffer:buffer offset:0 atIndex:0];
-
-	MTLSize perGrid;
-	perGrid.width = x;
-	perGrid.height = y;
-	perGrid.depth = z;
-	MTLSize perGroup;
-	perGroup.width = 16;
-	perGroup.height = 16;
-	perGroup.depth = 1;
-	[commandEncoder dispatchThreadgroups:perGrid threadsPerThreadgroup:perGroup];
-
-	[commandEncoder endEncoding];
-	[commandBuffer commit];
-	[commandBuffer waitUntilCompleted];
-
-	commandBuffer = [commandQueue commandBuffer];
-	commandEncoder = [commandBuffer computeCommandEncoder];
-}
diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/metalunit.m b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/metalunit.m
@@ -1,5 +1,16 @@
+#import <Metal/Metal.h>
+#import <MetalKit/MTKView.h>
+
+static id<MTLCommandBuffer> command_buffer;
+static id<MTLRenderCommandEncoder> render_command_encoder;
+static id<MTLComputeCommandEncoder> compute_command_encoder;
+
+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"

diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/pipeline.h b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/pipeline.h
@@ -20,4 +20,5 @@ typedef struct {
 typedef struct {
 	int vertexOffset;
 	int fragmentOffset;
+	int computeOffset;
 } ConstantLocation5Impl;
diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/pipeline.m.h b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/pipeline.m.h
@@ -371,7 +371,8 @@ kinc_g5_constant_location_t kinc_g5_pipeline_get_constant_location(kinc_g5_pipel
 	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) {

diff --git a/Backends/System/macOS/Sources/kinc/backend/BasicOpenGLView.m.h b/Backends/System/macOS/Sources/kinc/backend/BasicOpenGLView.m.h
@@ -416,15 +416,12 @@ static bool controlKeyMouseButton = false;
 }
 #else
 
-void initMetalCompute(id<MTLDevice> device, id<MTLCommandQueue> commandBuffer);
-
 - (id)initWithFrame:(NSRect)frameRect {
 	self = [super initWithFrame:frameRect];
 
 	device = MTLCreateSystemDefaultDevice();
 	commandQueue = [device newCommandQueue];
 	library = [device newDefaultLibrary];
-	initMetalCompute(device, commandQueue);
 
 	CAMetalLayer *metalLayer = (CAMetalLayer *)self.layer;
 

diff --git a/Sources/kinc/system.h b/Sources/kinc/system.h
@@ -562,10 +562,6 @@ void kinc_set_application_name(const char *name) {
 	strcpy(application_name, name);
 }
 
-#ifdef KORE_METAL
-void shutdownMetalCompute(void);
-#endif
-
 void kinc_stop(void) {
 	running = false;
 
@@ -574,10 +570,6 @@ void kinc_stop(void) {
 	// for (int windowIndex = 0; windowIndex < sizeof(windowIds) / sizeof(int); ++windowIndex) {
 	//	Graphics::destroy(windowIndex);
 	//}
-
-#ifdef KORE_METAL
-	shutdownMetalCompute();
-#endif
 }
 
 bool kinc_internal_frame(void) {