base: handle invalid shaders

base/sources/kong/kong.c

@@ -1,6 +1,7 @@
 #include "kong.h"
 #include "dir.h"
 #include <assert.h>
+#include <setjmp.h>
 #include <stdint.h>
 #include <stdio.h>
 #include <stdlib.h>
@@ -86,6 +87,8 @@ size_t                  allocated_globals_size = 0;
 uint64_t next_variable_id = 1;
 
 bool               kong_error          = false;
+static jmp_buf     kong_error_jmp;
+static bool        kong_error_jmp_active = false;
 static function   *functions           = NULL;
 static function_id functions_size      = 128;
 static function_id functions_zeroed    = 0;
@@ -352,12 +355,12 @@ void find_used_capabilities(function *f) {
 			type_id  to_type = to.type.type;
 
 			if (is_texture(to_type)) {
-				assert(get_type(to_type)->array_size == 0);
+				kong_assert(get_type(to_type)->array_size == 0);
 
 				f->used_capabilities.image_write = true;
 
 				global *g = find_global_by_var(to);
-				assert(g != NULL);
+				kong_assert(g != NULL);
 			}
 			break;
 		}
@@ -374,7 +377,7 @@ void find_used_capabilities(function *f) {
 				}
 				else {
 					global *g = find_global_by_var(from);
-					assert(g != NULL);
+					kong_assert(g != NULL);
 				}
 			}
 
