From 7074c780b2eb6fd977a8a0c636146ed0d8a2ff6c Mon Sep 17 00:00:00 2001 From: luboslenco Date: Wed, 23 Sep 2026 22:43:26 +0200 Subject: [PATCH] base: handle invalid shaders --- base/sources/kong/kong.c | 107 ++++++++++++++++++++------------ base/sources/kong/kong.h | 12 +++- base/sources/kong/kong_cstyle.c | 2 +- base/sources/kong/kong_hlsl.c | 10 +-- base/sources/kong/kong_metal.c | 4 +- base/sources/kong/kong_spirv.c | 100 ++++++++++++++--------------- base/sources/kong/kong_wgsl.c | 30 ++++----- 7 files changed, 151 insertions(+), 114 deletions(-) diff --git a/base/sources/kong/kong.c b/base/sources/kong/kong.c index 69dc082c3..d80bcd066 100644 --- a/base/sources/kong/kong.c +++ b/base/sources/kong/kong.c @@ -1,6 +1,7 @@ #include "kong.h" #include "dir.h" #include +#include #include #include #include @@ -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; } diff --git a/base/sources/kong/kong.h b/base/sources/kong/kong.h index f57ab1e7a..b5db0bb98 100644 --- a/base/sources/kong/kong.h +++ b/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); diff --git a/base/sources/kong/kong_cstyle.c b/base/sources/kong/kong_cstyle.c index 1718904e2..839a99f8a 100644 --- a/base/sources/kong/kong_cstyle.c +++ b/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; diff --git a/base/sources/kong/kong_hlsl.c b/base/sources/kong/kong_hlsl.c index b6841b7fe..f4937df4b 100644 --- a/base/sources/kong/kong_hlsl.c +++ b/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}; diff --git a/base/sources/kong/kong_metal.c b/base/sources/kong/kong_metal.c index 174466049..77eb745c8 100644 --- a/base/sources/kong/kong_metal.c +++ b/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; } diff --git a/base/sources/kong/kong_spirv.c b/base/sources/kong/kong_spirv.c index 8678481a8..4a2abd151 100644 --- a/base/sources/kong/kong_spirv.c +++ b/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; diff --git a/base/sources/kong/kong_wgsl.c b/base/sources/kong/kong_wgsl.c index 2050852c7..174f95183 100644 --- a/base/sources/kong/kong_wgsl.c +++ b/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, " @@ -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(_%" 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; }