[Metal] Make compute work

Kode · Jan 17, 2024 · 7caa822 · 7caa822
1 parent 6e50435
commit 7caa822
Show file tree

Hide file tree

Showing 5 changed files with 14 additions and 6 deletions.
diff --git a/Backends/Graphics4/G4onG5/Sources/kinc/backend/graphics4/G4.c.h b/Backends/Graphics4/G4onG5/Sources/kinc/backend/graphics4/G4.c.h
@@ -206,7 +206,10 @@ static void endDraw(bool compute) {
 			kinc_g5_command_list_set_pipeline(&commandList, current_state.pipeline);
 		}
 		if (current_state.compute_shader != NULL) {
+#ifndef KORE_METAL
+			// Metal still has some trouble switching between graphics and compute encoders
 			kinc_g5_command_list_set_compute_shader(&commandList, current_state.compute_shader);
+#endif
 		}
 		if (current_state.index_buffer != NULL) {
 			kinc_g5_command_list_set_index_buffer(&commandList, current_state.index_buffer);

diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/Metal.m.h b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/Metal.m.h
@@ -81,6 +81,7 @@ static void start_render_pass(void) {
 
 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) {

diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/commandlist.m.h b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/commandlist.m.h
@@ -281,7 +281,7 @@ void kinc_g5_command_list_set_fragment_constant_buffer(kinc_g5_command_list_t *l
 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];
+	[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) {
@@ -308,7 +308,9 @@ void kinc_g5_command_list_set_texture(kinc_g5_command_list_t *list, kinc_g5_text
 	}
 }
 
-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) {}
+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;
@@ -384,6 +386,8 @@ void kinc_g5_command_list_compute(kinc_g5_command_list_t *list, int x, int y, in
 	[compute_command_encoder dispatchThreadgroups:perGrid threadsPerThreadgroup:perGroup];
 
 	[compute_command_encoder endEncoding];
+
+	compute_command_encoder = nil;
 
 	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
@@ -39,7 +39,7 @@ void kinc_g5_compute_shader_init(kinc_g5_compute_shader *shader, void *_data, in
 		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(shader->impl._function != nil);
+	assert(function != nil);
 	shader->impl._function = (__bridge_retained void *)function;
 
 	id<MTLDevice> device = getMetalDevice();

diff --git a/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/metalunit.m b/Backends/Graphics5/Metal/Sources/kinc/backend/graphics5/metalunit.m
@@ -1,9 +1,9 @@
 #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 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);