@@ -389,14 +392,14 @@ void find_used_capabilities(function *f) {
 				global *g = NULL;
 
 				if (tex_parameter.kind == VARIABLE_INTERNAL) {
-					assert(last_base_texture_to.index == tex_parameter.index);
+					kong_assert(last_base_texture_to.index == tex_parameter.index);
 					g = find_global_by_var(last_base_texture_from);
 				}
 				else {
 					g = find_global_by_var(tex_parameter);
 				}
 
-				assert(g != NULL);
+				kong_assert(g != NULL);
 			}
 
 			for (function_id i = 0; get_function(i) != NULL; ++i) {
@@ -684,12 +687,12 @@ static void update_globals_in_descriptor_set_group(descriptor_set_group *group,
 }
 
 descriptor_set_group *get_descriptor_set_group(uint32_t descriptor_set_group_index) {
-	assert(descriptor_set_group_index < all_descriptor_set_groups.size);
+	kong_assert(descriptor_set_group_index < all_descriptor_set_groups.size);
 	return &all_descriptor_set_groups.values[descriptor_set_group_index];
 }
 
 static void assign_descriptor_set_group_index(function *f, uint32_t descriptor_set_group_index) {
-	assert(f->descriptor_set_group_index == UINT32_MAX || f->descriptor_set_group_index == descriptor_set_group_index);
+	kong_assert(f->descriptor_set_group_index == UINT32_MAX || f->descriptor_set_group_index == descriptor_set_group_index);
 	f->descriptor_set_group_index = descriptor_set_group_index;
 }
 
@@ -864,7 +867,7 @@ variable allocate_variable(type_ref type, variable_kind kind) {
 }
 
 opcode *emit_op(opcodes *code, opcode *o) {
-	assert(code->size + o->size < OPCODES_SIZE);
+	kong_assert(code->size + o->size < OPCODES_SIZE);
 
 	if (code->o == NULL) {
 		code->o = (uint8_t *)malloc(OPCODES_SIZE);
@@ -1097,7 +1100,7 @@ variable emit_expression(opcodes *code, block *parent, expression *e) {
 
 						break;
 					default:
-						assert(false);
+						kong_assert(false);
 						break;
 					}
 
@@ -1308,7 +1311,7 @@ variable emit_expression(opcodes *code, block *parent, expression *e) {
 
 				break;
 			default:
-				assert(false);
+				kong_assert(false);
 				break;
 			}
 
@@ -1693,11 +1696,22 @@ void error_args(debug_context context, const char *message, va_list args) {
 
 	iron_log_args(IRON_LOG_LEVEL_ERROR, buffer, args);
 	kong_error = true;
+	if (kong_error_jmp_active) {
+		longjmp(kong_error_jmp, 1);
+	}
 }
 
 void error_args_no_context(const char *message, va_list args) {
 	iron_log_args(IRON_LOG_LEVEL_ERROR, message, args);
 	kong_error = true;
+	if (kong_error_jmp_active) {
+		longjmp(kong_error_jmp, 1);
+	}
+}
+
+void kong_assert_failed(const char *test, const char *file, int line) {
+	error_no_context("Shader compiler assertion failed: %s (%s:%i)", test, file, line);
+	assert(false); // Not compiling a shader at runtime (amake)
 }
 
 void error(debug_context context, const char *message, ...) {
@@ -1726,6 +1740,7 @@ void check_function(bool test, debug_context context, const char *message, ...)
 		va_start(args, message);
 		error_args(context, message, args);
 		va_end(args);
+		assert(false); // Not compiling a shader at runtime (amake)
 	}
 }
 
@@ -3138,7 +3153,7 @@ static statement *parse_statement(state_t *state, block *parent_block) {
 			}
 
 			s->iffy.else_size += 1;
-			assert(s->iffy.else_size < 64);
+			kong_assert(s->iffy.else_size < 64);
 		}
 
 		return s;
@@ -4138,7 +4153,7 @@ size_t get_sets_count(void) {
 }
 
 void add_definition_to_set(descriptor_set *set, definition def) {
-	assert(def.kind != DEFINITION_FUNCTION && def.kind != DEFINITION_STRUCT);
+	kong_assert(def.kind != DEFINITION_FUNCTION && def.kind != DEFINITION_STRUCT);
 
 	for (size_t global_index = 0; global_index < set->globals.size; ++global_index) {
 		if (set->globals.globals[global_index] == def.global) {
@@ -4666,7 +4681,7 @@ static void copy_opcode(opcode *o) {
 
 	uint8_t *new_data = &new_code.o[new_code.size];
 
-	assert(new_code.size + o->size < OPCODES_SIZE);
+	kong_assert(new_code.size + o->size < OPCODES_SIZE);
 
 	memcpy(new_data, o, o->size);
 
@@ -4695,7 +4710,7 @@ void transform(uint32_t flags) {
 				kong_access a = o->op_store_access_list.access_list[o->op_store_access_list.access_list_size - 1];
 
 				if ((flags & TRANSFORM_FLAG_ONE_COMPONENT_SWIZZLE) != 0 && a.kind == ACCESS_SWIZZLE && a.access_swizzle.swizzle.size > 1) {
-					assert(is_vector(o->op_store_access_list.from.type.type));
+					kong_assert(is_vector(o->op_store_access_list.from.type.type));
 
 					type_id from_base_type = vector_base_type(o->op_store_access_list.from.type.type);
 
@@ -4752,7 +4767,7 @@ void transform(uint32_t flags) {
 				kong_access a = o->op_load_access_list.access_list[o->op_load_access_list.access_list_size - 1];
 
 				if ((flags & TRANSFORM_FLAG_ONE_COMPONENT_SWIZZLE) != 0 && a.kind == ACCESS_SWIZZLE && a.access_swizzle.swizzle.size > 1) {
-					assert(is_vector(o->op_load_access_list.to.type.type));
+					kong_assert(is_vector(o->op_load_access_list.to.type.type));
 
 					type_id to_type = vector_base_type(o->op_load_access_list.to.type.type);
 
@@ -5012,7 +5027,7 @@ static void resolve_types_in_element(statement *parent_block, expression *elemen
 
 	type_id of_type = element->element.of->type.type;
 
-	assert(of_type != NO_TYPE);
+	kong_assert(of_type != NO_TYPE);
 
 	type *of = get_type(of_type);
 
@@ -5036,7 +5051,7 @@ static void resolve_types_in_member(statement *parent_block, expression *member)
 	type_id of_type     = member->member.of->type.type;
 	name_id member_name = member->member.member_name;
 
-	assert(of_type != NO_TYPE);
+	kong_assert(of_type != NO_TYPE);
 
 	if (is_vector_or_scalar(of_type)) {
 		expression *of = member->member.of;
@@ -5102,7 +5117,7 @@ static void resolve_types_in_member(statement *parent_block, expression *member)
 				member->type.type = float4_id;
 				break;
 			default:
-				assert(false);
+				kong_assert(false);
 				break;
 			}
 		}
@@ -5121,7 +5136,7 @@ static void resolve_types_in_member(statement *parent_block, expression *member)
 				member->type.type = int4_id;
 				break;
 			default:
-				assert(false);
+				kong_assert(false);
 				break;
 			}
 		}
@@ -5140,7 +5155,7 @@ static void resolve_types_in_member(statement *parent_block, expression *member)
 				member->type.type = uint4_id;
 				break;
 			default:
-				assert(false);
+				kong_assert(false);
 				break;
 			}
 		}
@@ -5159,12 +5174,12 @@ static void resolve_types_in_member(statement *parent_block, expression *member)
 				member->type.type = bool4_id;
 				break;
 			default:
-				assert(false);
+				kong_assert(false);
 				break;
 			}
 		}
 		else {
-			assert(false);
+			kong_assert(false);
 		}
 	}
 	else {
@@ -5533,7 +5548,7 @@ void resolve_types_in_expression(statement *parent, expression *e) {
 		if (e->call.func_name == add_name("sample") || e->call.func_name == add_name("sample_lod")) {
 			if (e->call.parameters.e[0]->kind == EXPRESSION_VARIABLE) {
 				global *g = find_global(e->call.parameters.e[0]->variable);
-				assert(g != NULL);
+				kong_assert(g != NULL);
 				e->type.type = float4_id;
 			}
 			else {
@@ -5564,7 +5579,7 @@ void resolve_types_in_expression(statement *parent, expression *e) {
 		break;
 	}
 	case EXPRESSION_SWIZZLE:
-		assert(false); // swizzle is created in the typer
+		kong_assert(false); // swizzle is created in the typer
 		break;
 	}
 
@@ -5575,6 +5590,8 @@ void resolve_types_in_expression(statement *parent, expression *e) {
 	}
 }
 
+static function *resolve_types_function = NULL;
+
 void resolve_types_in_block(statement *parent, statement *block) {
 	debug_context context = {0};
 	check(block->kind == STATEMENT_BLOCK, context, "Malformed block");
@@ -5590,6 +5607,15 @@ void resolve_types_in_block(statement *parent, statement *block) {
 		}
 		case STATEMENT_RETURN_EXPRESSION: {
 			resolve_types_in_expression(block, s->expression);
+			type_id return_type = resolve_types_function->return_type.type;
+			type_id value_type  = s->expression->type.type;
+			bool    mismatch    = !types_compatible(value_type, return_type) ||
+			                      (is_vector_or_scalar(value_type) && is_vector_or_scalar(return_type) && vector_size(value_type) != vector_size(return_type));
+			if (value_type != NO_TYPE && return_type != NO_TYPE && mismatch) {
+				debug_context context = {0};
+				error(context, "Return type mismatch %s vs %s in %s", get_name(get_type(value_type)->name), get_name(get_type(return_type)->name),
+				      get_name(resolve_types_function->name));
+			}
 			break;
 		}
 		case STATEMENT_IF: {
@@ -5692,6 +5718,7 @@ void resolve_types(void) {
 			++f->block->block.vars.size;
 		}
 
+		resolve_types_function = f;
 		resolve_types_in_block(NULL, f->block);
 	}
 }
@@ -5974,7 +6001,7 @@ uint32_t vector_size(type_id t) {
 		return 4u;
 	}
 
-	assert(false);
+	kong_assert(false);
 	return 0;
 }
 
@@ -5992,7 +6019,7 @@ type_id vector_base_type(type_id vector_type) {
 		return bool_id;
 	}
 
-	assert(false);
+	kong_assert(false);
 	return float_id;
 }
 
@@ -6009,7 +6036,7 @@ type_id vector_to_size(type_id vector_type, uint32_t size) {
 		case 4u:
 			return float4_id;
 		default:
-			assert(false);
+			kong_assert(false);
 			return float_id;
 		}
 	}
@@ -6024,7 +6051,7 @@ type_id vector_to_size(type_id vector_type, uint32_t size) {
 		case 4u:
 			return int4_id;
 		default:
-			assert(false);
+			kong_assert(false);
 			return int_id;
 		}
 	}
@@ -6039,7 +6066,7 @@ type_id vector_to_size(type_id vector_type, uint32_t size) {
 		case 4u:
 			return uint4_id;
 		default:
-			assert(false);
+			kong_assert(false);
 			return uint_id;
 		}
 	}
@@ -6054,12 +6081,12 @@ type_id vector_to_size(type_id vector_type, uint32_t size) {
 		case 4u:
 			return bool4_id;
 		default:
-			assert(false);
+			kong_assert(false);
 			return bool_id;
 		}
 	}
 	else {
-		assert(false);
+		kong_assert(false);
 		return float_id;
 	}
 }
@@ -6172,13 +6199,9 @@ void gpu_create_shaders_from_kong(char *kong, char **vs, char **fs, int *vs_size
 		statement_index  = _statement_index;
 	}
 
-	kong_error    = false;
-	char  *from   = "";
-	tokens tokens = tokenize(from, kong);
-	parse(from, &tokens);
-	resolve_types();
-
-	if (kong_error) {
+	kong_error = false;
+	if (setjmp(kong_error_jmp) != 0) { // error() while compiling
+		kong_error_jmp_active = false;
 		console_info("Warning: Shader compilation failed");
 #if defined(__APPLE__)
 		*vs = "";
@@ -6186,6 +6209,12 @@ void gpu_create_shaders_from_kong(char *kong, char **vs, char **fs, int *vs_size
 #endif
 		return;
 	}
+	kong_error_jmp_active = true;
+
+	char  *from   = "";
+	tokens tokens = tokenize(from, kong);
+	parse(from, &tokens);
+	resolve_types();
 	allocate_globals();
 	for (function_id i = 0; get_function(i) != NULL; ++i) {
 		compile_function_block(&get_function(i)->code, get_function(i)->block);
@@ -6217,4 +6246,6 @@ void gpu_create_shaders_from_kong(char *kong, char **vs, char **fs, int *vs_size
 	spirv_export2(vs, fs, vs_size, fs_size, false);
 
 #endif
+
+	kong_error_jmp_active = false;
 }

base/sources/kong/kong.h

@@ -634,9 +634,15 @@ void                  error_args(debug_context context, const char *message, va_
 void                  error_args_no_context(const char *message, va_list args);
 void                  check_function(bool test, debug_context context, const char *message, ...);
 
-#define check(test, context, message, ...) \
-	assert(test);                          \
-	check_function(test, context, message, ##__VA_ARGS__)
+#define check(test, context, message, ...) check_function(test, context, message, ##__VA_ARGS__)
+
+void kong_assert_failed(const char *test, const char *file, int line);
+#define kong_assert(test)                                      \
+	do {                                                       \
+		if (!(test)) {                                         \
+			kong_assert_failed(#test, __FILE__, __LINE__);     \
+		}                                                      \
+	} while (0)
 
 void        check_args(bool test, debug_context context, const char *message, va_list args);
 void        functions_init(void);

base/sources/kong/kong_cstyle.c

@@ -132,7 +132,7 @@ void cstyle_write_opcode(char *code, size_t *offset, opcode *o, type_string_func
 			*offset += sprintf(&code[*offset], " *= _%" PRIu64 ";\n", o->op_store_access_list.from.index);
 			break;
 		default:
-			assert(false);
+			kong_assert(false);
 			break;
 		}
 		break;

base/sources/kong/kong_hlsl.c

@@ -62,7 +62,7 @@ static char *type_string(type_id type) {
 		}
 		else {
 			// TODO
-			assert(false);
+			kong_assert(false);
 		}
 	}
 	return get_name(get_type(type)->name);
@@ -701,7 +701,7 @@ static void write_functions(char *hlsl, size_t *offset, shader_stage stage, func
 
 	for (size_t i = 0; i < functions_size; ++i) {
 		function *f = functions[i];
-		assert(f != NULL);
+		kong_assert(f != NULL);
 
 		debug_context context = {0};
 		check(f->block != NULL, context, "Function block missing");
@@ -973,7 +973,7 @@ static void write_functions(char *hlsl, size_t *offset, shader_stage stage, func
 					*offset += sprintf(&hlsl[*offset], " *= _%" PRIu64 ";\n", o->op_store_access_list.from.index);
 					break;
 				default:
-					assert(false);
+					kong_assert(false);
 					break;
 				}
 				break;
@@ -1240,7 +1240,7 @@ static char *hlsl_export_vertex2(api_kind d3d, function *main, bool debug) {
 	char       *hlsl   = &_buffer[0];
 	size_t      offset = 0;
 
-	assert(main->parameters_size > 0);
+	kong_assert(main->parameters_size > 0);
 	type_id vertex_inputs[64];
 	for (size_t input_index = 0; input_index < main->parameters_size; ++input_index) {
 		vertex_inputs[input_index] = main->parameter_types[input_index].type;
@@ -1262,7 +1262,7 @@ static char *hlsl_export_fragment2(api_kind d3d, function *main, bool debug) {
 	char       *hlsl   = &_buffer[0];
 	size_t      offset = 0;
 
-	assert(main->parameters_size > 0);
+	kong_assert(main->parameters_size > 0);
 	type_id pixel_input = main->parameter_types[0].type;
 
 	debug_context context = {0};

base/sources/kong/kong_metal.c

@@ -574,7 +574,7 @@ static void write_functions(char *code, size_t *offset) {
 						*offset += sprintf(&code[*offset], " *= _%" PRIu64 ";\n", o->op_store_access_list.from.index);
 						break;
 					default:
-						assert(false);
+						kong_assert(false);
 						break;
 					}
 				}
@@ -900,7 +900,7 @@ char *metal_export(char *directory) {
 					fragment_functions[fragment_functions_size] = i;
 					fragment_functions_size += 1;
 
-					assert(f->parameters_size > 0);
+					kong_assert(f->parameters_size > 0);
 					fragment_inputs[fragment_inputs_size] = f->parameter_types[0].type;
 					fragment_inputs_size += 1;
 				}

base/sources/kong/kong_spirv.c

@@ -334,7 +334,7 @@ static void write_buffer(FILE *file, uint8_t *output, size_t output_size) {
 
 static type_id find_access_type(int *indices, access_kind *access_kinds, int indices_size, type_id base_type) {
 	if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D) {
-		assert(indices_size == 1);
+		kong_assert(indices_size == 1);
 		return float4_id;
 	}
 
@@ -348,12 +348,12 @@ static type_id find_access_type(int *indices, access_kind *access_kinds, int ind
 			return vector_base_type(base_type);
 		case ACCESS_MEMBER: {
 			type *t = get_type(base_type);
-			assert(indices[0] < t->members.size);
+			kong_assert(indices[0] < t->members.size);
 			return t->members.m[indices[0]].type.type;
 		}
 		}
 
-		assert(false);
+		kong_assert(false);
 		return base_type;
 	}
 	else {
@@ -366,12 +366,12 @@ static type_id find_access_type(int *indices, access_kind *access_kinds, int ind
 			return base_type;
 		case ACCESS_MEMBER: {
 			type *t = get_type(base_type);
-			assert(indices[0] < t->members.size);
+			kong_assert(indices[0] < t->members.size);
 			return find_access_type(&indices[1], &access_kinds[1], indices_size - 1, t->members.m[indices[0]].type.type);
 		}
 		}
 
-		assert(false);
+		kong_assert(false);
 		return base_type;
 	}
 }
@@ -644,7 +644,7 @@ static struct {
 } *type_map = NULL;
 
 static void add_to_type_map(type_id kong_type, spirv_id spirv_type, bool readwrite, storage_class storage) {
-	assert(kong_type != NO_TYPE);
+	kong_assert(kong_type != NO_TYPE);
 
 	complex_type ct = {
 	    .type      = kong_type,
@@ -793,7 +793,7 @@ static void write_types(instructions_buffer *buffer, function *main) {
 			for (size_t j = 0; j < t->members.size; ++j) {
 				member_types[member_types_size] = convert_type_to_spirv_id(t->members.m[j].type.type);
 				member_types_size += 1;
-				assert(member_types_size < 256);
+				kong_assert(member_types_size < 256);
 			}
 
 			spirv_id struct_type = write_type_struct(buffer, member_types, member_types_size);
@@ -817,7 +817,7 @@ static void write_types(instructions_buffer *buffer, function *main) {
 				pointer_relation *previous_relation = &written_pointer_relations.values[relation_index];
 
 				if (previous_relation->pointer_type_id.id == type_map[i].value.id) {
-					assert(previous_relation->non_pointer_type_id.id == non_pointer_type_id.id);
+					kong_assert(previous_relation->non_pointer_type_id.id == non_pointer_type_id.id);
 					found = true;
 					break;
 				}
@@ -991,7 +991,7 @@ static spirv_id get_int_constant(int value) {
 
 	if (container == NULL) {
 		container = (int_constant_container *)malloc(sizeof(int_constant_container));
-		assert(container != NULL);
+		kong_assert(container != NULL);
 		container->container.key = value;
 		container->value         = allocate_index();
 		hash_map_add(int_constants, (struct container *)container);
@@ -1823,8 +1823,8 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 			uint16_t indices_size = o->op_load_access_list.access_list_size;
 
 			if (get_type(o->op_load_access_list.from.type.type)->tex_kind == TEXTURE_KIND_2D) {
-				assert(indices_size == 1);
-				assert(o->op_load_access_list.access_list[0].kind == ACCESS_ELEMENT);
+				kong_assert(indices_size == 1);
+				kong_assert(o->op_load_access_list.access_list[0].kind == ACCESS_ELEMENT);
 
 				////
 				// spirv_id image = write_op_load(instructions, spirv_readwrite_image_type, convert_kong_index_to_spirv_id(o->op_load_access_list.from.index));
@@ -1849,7 +1849,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				for (uint16_t i = 0; i < indices_size; ++i) {
 					switch (o->op_load_access_list.access_list[i].kind) {
 					case ACCESS_ELEMENT:
-						assert(false);
+						kong_assert(false);
 						break;
 					case ACCESS_MEMBER: {
 						uint32_t member_index = 0;
@@ -1862,14 +1862,14 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 							}
 						}
 
-						assert(found);
+						kong_assert(found);
 
 						indices[i] = member_index;
 
 						break;
 					}
 					case ACCESS_SWIZZLE: {
-						assert(o->op_load_access_list.access_list[i].access_swizzle.swizzle.size == 1);
+						kong_assert(o->op_load_access_list.access_list[i].access_swizzle.swizzle.size == 1);
 
 						indices[i] = o->op_load_access_list.access_list[i].access_swizzle.swizzle.indices[0];
 
@@ -1910,7 +1910,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 							}
 						}
 
-						assert(found);
+						kong_assert(found);
 
 						access_kinds[i]  = ACCESS_MEMBER;
 						plain_indices[i] = member_index;
@@ -1919,7 +1919,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 						break;
 					}
 					case ACCESS_SWIZZLE: {
-						assert(o->op_load_access_list.access_list[i].access_swizzle.swizzle.size == 1);
+						kong_assert(o->op_load_access_list.access_list[i].access_swizzle.swizzle.size == 1);
 
 						access_kinds[i]  = ACCESS_SWIZZLE;
 						plain_indices[i] = 0; // unused
@@ -1933,7 +1933,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				}
 
 				type_id access_kong_type = find_access_type(plain_indices, access_kinds, indices_size, o->op_load_access_list.from.type.type);
-				assert(access_kong_type != NO_TYPE);
+				kong_assert(access_kong_type != NO_TYPE);
 
 				spirv_id access_type = {0};
 
@@ -2046,7 +2046,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				}
 				////
 				else {
-					assert(false);
+					kong_assert(false);
 				}
 			}
 			else if (func == add_name("float2")) {
@@ -2061,7 +2061,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 						hmput(index_map, o->op_call.var.index, id);
 					}
 					else {
-						assert(false);
+						kong_assert(false);
 					}
 				}
 				else if (o->op_call.parameters_size == 2) {
@@ -2073,7 +2073,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 					hmput(index_map, o->op_call.var.index, id);
 				}
 				else {
-					assert(false);
+					kong_assert(false);
 				}
 			}
 			else if (func == add_name("float3")) {
@@ -2114,7 +2114,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 					hmput(index_map, o->op_call.var.index, id);
 				}
 				else {
-					assert(false);
+					kong_assert(false);
 				}
 			}
 			else if (func == add_name("int2")) {
@@ -2129,11 +2129,11 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 						hmput(index_map, o->op_call.var.index, id);
 					}
 					else {
-						assert(false);
+						kong_assert(false);
 					}
 				}
 				else {
-					assert(o->op_call.parameters_size == 2);
+					kong_assert(o->op_call.parameters_size == 2);
 					spirv_id constituents[2];
 					for (int i = 0; i < o->op_call.parameters_size; ++i) {
 						constituents[i] = get_var(instructions, o->op_call.parameters[i]);
@@ -2168,7 +2168,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 					hmput(index_map, o->op_call.var.index, id);
 				}
 				else {
-					assert(false);
+					kong_assert(false);
 				}
 			}
 			else if (func == add_name("uint2")) {
@@ -2568,8 +2568,8 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 			type *s = get_type(o->op_store_access_list.to.type.type);
 
 			if (get_type(o->op_store_access_list.to.type.type)->tex_kind == TEXTURE_KIND_2D) {
-				assert(indices_size == 1);
-				assert(o->op_store_access_list.access_list[0].kind == ACCESS_ELEMENT);
+				kong_assert(indices_size == 1);
+				kong_assert(o->op_store_access_list.access_list[0].kind == ACCESS_ELEMENT);
 
 				spirv_id image = write_op_load(instructions, spirv_readwrite_image_type, convert_kong_index_to_spirv_id(o->op_store_access_list.to.index));
 
@@ -2600,7 +2600,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 							}
 						}
 
-						assert(found);
+						kong_assert(found);
 
 						access_kinds[i]  = ACCESS_MEMBER;
 						plain_indices[i] = member_index;
@@ -2610,7 +2610,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 						break;
 					}
 					case ACCESS_SWIZZLE: {
-						assert(o->op_store_access_list.access_list[i].access_swizzle.swizzle.size == 1);
+						kong_assert(o->op_store_access_list.access_list[i].access_swizzle.swizzle.size == 1);
 
 						access_kinds[i]  = ACCESS_SWIZZLE;
 						plain_indices[i] = 0; // unused
@@ -2625,7 +2625,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				}
 
 				type_id access_kong_type = find_access_type(plain_indices, access_kinds, indices_size, o->op_store_access_list.to.type.type);
-				assert(access_kong_type != NO_TYPE);
+				kong_assert(access_kong_type != NO_TYPE);
 
 				spirv_id access_type = {0};
 
@@ -2637,7 +2637,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 					access_type = convert_pointer_type_to_spirv_id(access_kong_type, STORAGE_CLASS_OUTPUT);
 					break;
 				case VARIABLE_INTERNAL:
-					assert(false);
+					kong_assert(false);
 					break;
 				}
 
@@ -2789,7 +2789,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				break;
 			}
 			default:
-				assert(false);
+				kong_assert(false);
 				break;
 			}
 
@@ -2871,7 +2871,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 			break;
 		}
 		case OPCODE_LESS: {
-			assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
+			kong_assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
 
 			spirv_id left  = get_var(instructions, o->op_binary.left);
 			spirv_id right = get_var(instructions, o->op_binary.right);
@@ -2888,7 +2888,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				result = write_op_u_less_than(instructions, spirv_bool_type, left, right);
 			}
 			else {
-				assert(false);
+				kong_assert(false);
 			}
 
 			hmput(index_map, o->op_binary.result.index, result);
@@ -2896,7 +2896,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 			break;
 		}
 		case OPCODE_LESS_EQUAL: {
-			assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
+			kong_assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
 
 			spirv_id left  = get_var(instructions, o->op_binary.left);
 			spirv_id right = get_var(instructions, o->op_binary.right);
@@ -2913,7 +2913,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				result = write_op_u_less_than_equal(instructions, spirv_bool_type, left, right);
 			}
 			else {
-				assert(false);
+				kong_assert(false);
 			}
 
 			hmput(index_map, o->op_binary.result.index, result);
@@ -2921,7 +2921,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 			break;
 		}
 		case OPCODE_GREATER: {
-			assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
+			kong_assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
 
 			spirv_id left  = get_var(instructions, o->op_binary.left);
 			spirv_id right = get_var(instructions, o->op_binary.right);
@@ -2938,7 +2938,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				result = write_op_u_greater_than(instructions, spirv_bool_type, left, right);
 			}
 			else {
-				assert(false);
+				kong_assert(false);
 			}
 
 			hmput(index_map, o->op_binary.result.index, result);
@@ -2946,7 +2946,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 			break;
 		}
 		case OPCODE_GREATER_EQUAL: {
-			assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
+			kong_assert(o->op_binary.left.type.type == o->op_binary.right.type.type);
 
 			spirv_id left  = get_var(instructions, o->op_binary.left);
 			spirv_id right = get_var(instructions, o->op_binary.right);
@@ -2963,7 +2963,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				result = write_op_u_greater_than_equal(instructions, spirv_bool_type, left, right);
 			}
 			else {
-				assert(false);
+				kong_assert(false);
 			}
 
 			hmput(index_map, o->op_binary.result.index, result);
@@ -2983,7 +2983,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				hmput(index_map, o->op_binary.result.index, result);
 			}
 			else {
-				assert(false);
+				kong_assert(false);
 			}
 
 			break;
@@ -3176,7 +3176,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 
 	if (!ends_with_return) {
 		if (main) {
-			assert(stage == SHADER_STAGE_COMPUTE);
+			kong_assert(stage == SHADER_STAGE_COMPUTE);
 		}
 		write_op_return(instructions);
 	}
@@ -3436,7 +3436,7 @@ static void write_globals(instructions_buffer *decorations, instructions_buffer
 		}
 		else if (get_type(base_type)->tex_kind != TEXTURE_KIND_NONE) {
 			if (t->array_size == UINT32_MAX) {
-				assert(false);
+				kong_assert(false);
 			}
 			else {
 				spirv_id image_pointer_type;
@@ -3460,20 +3460,20 @@ static void write_globals(instructions_buffer *decorations, instructions_buffer
 			}
 		}
 		else if (base_type == bvh_type_id) {
-			assert(false);
+			kong_assert(false);
 		}
 		else if (base_type == float_id) {
 			spirv_id id = get_float_constant(g->value.value.floats[0]);
 			hmput(index_map, g->var_index, id);
 		}
 		else if (base_type == float2_id) {
-			assert(false);
+			kong_assert(false);
 		}
 		else if (base_type == float3_id) {
-			assert(false);
+			kong_assert(false);
 		}
 		else if (base_type == float4_id) {
-			assert(false);
+			kong_assert(false);
 		}
 		else {
 			bool root_constant = binding == 0xffffffff;
@@ -3494,7 +3494,7 @@ static void write_globals(instructions_buffer *decorations, instructions_buffer
 				add_to_type_map(member_type, member_pointer_type, false, storage);
 
 				member_types_size += 1;
-				assert(member_types_size < 256);
+				kong_assert(member_types_size < 256);
 			}
 
 			spirv_id struct_type = write_type_struct(aggregate_types_block, member_types, member_types_size);
@@ -3697,7 +3697,7 @@ static char *spirv_export_vertex2(function *main, bool debug, int *size_out) {
 	    .instructions = (uint32_t *)calloc(1024 * 1024, 1),
 	};
 
-	assert(main->parameters_size > 0);
+	kong_assert(main->parameters_size > 0);
 	type_id vertex_input  = main->parameter_types[0].type;
 	type_id vertex_output = main->return_type.type;
 
@@ -3861,7 +3861,7 @@ static char *spirv_export_fragment2(function *main, bool debug, int *size_out) {
 	    .instructions = (uint32_t *)calloc(1024 * 1024, 1),
 	};
 
-	assert(main->parameters_size > 0);
+	kong_assert(main->parameters_size > 0);
 	type_id pixel_input  = main->parameter_types[0].type;
 	type_id pixel_output = main->return_type.type;
 
@@ -3883,7 +3883,7 @@ static char *spirv_export_fragment2(function *main, bool debug, int *size_out) {
 		input_vars[input_var_index] = allocate_index();
 	}
 
-	assert(output->built_in); // has to be a float4 or a float4[]
+	kong_assert(output->built_in); // has to be a float4 or a float4[]
 
 	if (output->array_size > 0) {
 		output_vars_count = output->array_size;

base/sources/kong/kong_wgsl.c

@@ -326,18 +326,18 @@ static void write_globals(char *wgsl, size_t *offset, function *main, bool *fram
 				binding += 1;
 			}
 			else if (base_type == bvh_type_id) {
-				assert(false);
+				kong_assert(false);
 				binding += 1;
 			}
 			else if (get_type(g->type)->built_in) {
 				if (get_type(g->type)->array_size > 0) {
-					assert(false);
+					kong_assert(false);
 					binding += 1;
 				}
 			}
 			else {
 				if (get_type(g->type)->array_size > 0) {
-					assert(false);
+					kong_assert(false);
 					binding += 1;
 				}
 				else {
@@ -427,7 +427,7 @@ static void write_functions(char *code, size_t *offset, shader_stage stage, func
 
 	for (size_t i = 0; i < functions_size; ++i) {
 		function *f = functions[i];
-		assert(f != NULL);
+		kong_assert(f != NULL);
 
 		debug_context context = {0};
 		check(f->block != NULL, context, "Function block missing");
@@ -503,11 +503,11 @@ static void write_functions(char *code, size_t *offset, shader_stage stage, func
 				}
 			}
 			else if (stage == SHADER_STAGE_COMPUTE) {
-				assert(f->parameters_size == 0);
-				assert(f->return_type.type == void_id);
+				kong_assert(f->parameters_size == 0);
+				kong_assert(f->return_type.type == void_id);
 
 				attribute *threads = find_attribute(&f->attributes, add_name("threads"));
-				assert(threads != NULL && threads->paramters_count == 3);
+				kong_assert(threads != NULL && threads->paramters_count == 3);
 
 				*offset += sprintf(&code[*offset],
 				                   "@compute @workgroup_size(%u, %u, %u) fn main(@builtin(local_invocation_id) _kong_group_thread_id: vec3<u32>, "
@@ -555,8 +555,8 @@ static void write_functions(char *code, size_t *offset, shader_stage stage, func
 				type_id from_type = o->op_load_access_list.from.type.type;
 
 				if (is_texture(from_type)) {
-					assert(o->op_load_access_list.access_list_size == 1);
-					assert(o->op_load_access_list.access_list[0].kind == ACCESS_ELEMENT);
+					kong_assert(o->op_load_access_list.access_list_size == 1);
+					kong_assert(o->op_load_access_list.access_list[0].kind == ACCESS_ELEMENT);
 
 					*offset += sprintf(&code[*offset], "var %s: %s = ", get_var(o->op_load_access_list.to, f, main).str,
 					                   type_string(o->op_load_access_list.to.type.type));
@@ -612,9 +612,9 @@ static void write_functions(char *code, size_t *offset, shader_stage stage, func
 				type_id to_type = o->op_store_access_list.to.type.type;
 
 				if (is_texture(to_type)) {
-					assert(o->type == OPCODE_STORE_ACCESS_LIST);
-					assert(o->op_store_access_list.access_list_size == 1);
-					assert(o->op_store_access_list.access_list[0].kind == ACCESS_ELEMENT);
+					kong_assert(o->type == OPCODE_STORE_ACCESS_LIST);
+					kong_assert(o->op_store_access_list.access_list_size == 1);
+					kong_assert(o->op_store_access_list.access_list[0].kind == ACCESS_ELEMENT);
 
 					*offset += sprintf(&code[*offset], "textureStore(%s, vec2<u32>(u32(_%" PRIu64 ".x), u32(_%" PRIu64 ".y)), _%" PRIu64 ");\n",
 					                   get_var(o->op_store_access_list.to, f, main).str, o->op_store_access_list.access_list[0].access_element.index.index,
@@ -668,7 +668,7 @@ static void write_functions(char *code, size_t *offset, shader_stage stage, func
 						*offset += sprintf(&code[*offset], " *= %s;\n", get_var(o->op_store_access_list.from, f, main).str);
 						break;
 					default:
-						assert(false);
+						kong_assert(false);
 						break;
 					}
 				}
@@ -1093,7 +1093,7 @@ static char *wgsl_export_fragment2(function *main) {
 
 	size_t offset = 0;
 
-	assert(main->parameters_size > 0);
+	kong_assert(main->parameters_size > 0);
 	type_id pixel_input = main->parameter_types[0].type;
 
 	check(pixel_input != NO_TYPE, context, "fragment input missing");
@@ -1161,7 +1161,7 @@ void wgsl_export2(char **vs, char **fs) {
 					fragment_functions[fragment_functions_size] = i;
 					fragment_functions_size += 1;
 
-					assert(f->parameters_size > 0);
+					kong_assert(f->parameters_size > 0);
 					fragment_inputs[fragment_inputs_size] = f->parameter_types[0].type;
 					fragment_inputs_size += 1;
 				}