Update kong to cfc045c2c43e5cd914935a29d6f6659b5f360684

base/sources/libs/kong/analyzer.c

@@ -923,6 +923,43 @@ descriptor_set_group *find_descriptor_set_group_for_function(function *f) {
 	}
 }
 
+void find_sampler_use() {
+	for (function_id i = 0; get_function(i) != NULL; ++i) {
+		function *f = get_function(i);
+
+		if (f->block == NULL) {
+			continue;
+		}
+
+		uint8_t *data = f->code.o;
+		size_t   size = f->code.size;
+
+		size_t index = 0;
+		while (index < size) {
+			opcode *o = (opcode *)&data[index];
+
+			switch (o->type) {
+			case OPCODE_CALL:
+				if (o->op_call.func == add_name("sample") || o->op_call.func == add_name("sample_lod")) {
+					if (is_depth(get_type(o->op_call.parameters[0].type.type)->tex_format)) {
+
+						type *sampler_type = get_type(o->op_call.parameters[1].type.type);
+
+						global *g = find_global_by_var(o->op_call.parameters[1]);
+						assert(g != NULL);
+						g->usage |= GLOBAL_USAGE_SAMPLE_DEPTH;
+					}
+				}
+				break;
+			default:
+				break;
+			}
+
+			index += o->size;
+		}
+	}
+}
+
 void analyze(void) {
 	find_all_render_pipelines();
 	find_render_pipeline_groups();
@@ -933,4 +970,6 @@ void analyze(void) {
 	find_raytracing_pipeline_groups();
 
 	find_descriptor_set_groups();
+
+	find_sampler_use();
 }

base/sources/libs/kong/backends/cpu.c

@@ -180,24 +180,30 @@ static void write_globals(char *code, size_t *offset, char *header_code, size_t
 		if (base_type == sampler_type_id) {
 			*offset += sprintf(&code[*offset], "SamplerState _%" PRIu64 ";\n\n", g->var_index);
 		}
-		else if (base_type == tex2d_type_id) {
-			if (has_attribute(&g->attributes, add_name("write"))) {
-				*offset += sprintf(&code[*offset], "RWTexture2D<float4> _%" PRIu64 ";\n\n", g->var_index);
-			}
-			else {
-				if (t->array_size > 0 && t->array_size == UINT32_MAX) {
-					*offset += sprintf(&code[*offset], "Texture2D<float4> _%" PRIu64 "[];\n\n", g->var_index);
+		else if (get_type(base_type)->tex_kind != TEXTURE_KIND_NONE) {
+			if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D) {
+				if (has_attribute(&g->attributes, add_name("write"))) {
+					*offset += sprintf(&code[*offset], "RWTexture2D<float4> _%" PRIu64 ";\n\n", g->var_index);
 				}
 				else {
-					*offset += sprintf(&code[*offset], "Texture2D<float4> _%" PRIu64 ";\n\n", g->var_index);
+					if (t->array_size > 0 && t->array_size == UINT32_MAX) {
+						*offset += sprintf(&code[*offset], "Texture2D<float4> _%" PRIu64 "[];\n\n", g->var_index);
+					}
+					else {
+						*offset += sprintf(&code[*offset], "Texture2D<float4> _%" PRIu64 ";\n\n", g->var_index);
+					}
 				}
 			}
-		}
-		else if (base_type == tex2darray_type_id) {
-			*offset += sprintf(&code[*offset], "Texture2DArray<float4> _%" PRIu64 ";\n\n", g->var_index);
-		}
-		else if (base_type == texcube_type_id) {
-			*offset += sprintf(&code[*offset], "TextureCube<float4> _%" PRIu64 ";\n\n", g->var_index);
+			else if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+				*offset += sprintf(&code[*offset], "Texture2DArray<float4> _%" PRIu64 ";\n\n", g->var_index);
+			}
+			else if (get_type(base_type)->tex_kind == TEXTURE_KIND_CUBE) {
+				*offset += sprintf(&code[*offset], "TextureCube<float4> _%" PRIu64 ";\n\n", g->var_index);
+			}
+			else {
+				// TODO
+				assert(false);
+			}
 		}
 		else if (base_type == bvh_type_id) {
 			*offset += sprintf(&code[*offset], "RaytracingAccelerationStructure  _%" PRIu64 ";\n\n", g->var_index);

base/sources/libs/kong/backends/glsl.c

@@ -89,19 +89,30 @@ static void write_code(char *glsl, char *directory, const char *filename, const
 	}
 }
 
-static void write_types(char *glsl, size_t *offset, shader_stage stage, type_id input, type_id output, function *main) {
+static void write_types(char *glsl, size_t *offset, shader_stage stage, type_id inputs[64], size_t inputs_count, type_id output, function *main) {
 	type_id types[256];
 	size_t  types_size = 0;
 	find_referenced_types(main, types, &types_size);
 
+	size_t input_location = 0;
+
 	for (size_t i = 0; i < types_size; ++i) {
 		type *t = get_type(types[i]);
 
 		if (!t->built_in && !has_attribute(&t->attributes, add_name("pipe"))) {
-			if (stage == SHADER_STAGE_VERTEX && types[i] == input) {
+			bool type_is_input = false;
+			for (size_t input_index = 0; input_index < inputs_count; ++input_index) {
+				if (types[i] == inputs[input_index]) {
+					type_is_input = true;
+					break;
+				}
+			}
+
+			if (stage == SHADER_STAGE_VERTEX && type_is_input) {
 				for (size_t j = 0; j < t->members.size; ++j) {
-					*offset += sprintf(&glsl[*offset], "layout(location = %zu) in %s %s_%s;\n", j, type_string(t->members.m[j].type.type), get_name(t->name),
-					                   get_name(t->members.m[j].name));
+					*offset += sprintf(&glsl[*offset], "layout(location = %zu) in %s %s_%s;\n", input_location, type_string(t->members.m[j].type.type),
+					                   get_name(t->name), get_name(t->members.m[j].name));
+					++input_location;
 				}
 			}
 			else if (stage == SHADER_STAGE_VERTEX && types[i] == output) {
@@ -112,7 +123,7 @@ static void write_types(char *glsl, size_t *offset, shader_stage stage, type_id
 					}
 				}
 			}
-			else if (stage == SHADER_STAGE_FRAGMENT && types[i] == input) {
+			else if (stage == SHADER_STAGE_FRAGMENT && type_is_input) {
 				for (size_t j = 0; j < t->members.size; ++j) {
 					if (j != 0) {
 						*offset += sprintf(&glsl[*offset], "in %s %s_%s;\n", type_string(t->members.m[j].type.type), get_name(t->name),
@@ -150,14 +161,20 @@ static void write_globals(char *glsl, size_t *offset, function *main) {
 
 		if (g->type == sampler_type_id) {
 		}
-		else if (g->type == tex2d_type_id) {
-			*offset += sprintf(&glsl[*offset], "uniform sampler2D _%" PRIu64 ";\n\n", g->var_index);
-		}
-		else if (g->type == tex2darray_type_id) {
-			*offset += sprintf(&glsl[*offset], "uniform sampler2DArray _%" PRIu64 ";\n\n", g->var_index);
-		}
-		else if (g->type == texcube_type_id) {
-			*offset += sprintf(&glsl[*offset], "uniform samplerCube _%" PRIu64 ";\n\n", g->var_index);
+		else if (get_type(g->type)->tex_kind != TEXTURE_KIND_NONE) {
+			if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D) {
+				*offset += sprintf(&glsl[*offset], "uniform sampler2D _%" PRIu64 ";\n\n", g->var_index);
+			}
+			else if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+				*offset += sprintf(&glsl[*offset], "uniform sampler2DArray _%" PRIu64 ";\n\n", g->var_index);
+			}
+			else if (get_type(g->type)->tex_kind == TEXTURE_KIND_CUBE) {
+				*offset += sprintf(&glsl[*offset], "uniform samplerCube _%" PRIu64 ";\n\n", g->var_index);
+			}
+			else {
+				// TODO
+				assert(false);
+			}
 		}
 		else if (g->type == float_id) {
 		}
@@ -173,7 +190,8 @@ static void write_globals(char *glsl, size_t *offset, function *main) {
 	}
 }
 
-static void write_functions(char *code, size_t *offset, shader_stage stage, type_id input, type_id output, function *main, bool flip) {
+static void write_functions(char *code, size_t *offset, shader_stage stage, type_id inputs[64], size_t inputs_count, type_id output, function *main,
+                            bool flip) {
 	function *functions[256];
 	size_t    functions_size = 0;
 
@@ -325,7 +343,17 @@ static void write_functions(char *code, size_t *offset, shader_stage stage, type
 
 				indent(code, offset, indentation);
 
-				if (f == main && o->op_load_access_list.from.type.type == input) {
+				bool from_is_input = false;
+				if (f == main) {
+					for (size_t input_index = 0; input_index < inputs_count; ++input_index) {
+						if (o->op_load_access_list.from.type.type == inputs[input_index]) {
+							from_is_input = true;
+							break;
+						}
+					}
+				}
+
+				if (from_is_input) {
 					*offset += sprintf(&code[*offset], "%s _%" PRIu64 " = %s", type_string(o->op_load_access_list.to.type.type),
 					                   o->op_load_access_list.to.index, type_string(o->op_load_access_list.from.type.type));
 				}
@@ -342,7 +370,7 @@ static void write_functions(char *code, size_t *offset, shader_stage stage, type
 						*offset += sprintf(&code[*offset], "[_%" PRIu64 "]", o->op_load_access_list.access_list[i].access_element.index.index);
 						break;
 					case ACCESS_MEMBER:
-						if (global_var_index != 0 || (f == main && o->op_load_access_list.from.type.type == input && i == 0)) {
+						if (global_var_index != 0 || (from_is_input && i == 0)) {
 							*offset += sprintf(&code[*offset], "_%s", get_name(o->op_load_access_list.access_list[i].access_member.name));
 						}
 						else {
@@ -457,20 +485,21 @@ static void glsl_export_vertex(char *directory, function *main, bool flip) {
 
 	size_t offset = 0;
 
-	assert(main->parameters_size > 0);
-	type_id vertex_input  = main->parameter_types[0].type;
+	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;
+	}
 	type_id vertex_output = main->return_type.type;
 
-	check(vertex_input != NO_TYPE, context, "vertex input missing");
 	check(vertex_output != NO_TYPE, context, "vertex output missing");
 
 	offset += sprintf(&glsl[offset], "#version 330\n\n");
 
-	write_types(glsl, &offset, SHADER_STAGE_VERTEX, vertex_input, vertex_output, main);
+	write_types(glsl, &offset, SHADER_STAGE_VERTEX, vertex_inputs, main->parameters_size, vertex_output, main);
 
 	write_globals(glsl, &offset, main);
 
-	write_functions(glsl, &offset, SHADER_STAGE_VERTEX, vertex_input, vertex_output, main, flip);
+	write_functions(glsl, &offset, SHADER_STAGE_VERTEX, vertex_inputs, main->parameters_size, vertex_output, main, flip);
 
 	char *name = get_name(main->name);
 
@@ -507,11 +536,11 @@ static void glsl_export_fragment(char *directory, function *main) {
 
 	offset += sprintf(&glsl[offset], "#version 330\n\n");
 
-	write_types(glsl, &offset, SHADER_STAGE_FRAGMENT, pixel_input, NO_TYPE, main);
+	write_types(glsl, &offset, SHADER_STAGE_FRAGMENT, &pixel_input, 1, NO_TYPE, main);
 
 	write_globals(glsl, &offset, main);
 
-	write_functions(glsl, &offset, SHADER_STAGE_FRAGMENT, pixel_input, NO_TYPE, main, false);
+	write_functions(glsl, &offset, SHADER_STAGE_FRAGMENT, &pixel_input, 1, NO_TYPE, main, false);
 
 	char *name = get_name(main->name);
 
@@ -535,11 +564,11 @@ static void glsl_export_compute(char *directory, function *main) {
 
 	offset += sprintf(&glsl[offset], "#version 330\n\n");
 
-	write_types(glsl, &offset, SHADER_STAGE_COMPUTE, NO_TYPE, NO_TYPE, main);
+	write_types(glsl, &offset, SHADER_STAGE_COMPUTE, NULL, 0, NO_TYPE, main);
 
 	write_globals(glsl, &offset, main);
 
-	write_functions(glsl, &offset, SHADER_STAGE_COMPUTE, NO_TYPE, NO_TYPE, main, false);
+	write_functions(glsl, &offset, SHADER_STAGE_COMPUTE, NULL, 0, NO_TYPE, main, false);
 
 	char *name = get_name(main->name);
 

base/sources/libs/kong/backends/hlsl.c

@@ -66,8 +66,14 @@ static char *type_string(type_id type) {
 	if (type == bvh_type_id) {
 		return "RaytracingAccelerationStructure";
 	}
-	if (type == tex2d_type_id) {
-		return "Texture2D<float4>";
+	if (get_type(type)->tex_kind != TEXTURE_KIND_NONE) {
+		if (get_type(type)->tex_kind == TEXTURE_KIND_2D) {
+			return "Texture2D<float4>";
+		}
+		else {
+			// TODO
+			assert(false);
+		}
 	}
 	return get_name(get_type(type)->name);
 }
@@ -319,7 +325,7 @@ static void assign_register_indices(uint32_t *register_indices, function *shader
 				register_indices[global_index] = sampler_index;
 				sampler_index += 1;
 			}
-			else if (base_type == tex2d_type_id) {
+			else if (get_type(base_type)->tex_kind != TEXTURE_KIND_NONE) {
 				if (t->array_size == UINT32_MAX) {
 					register_indices[global_index] = 0;
 				}
@@ -332,7 +338,7 @@ static void assign_register_indices(uint32_t *register_indices, function *shader
 					srv_index += 1;
 				}
 			}
-			else if (base_type == texcube_type_id || base_type == tex2darray_type_id || base_type == bvh_type_id) {
+			else if (base_type == bvh_type_id) {
 				register_indices[global_index] = srv_index;
 				srv_index += 1;
 			}
@@ -402,24 +408,30 @@ static void write_globals(char *hlsl, size_t *offset, function *main, function *
 		if (base_type == sampler_type_id) {
 			*offset += sprintf(&hlsl[*offset], "SamplerState _%" PRIu64 " : register(s%i);\n\n", g->var_index, register_index);
 		}
-		else if (base_type == tex2d_type_id) {
-			if (writable) {
-				*offset += sprintf(&hlsl[*offset], "RWTexture2D<float4> _%" PRIu64 " : register(u%i);\n\n", g->var_index, register_index);
-			}
-			else {
-				if (t->array_size > 0 && t->array_size == UINT32_MAX) {
-					*offset += sprintf(&hlsl[*offset], "Texture2D<float4> _%" PRIu64 "[] : register(t%i, space1);\n\n", g->var_index, register_index);
+		else if (get_type(base_type)->tex_kind != TEXTURE_KIND_NONE) {
+			if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D) {
+				if (writable) {
+					*offset += sprintf(&hlsl[*offset], "RWTexture2D<float4> _%" PRIu64 " : register(u%i);\n\n", g->var_index, register_index);
 				}
 				else {
-					*offset += sprintf(&hlsl[*offset], "Texture2D<float4> _%" PRIu64 " : register(t%i);\n\n", g->var_index, register_index);
+					if (t->array_size > 0 && t->array_size == UINT32_MAX) {
+						*offset += sprintf(&hlsl[*offset], "Texture2D<float4> _%" PRIu64 "[] : register(t%i, space1);\n\n", g->var_index, register_index);
+					}
+					else {
+						*offset += sprintf(&hlsl[*offset], "Texture2D<float4> _%" PRIu64 " : register(t%i);\n\n", g->var_index, register_index);
+					}
 				}
 			}
-		}
-		else if (base_type == tex2darray_type_id) {
-			*offset += sprintf(&hlsl[*offset], "Texture2DArray<float4> _%" PRIu64 " : register(t%i);\n\n", g->var_index, register_index);
-		}
-		else if (base_type == texcube_type_id) {
-			*offset += sprintf(&hlsl[*offset], "TextureCube<float4> _%" PRIu64 " : register(t%i);\n\n", g->var_index, register_index);
+			else if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+				*offset += sprintf(&hlsl[*offset], "Texture2DArray<float4> _%" PRIu64 " : register(t%i);\n\n", g->var_index, register_index);
+			}
+			else if (get_type(base_type)->tex_kind == TEXTURE_KIND_CUBE) {
+				*offset += sprintf(&hlsl[*offset], "TextureCube<float4> _%" PRIu64 " : register(t%i);\n\n", g->var_index, register_index);
+			}
+			else {
+				// TODO
+				assert(false);
+			}
 		}
 		else if (base_type == bvh_type_id) {
 			*offset += sprintf(&hlsl[*offset], "RaytracingAccelerationStructure  _%" PRIu64 " : register(t%i);\n\n", g->var_index, register_index);
@@ -1139,7 +1151,7 @@ static void write_functions(char *hlsl, size_t *offset, shader_stage stage, func
 					case ACCESS_ELEMENT: {
 						type *from_type = get_type(o->op_load_access_list.from.type.type);
 
-						if (from_type->array_size == UINT32_MAX && from_type->base == tex2d_type_id) {
+						if (from_type->array_size == UINT32_MAX && get_type(from_type->base)->tex_kind != TEXTURE_KIND_NONE) {
 							*offset += sprintf(&hlsl[*offset], "[NonUniformResourceIndex(_%" PRIu64 ")]",
 							                   o->op_load_access_list.access_list[i].access_element.index.index);
 						}

base/sources/libs/kong/backends/metal.c

@@ -251,13 +251,7 @@ static void write_argument_buffers(char *code, size_t *offset) {
 				}
 			}
 			else if (is_texture(g->type)) {
-				if (g->type == tex2darray_type_id) {
-					*offset += sprintf(&code[*offset], "\ttexture2d_array<float> _%" PRIu64 " [[id(%zu)]];\n", g->var_index, global_index);
-				}
-				else if (g->type == texcube_type_id) {
-					*offset += sprintf(&code[*offset], "\ttexturecube<float> _%" PRIu64 " [[id(%zu)]];\n", g->var_index, global_index);
-				}
-				else {
+				if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D) {
 					if (writable) {
 						*offset += sprintf(&code[*offset], "\ttexture2d<float, access::write> _%" PRIu64 " [[id(%zu)]];\n", g->var_index, global_index);
 					}
@@ -265,6 +259,16 @@ static void write_argument_buffers(char *code, size_t *offset) {
 						*offset += sprintf(&code[*offset], "\ttexture2d<float> _%" PRIu64 " [[id(%zu)]];\n", g->var_index, global_index);
 					}
 				}
+				else if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+					*offset += sprintf(&code[*offset], "\ttexture2d_array<float> _%" PRIu64 " [[id(%zu)]];\n", g->var_index, global_index);
+				}
+				else if (get_type(g->type)->tex_kind == TEXTURE_KIND_CUBE) {
+					*offset += sprintf(&code[*offset], "\ttexturecube<float> _%" PRIu64 " [[id(%zu)]];\n", g->var_index, global_index);
+				}
+				else {
+					// TODO
+					assert(false);
+				}
 			}
 			else if (is_sampler(g->type)) {
 				*offset += sprintf(&code[*offset], "\tsampler _%" PRIu64 " [[id(%zu)]];\n", g->var_index, global_index);
@@ -666,18 +670,28 @@ static void write_functions(char *code, size_t *offset) {
 				if (o->op_call.func == add_name("sample")) {
 					check(o->op_call.parameters_size == 3, context, "sample requires three parameters");
 
-					if (o->op_call.parameters[0].type.type == tex2darray_type_id) {
+					variable image_var = o->op_call.parameters[0];
+
+					if (get_type(o->op_call.parameters[0].type.type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
 						*offset +=
 						    sprintf(&code[*offset],
 						            "%s _%" PRIu64 " = argument_buffer0._%" PRIu64 ".sample(argument_buffer0._%" PRIu64 ", _%" PRIu64 ".xy, _%" PRIu64 ".z);\n",
-						            type_string(o->op_call.var.type.type), o->op_call.var.index, o->op_call.parameters[0].index, o->op_call.parameters[1].index,
+						            type_string(o->op_call.var.type.type), o->op_call.var.index, image_var.index, o->op_call.parameters[1].index,
 						            o->op_call.parameters[2].index, o->op_call.parameters[2].index);
 					}
 					else {
-						*offset +=
-						    sprintf(&code[*offset], "%s _%" PRIu64 " = argument_buffer0._%" PRIu64 ".sample(argument_buffer0._%" PRIu64 ", _%" PRIu64 ");\n",
-						            type_string(o->op_call.var.type.type), o->op_call.var.index, o->op_call.parameters[0].index, o->op_call.parameters[1].index,
-						            o->op_call.parameters[2].index);
+						if (is_depth(get_type(image_var.type.type)->tex_format)) {
+							*offset += sprintf(&code[*offset],
+							                   "%s _%" PRIu64 " = argument_buffer0._%" PRIu64 ".sample(argument_buffer0._%" PRIu64 ", _%" PRIu64 ").r;\n",
+							                   type_string(o->op_call.var.type.type), o->op_call.var.index, image_var.index, o->op_call.parameters[1].index,
+							                   o->op_call.parameters[2].index);
+						}
+						else {
+							*offset += sprintf(&code[*offset],
+							                   "%s _%" PRIu64 " = argument_buffer0._%" PRIu64 ".sample(argument_buffer0._%" PRIu64 ", _%" PRIu64 ");\n",
+							                   type_string(o->op_call.var.type.type), o->op_call.var.index, image_var.index, o->op_call.parameters[1].index,
+							                   o->op_call.parameters[2].index);
+						}
 					}
 				}
 				else if (o->op_call.func == add_name("sample_lod")) {
@@ -889,7 +903,7 @@ char *metal_export(char *directory) {
 			global_register_indices[i] = sampler_index;
 			sampler_index += 1;
 		}
-		else if (g->type == tex2d_type_id || g->type == texcube_type_id) {
+		else if (is_texture(g->type)) {
 			global_register_indices[i] = texture_index;
 			texture_index += 1;
 		}

base/sources/libs/kong/backends/spirv.c

@@ -277,7 +277,7 @@ typedef enum spirv_glsl_std {
 } spirv_glsl_std;
 
 static type_id find_access_type(int *indices, access_kind *access_kinds, int indices_size, type_id base_type) {
-	if (base_type == tex2d_type_id) {
+	if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D) {
 		assert(indices_size == 1);
 		return float4_id;
 	}
@@ -827,6 +827,14 @@ static void write_base_types(instructions_buffer *buffer) {
 
 static spirv_id get_int_constant(int value);
 
+typedef struct pointer_relation {
+	spirv_id non_pointer_type_id;
+	spirv_id pointer_type_id;
+} pointer_relation;
+
+static_array(pointer_relation, written_pointers, 256);
+static written_pointers written_pointer_relations;
+
 static void write_types(instructions_buffer *buffer, function *main) {
 	type_id types[256];
 	size_t  types_size = 0;
@@ -856,6 +864,8 @@ static void write_types(instructions_buffer *buffer, function *main) {
 		}
 	}
 
+	static_array_init(written_pointer_relations);
+
 	size_t size = hmlenu(type_map);
 	for (size_t i = 0; i < size; ++i) {
 		complex_type type = type_map[i].key;
@@ -864,7 +874,27 @@ static void write_types(instructions_buffer *buffer, function *main) {
 			non_pointer_type.storage      = STORAGE_CLASS_NONE;
 			spirv_id non_pointer_type_id  = convert_complex_type_to_spirv_id(non_pointer_type);
 
-			write_type_pointer_preallocated(buffer, type.storage, non_pointer_type_id, type_map[i].value);
+			bool found = false;
+
+			for (size_t relation_index = 0; relation_index < written_pointer_relations.size; ++relation_index) {
+				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);
+					found = true;
+					break;
+				}
+			}
+
+			if (!found) {
+				pointer_relation relation = {
+				    .non_pointer_type_id = non_pointer_type_id,
+				    .pointer_type_id     = type_map[i].value,
+				};
+				static_array_push(written_pointer_relations, relation);
+
+				write_type_pointer_preallocated(buffer, type.storage, non_pointer_type_id, type_map[i].value);
+			}
 		}
 	}
 }
@@ -1806,9 +1836,9 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 		}
 	}
 
-	bool ends_with_return = false;
+	bool     ends_with_return     = false;
 	uint64_t next_block_branch_id = 0;
-	uint64_t next_block_label_id = 0;
+	uint64_t next_block_label_id  = 0;
 
 	index = 0;
 	while (index < size) {
@@ -1821,7 +1851,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 		case OPCODE_LOAD_ACCESS_LIST: {
 			uint16_t indices_size = o->op_load_access_list.access_list_size;
 
-			if (o->op_load_access_list.from.type.type == tex2d_type_id) {
+			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);
 
@@ -1995,17 +2025,23 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				spirv_id image_type;
 				spirv_id sampled_image_type;
 
-				if (image_var.type.type == tex2d_type_id) {
-					image_type         = spirv_image_type;
-					sampled_image_type = spirv_sampled_image_type;
-				}
-				else if (image_var.type.type == tex2darray_type_id) {
-					image_type         = spirv_image2darray_type;
-					sampled_image_type = spirv_sampled_image2darray_type;
-				}
-				else if (image_var.type.type == texcube_type_id) {
-					image_type         = spirv_imagecube_type;
-					sampled_image_type = spirv_sampled_imagecube_type;
+				if (get_type(image_var.type.type)->tex_kind != TEXTURE_KIND_NONE) {
+					if (get_type(image_var.type.type)->tex_kind == TEXTURE_KIND_2D) {
+						image_type         = spirv_image_type;
+						sampled_image_type = spirv_sampled_image_type;
+					}
+					else if (get_type(image_var.type.type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+						image_type         = spirv_image2darray_type;
+						sampled_image_type = spirv_sampled_image2darray_type;
+					}
+					else if (get_type(image_var.type.type)->tex_kind == TEXTURE_KIND_CUBE) {
+						image_type         = spirv_imagecube_type;
+						sampled_image_type = spirv_sampled_imagecube_type;
+					}
+					else {
+						// TODO
+						assert(false);
+					}
 				}
 
 				spirv_id image         = write_op_load(instructions, image_type, convert_kong_index_to_spirv_id(image_var.index));
@@ -2014,6 +2050,13 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				spirv_id coordinate    = get_var(instructions, o->op_call.parameters[2]);
 
 				spirv_id id = write_op_image_sample_implicit_lod(instructions, spirv_float4_type, sampled_image, coordinate);
+
+				if (is_depth(get_type(image_var.type.type)->tex_format)) {
+					uint32_t index = 0;
+
+					id = write_op_composite_extract(instructions, spirv_float_type, id, &index, 1);
+				}
+
 				hmput(index_map, o->op_call.var.index, id);
 			}
 			else if (func == add_name("sample_lod")) {
@@ -2022,17 +2065,23 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				spirv_id image_type;
 				spirv_id sampled_image_type;
 
-				if (image_var.type.type == tex2d_type_id) {
-					image_type         = spirv_image_type;
-					sampled_image_type = spirv_sampled_image_type;
-				}
-				else if (image_var.type.type == tex2darray_type_id) {
-					image_type         = spirv_image2darray_type;
-					sampled_image_type = spirv_sampled_image2darray_type;
-				}
-				else if (image_var.type.type == texcube_type_id) {
-					image_type         = spirv_imagecube_type;
-					sampled_image_type = spirv_sampled_imagecube_type;
+				if (get_type(image_var.type.type)->tex_kind != TEXTURE_KIND_NONE) {
+					if (get_type(image_var.type.type)->tex_kind == TEXTURE_KIND_2D) {
+						image_type         = spirv_image_type;
+						sampled_image_type = spirv_sampled_image_type;
+					}
+					else if (get_type(image_var.type.type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+						image_type         = spirv_image2darray_type;
+						sampled_image_type = spirv_sampled_image2darray_type;
+					}
+					else if (get_type(image_var.type.type)->tex_kind == TEXTURE_KIND_CUBE) {
+						image_type         = spirv_imagecube_type;
+						sampled_image_type = spirv_sampled_imagecube_type;
+					}
+					else {
+						// TODO
+						assert(false);
+					}
 				}
 
 				spirv_id image         = write_op_load(instructions, image_type, convert_kong_index_to_spirv_id(image_var.index));
@@ -2098,20 +2147,20 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 				spirv_id id = write_op_composite_construct(instructions, spirv_float3_type, constituents, o->op_call.parameters_size);
 				hmput(index_map, o->op_call.var.index, id);
 			}
-			else if (func == add_name("float3x3")) {
-				spirv_id constituents[3];
+			else if (func == add_name("float4")) {
+				spirv_id constituents[4];
 				for (int i = 0; i < o->op_call.parameters_size; ++i) {
 					constituents[i] = get_var(instructions, o->op_call.parameters[i]);
 				}
-				spirv_id id = write_op_composite_construct(instructions, spirv_float3x3_type, constituents, o->op_call.parameters_size);
+				spirv_id id = write_op_composite_construct(instructions, spirv_float4_type, constituents, o->op_call.parameters_size);
 				hmput(index_map, o->op_call.var.index, id);
 			}
-			else if (func == add_name("float4")) {
-				spirv_id constituents[4];
+			else if (func == add_name("float3x3")) {
+				spirv_id constituents[3];
 				for (int i = 0; i < o->op_call.parameters_size; ++i) {
 					constituents[i] = get_var(instructions, o->op_call.parameters[i]);
 				}
-				spirv_id id = write_op_composite_construct(instructions, spirv_float4_type, constituents, o->op_call.parameters_size);
+				spirv_id id = write_op_composite_construct(instructions, spirv_float3x3_type, constituents, o->op_call.parameters_size);
 				hmput(index_map, o->op_call.var.index, id);
 			}
 			else if (func == add_name("float4x4")) {
@@ -2520,7 +2569,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 
 			type *s = get_type(o->op_store_access_list.to.type.type);
 
-			if (o->op_store_access_list.to.type.type == tex2d_type_id) {
+			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);
 
@@ -2752,7 +2801,6 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 						write_op_store(instructions, output_vars[i], value);
 					}
 				}
-
 				write_op_return(instructions);
 			}
 			else if (stage == SHADER_STAGE_FRAGMENT && main) {
@@ -2773,20 +2821,19 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 					spirv_id loaded = get_var(instructions, o->op_return.var);
 					write_op_store(instructions, output_vars[0], loaded);
 				}
-
 				write_op_return(instructions);
 			}
 			else {
 				spirv_id return_value = get_var(instructions, o->op_return.var);
 				write_op_return_value(instructions, return_value);
 			}
-			ends_with_return = true;
+			ends_with_return     = true;
 			next_block_branch_id = 0;
 			break;
 		}
 		case OPCODE_DISCARD: {
 			write_op_discard(instructions);
-			ends_with_return = true;
+			ends_with_return     = true;
 			next_block_branch_id = 0;
 			break;
 		}
@@ -2999,7 +3046,7 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 		}
 		case OPCODE_IF: {
 			next_block_branch_id = o->op_if.end_id;
-			next_block_label_id = o->op_if.end_id;
+			next_block_label_id  = o->op_if.end_id;
 			write_op_selection_merge(instructions, convert_kong_index_to_spirv_id(o->op_if.end_id), SELECTION_CONTROL_NONE);
 
 			write_op_branch_conditional(instructions, convert_kong_index_to_spirv_id(o->op_if.condition.index),
@@ -3056,7 +3103,6 @@ static void write_function(instructions_buffer *instructions, function *f, spirv
 			if (o->op_block.id == next_block_label_id) {
 				write_op_label_preallocated(instructions, convert_kong_index_to_spirv_id(o->op_block.id));
 			}
-
 			break;
 		}
 		default: {
@@ -3127,7 +3173,7 @@ static void write_functions(instructions_buffer *instructions, function *main, s
 					}
 				}
 				if (parameters_match) {
-					function_type_index = j;
+					function_type_index = (int)j;
 					break;
 				}
 			}
@@ -3216,7 +3262,7 @@ static void assign_bindings(uint32_t *bindings, function *shader) {
 				bindings[global_index] = binding;
 				binding += 1;
 			}
-			else if (base_type == tex2d_type_id) {
+			else if (get_type(base_type)->tex_kind != TEXTURE_KIND_NONE) {
 				if (t->array_size == UINT32_MAX) {
 					bindings[global_index] = 0;
 				}
@@ -3225,7 +3271,7 @@ static void assign_bindings(uint32_t *bindings, function *shader) {
 					binding += 1;
 				}
 			}
-			else if (base_type == texcube_type_id || base_type == tex2darray_type_id || base_type == bvh_type_id) {
+			else if (base_type == bvh_type_id) {
 				bindings[global_index] = binding;
 				binding += 1;
 			}
@@ -3329,75 +3375,81 @@ static void write_globals(instructions_buffer *decorations, instructions_buffer
 			write_op_decorate_value(decorations, spirv_var_id, DECORATION_DESCRIPTOR_SET, 0);
 			write_op_decorate_value(decorations, spirv_var_id, DECORATION_BINDING, binding);
 		}
-		else if (base_type == tex2d_type_id) {
-			if (t->array_size == UINT32_MAX) {
-				assert(false);
-			}
-			else {
-				spirv_id image_pointer_type;
-
-				if (readable || writable) {
-					add_to_type_map(g->type, spirv_readwrite_image_type, true, STORAGE_CLASS_NONE);
-					image_pointer_type = spirv_readwrite_image_pointer_type;
+		else if (get_type(base_type)->tex_kind != TEXTURE_KIND_NONE) {
+			if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D) {
+				if (t->array_size == UINT32_MAX) {
+					assert(false);
 				}
 				else {
-					add_to_type_map(g->type, spirv_image_type, false, STORAGE_CLASS_NONE);
-					image_pointer_type = spirv_image_pointer_type;
-				}
+					spirv_id image_pointer_type;
 
-				add_to_type_map(g->type, image_pointer_type, readable || writable, STORAGE_CLASS_UNIFORM_CONSTANT);
+					if (readable || writable) {
+						add_to_type_map(g->type, spirv_readwrite_image_type, true, STORAGE_CLASS_NONE);
+						image_pointer_type = spirv_readwrite_image_pointer_type;
+					}
+					else {
+						add_to_type_map(g->type, spirv_image_type, false, STORAGE_CLASS_NONE);
+						image_pointer_type = spirv_image_pointer_type;
+					}
 
-				spirv_id spirv_var_id = convert_kong_index_to_spirv_id(g->var_index);
-				write_op_variable_preallocated(global_vars_block, image_pointer_type, spirv_var_id, STORAGE_CLASS_UNIFORM_CONSTANT);
+					add_to_type_map(g->type, image_pointer_type, readable || writable, STORAGE_CLASS_UNIFORM_CONSTANT);
 
-				write_op_decorate_value(decorations, spirv_var_id, DECORATION_DESCRIPTOR_SET, 0);
-				write_op_decorate_value(decorations, spirv_var_id, DECORATION_BINDING, binding);
-			}
-		}
-		else if (base_type == tex2darray_type_id) {
-			if (t->array_size == UINT32_MAX) {
-				assert(false);
-			}
-			else {
-				spirv_id image_pointer_type;
+					spirv_id spirv_var_id = convert_kong_index_to_spirv_id(g->var_index);
+					write_op_variable_preallocated(global_vars_block, image_pointer_type, spirv_var_id, STORAGE_CLASS_UNIFORM_CONSTANT);
 
-				if (writable) {
+					write_op_decorate_value(decorations, spirv_var_id, DECORATION_DESCRIPTOR_SET, 0);
+					write_op_decorate_value(decorations, spirv_var_id, DECORATION_BINDING, binding);
+				}
+			}
+			else if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+				if (t->array_size == UINT32_MAX) {
 					assert(false);
 				}
 				else {
-					add_to_type_map(g->type, spirv_image2darray_type, false, STORAGE_CLASS_NONE);
-					add_to_type_map(g->type, spirv_image2darray_pointer_type, false, STORAGE_CLASS_UNIFORM_CONSTANT);
-					image_pointer_type = spirv_image2darray_pointer_type;
-				}
+					spirv_id image_pointer_type;
 
-				spirv_id spirv_var_id = convert_kong_index_to_spirv_id(g->var_index);
-				write_op_variable_preallocated(global_vars_block, image_pointer_type, spirv_var_id, STORAGE_CLASS_UNIFORM_CONSTANT);
+					if (writable) {
+						assert(false);
+					}
+					else {
+						add_to_type_map(g->type, spirv_image2darray_type, false, STORAGE_CLASS_NONE);
+						add_to_type_map(g->type, spirv_image2darray_pointer_type, false, STORAGE_CLASS_UNIFORM_CONSTANT);
+						image_pointer_type = spirv_image2darray_pointer_type;
+					}
 
-				write_op_decorate_value(decorations, spirv_var_id, DECORATION_DESCRIPTOR_SET, 0);
-				write_op_decorate_value(decorations, spirv_var_id, DECORATION_BINDING, binding);
-			}
-		}
-		else if (base_type == texcube_type_id) {
-			if (t->array_size == UINT32_MAX) {
-				assert(false);
-			}
-			else {
-				spirv_id image_pointer_type;
+					spirv_id spirv_var_id = convert_kong_index_to_spirv_id(g->var_index);
+					write_op_variable_preallocated(global_vars_block, image_pointer_type, spirv_var_id, STORAGE_CLASS_UNIFORM_CONSTANT);
 
-				if (writable) {
+					write_op_decorate_value(decorations, spirv_var_id, DECORATION_DESCRIPTOR_SET, 0);
+					write_op_decorate_value(decorations, spirv_var_id, DECORATION_BINDING, binding);
+				}
+			}
+			else if (get_type(base_type)->tex_kind == TEXTURE_KIND_CUBE) {
+				if (t->array_size == UINT32_MAX) {
 					assert(false);
 				}
 				else {
-					add_to_type_map(g->type, spirv_imagecube_type, false, STORAGE_CLASS_NONE);
-					add_to_type_map(g->type, spirv_imagecube_pointer_type, false, STORAGE_CLASS_UNIFORM_CONSTANT);
-					image_pointer_type = spirv_imagecube_pointer_type;
-				}
+					spirv_id image_pointer_type;
 
-				spirv_id spirv_var_id = convert_kong_index_to_spirv_id(g->var_index);
-				write_op_variable_preallocated(global_vars_block, image_pointer_type, spirv_var_id, STORAGE_CLASS_UNIFORM_CONSTANT);
+					if (writable) {
+						assert(false);
+					}
+					else {
+						add_to_type_map(g->type, spirv_imagecube_type, false, STORAGE_CLASS_NONE);
+						add_to_type_map(g->type, spirv_imagecube_pointer_type, false, STORAGE_CLASS_UNIFORM_CONSTANT);
+						image_pointer_type = spirv_imagecube_pointer_type;
+					}
 
-				write_op_decorate_value(decorations, spirv_var_id, DECORATION_DESCRIPTOR_SET, 0);
-				write_op_decorate_value(decorations, spirv_var_id, DECORATION_BINDING, binding);
+					spirv_id spirv_var_id = convert_kong_index_to_spirv_id(g->var_index);
+					write_op_variable_preallocated(global_vars_block, image_pointer_type, spirv_var_id, STORAGE_CLASS_UNIFORM_CONSTANT);
+
+					write_op_decorate_value(decorations, spirv_var_id, DECORATION_DESCRIPTOR_SET, 0);
+					write_op_decorate_value(decorations, spirv_var_id, DECORATION_BINDING, binding);
+				}
+			}
+			else {
+				// TODO
+				assert(false);
 			}
 		}
 		else if (base_type == bvh_type_id) {
@@ -3615,7 +3667,7 @@ static void spirv_export_vertex(char *directory, function *main, bool debug) {
 			input_types[input_type_index] = input->members.m[member_index].type.type;
 
 			vertex_parameter_indices[input_type_index]        = input_index;
-			vertex_parameter_member_indices[input_type_index] = member_index;
+			vertex_parameter_member_indices[input_type_index] = (uint32_t)member_index;
 
 			++input_type_index;
 		}
@@ -3845,7 +3897,8 @@ static void spirv_export_fragment(char *directory, function *main, bool debug) {
 	type_id output_type = output->array_size > 0 ? output->base : pixel_output;
 
 	for (size_t i = 0; i < output_vars_count; ++i) {
-		write_op_variable_preallocated(&instructions, convert_pointer_type_to_spirv_id(output_type, STORAGE_CLASS_OUTPUT), output_vars[i], STORAGE_CLASS_OUTPUT);
+		write_op_variable_preallocated(&instructions, convert_pointer_type_to_spirv_id(output_type, STORAGE_CLASS_OUTPUT), output_vars[i],
+		                               STORAGE_CLASS_OUTPUT);
 	}
 
 	write_functions(&instructions, main, entry_point, SHADER_STAGE_FRAGMENT, NO_TYPE);

base/sources/libs/kong/backends/wgsl.c

@@ -95,25 +95,27 @@ static char *type_string(type_id type) {
 //	return get_name(func);
 // }
 
-static void write_code(char *wgsl, char *directory, const char *filename) {
+static void write_code(char *wgsl, char *directory, const char *filename, const char *name, bool framebuffer_texture_format) {
 	char full_filename[512];
 
 	{
 		sprintf(full_filename, "%s/%s.h", directory, filename);
 		FILE *file = fopen(full_filename, "wb");
+		fprintf(file, "#include <stdbool.h>\n");
 		fprintf(file, "#include <stddef.h>\n\n");
-		fprintf(file, "extern const char *wgsl;\n");
-		fprintf(file, "extern size_t wgsl_size;\n");
+		fprintf(file, "extern const char *%s;\n", name);
+		fprintf(file, "extern size_t %s_size;\n", name);
+		fprintf(file, "extern bool %s_uses_framebuffer_texture_format;\n", name);
 		fclose(file);
 	}
 
 	{
 		sprintf(full_filename, "%s/%s.c", directory, filename);
-		FILE *file = fopen(full_filename, "wb");
 
+		FILE *file = fopen(full_filename, "wb");
 		fprintf(file, "#include \"%s.h\"\n\n", filename);
 
-		fprintf(file, "const char *wgsl = \"");
+		fprintf(file, "const char *%s = \"", name);
 
 		size_t length = strlen(wgsl);
 
@@ -137,7 +139,9 @@ static void write_code(char *wgsl, char *directory, const char *filename) {
 
 		fprintf(file, "\";\n\n");
 
-		fprintf(file, "size_t wgsl_size = %zu;\n\n", length);
+		fprintf(file, "size_t %s_size = %zu;\n\n", name, length);
+
+		fprintf(file, "bool %s_uses_framebuffer_texture_format = %s;\n\n", name, framebuffer_texture_format ? "true" : "false");
 
 		fprintf(file, "/*\n%s*/\n", wgsl);
 
@@ -178,18 +182,51 @@ static bool is_fragment_input(type_id t) {
 	return false;
 }
 
-static void write_types(char *wgsl, size_t *offset) {
-	for (type_id i = 0; get_type(i) != NULL; ++i) {
-		type *t = get_type(i);
+static void add_found_type(type_id t, type_id *types, size_t *types_size) {
+	for (size_t i = 0; i < *types_size; ++i) {
+		if (types[i] == t) {
+			return;
+		}
+	}
+
+	types[*types_size] = t;
+	*types_size += 1;
+}
+
+static void write_types(char *wgsl, size_t *offset, shader_stage stage, type_id inputs[64], size_t inputs_count, type_id output, function *main) {
+	type_id types[256];
+	size_t  types_size = 0;
+	find_referenced_types(main, types, &types_size);
+
+	global_array globals = {0};
+
+	find_referenced_globals(main, &globals);
+
+	for (size_t i = 0; i < globals.size; ++i) {
+		global *g = get_global(globals.globals[i]);
+
+		if (g->type == sampler_type_id) {
+		}
+		else if (get_type(g->type)->tex_kind != TEXTURE_KIND_NONE) {
+		}
+		else if (g->type == float_id) {
+		}
+		else {
+			add_found_type(g->type, types, &types_size);
+		}
+	}
+
+	for (size_t i = 0; i < types_size; ++i) {
+		type *t = get_type(types[i]);
 
 		if (!t->built_in && !has_attribute(&t->attributes, add_name("pipe"))) {
 			if (t->name == NO_NAME) {
 				char name[256];
 
 				bool found = false;
-				for (global_id j = 0; get_global(j)->type != NO_TYPE; ++j) {
+				for (global_id j = 0; get_global(j) != NULL && get_global(j)->type != NO_TYPE; ++j) {
 					global *g = get_global(j);
-					if (g->type == i) {
+					if (g->type == types[i]) {
 						sprintf(name, "_%" PRIu64, g->var_index);
 						found = true;
 						break;
@@ -206,8 +243,8 @@ static void write_types(char *wgsl, size_t *offset) {
 				*offset += sprintf(&wgsl[*offset], "struct %s {\n", get_name(t->name));
 			}
 
-			if (is_vertex_input(i)) {
-				size_t location = find_vertex_location_offset(i);
+			if (is_vertex_input(types[i])) {
+				size_t location = find_vertex_location_offset(types[i]);
 
 				for (size_t j = 0; j < t->members.size; ++j) {
 					*offset +=
@@ -215,7 +252,7 @@ static void write_types(char *wgsl, size_t *offset) {
 					++location;
 				}
 			}
-			else if (is_fragment_input(i)) {
+			else if (is_fragment_input(types[i])) {
 				for (size_t j = 0; j < t->members.size; ++j) {
 					if (j == 0) {
 						*offset +=
@@ -237,146 +274,108 @@ static void write_types(char *wgsl, size_t *offset) {
 	}
 }
 
-static void assign_bindings(uint32_t *bindings) {
-	for (size_t set_index = 0; set_index < get_sets_count(); ++set_index) {
-		uint32_t binding = 0;
-
-		descriptor_set *set = get_set(set_index);
+static void format_to_string(texture_format format, char *str) {
+	switch (format) {
+	case TEXTURE_FORMAT_FRAMEBUFFER:
+		strcpy(str, "$                    ");
+		break;
+	case TEXTURE_FORMAT_RGBA32_FLOAT:
+		strcpy(str, "rgba32float");
+		break;
+	case TEXTURE_FORMAT_RGBA8_UNORM:
+		strcpy(str, "rgba8unorm");
+		break;
+	default:
+		assert(false);
+		break;
+	}
+}
 
-		if (set->name == add_name("root_constants")) {
-			if (set->globals.size != 1) {
-				debug_context context = {0};
-				error(context, "More than one root constants struct found");
-			}
+static void write_globals(char *wgsl, size_t *offset, function *main, bool *framebuffer_format_texture) {
+	global_array referenced_globals = {0};
 
-			global_id g_id = set->globals.globals[0];
-			global   *g    = get_global(g_id);
+	find_referenced_globals(main, &referenced_globals);
 
-			if (get_type(g->type)->built_in) {
-				debug_context context = {0};
-				error(context, "Unsupported type for a root constant");
-			}
+	descriptor_set_group *group = find_descriptor_set_group_for_function(main);
 
-			bindings[g_id] = binding;
-			binding += 1;
+	for (size_t set_index = 0; set_index < group->size; ++set_index) {
+		uint32_t binding = 0;
 
-			continue;
-		}
+		descriptor_set *set = group->values[set_index];
 
 		for (size_t g_index = 0; g_index < set->globals.size; ++g_index) {
 			global_id global_index = set->globals.globals[g_index];
 			bool      writable     = set->globals.writable[g_index];
 
+			bool referenced = false;
+
+			for (size_t i = 0; i < referenced_globals.size; ++i) {
+				if (global_index == referenced_globals.globals[i]) {
+					referenced = true;
+					break;
+				}
+			}
+
 			global *g = get_global(global_index);
 
 			type   *t         = get_type(g->type);
 			type_id base_type = t->array_size > 0 ? t->base : g->type;
 
 			if (base_type == sampler_type_id) {
-				bindings[global_index] = binding;
-				binding += 1;
-			}
-			else if (base_type == tex2d_type_id) {
-				if (t->array_size == UINT32_MAX) {
-					bindings[global_index] = 0;
-				}
-				else if (writable) {
-					bindings[global_index] = binding;
-					binding += 1;
-				}
-				else {
-					bindings[global_index] = binding;
-					binding += 1;
+				if (referenced) {
+					*offset +=
+					    sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": sampler;\n\n", set_index, binding, set_index, g->var_index);
 				}
-			}
-			else if (base_type == texcube_type_id || base_type == tex2darray_type_id || base_type == bvh_type_id) {
-				bindings[global_index] = binding;
 				binding += 1;
 			}
-			else if (get_type(g->type)->built_in) {
-				if (get_type(g->type)->array_size > 0) {
-					bindings[global_index] = binding;
+			else if (get_type(base_type)->tex_kind != TEXTURE_KIND_NONE) {
+				if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D) {
+					if (referenced) {
+						if (t->array_size == UINT32_MAX) {
+							assert(false);
+						}
+						else if (writable) {
+							if (get_type(base_type)->tex_format == TEXTURE_FORMAT_FRAMEBUFFER) {
+								*framebuffer_format_texture = true;
+							}
+
+							char format[64];
+							format_to_string(get_type(base_type)->tex_format, format);
+
+							*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_storage_2d<%s, write>;\n\n",
+							                   set_index, binding, set_index, g->var_index, format);
+						}
+						else {
+							if (is_depth(get_type(base_type)->tex_format)) {
+								*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_depth_2d;\n\n", set_index,
+								                   binding, set_index, g->var_index);
+							}
+							else {
+								*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_2d<f32>;\n\n", set_index, binding,
+								                   set_index, g->var_index);
+							}
+						}
+					}
 					binding += 1;
 				}
-			}
-			else {
-				if (get_type(g->type)->array_size > 0) {
-					bindings[global_index] = binding;
+				else if (get_type(base_type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+					if (referenced) {
+						*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_2d_array<f32>;\n\n", set_index, binding,
+						                   set_index, g->var_index);
+					}
 					binding += 1;
 				}
-				else {
-					bindings[global_index] = binding;
+				else if (get_type(base_type)->tex_kind == TEXTURE_KIND_CUBE) {
+					if (referenced) {
+						*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_cube<f32>;\n\n", set_index, binding,
+						                   set_index, g->var_index);
+					}
 					binding += 1;
 				}
-			}
-		}
-	}
-}
-
-static void write_globals(char *wgsl, size_t *offset) {
-	for (size_t set_index = 0; set_index < get_sets_count(); ++set_index) {
-		uint32_t binding = 0;
-
-		descriptor_set *set = get_set(set_index);
-
-		if (set->name == add_name("root_constants")) {
-			if (set->globals.size != 1) {
-				debug_context context = {0};
-				error(context, "More than one root constants struct found");
-			}
-
-			global_id g_id = set->globals.globals[0];
-			global   *g    = get_global(g_id);
-
-			if (get_type(g->type)->built_in) {
-				debug_context context = {0};
-				error(context, "Unsupported type for a root constant");
-			}
-
-			assert(false);
-
-			binding += 1;
-
-			continue;
-		}
-
-		for (size_t g_index = 0; g_index < set->globals.size; ++g_index) {
-			global_id global_index = set->globals.globals[g_index];
-			bool      writable     = set->globals.writable[g_index];
-
-			global *g = get_global(global_index);
-
-			type   *t         = get_type(g->type);
-			type_id base_type = t->array_size > 0 ? t->base : g->type;
-
-			if (base_type == sampler_type_id) {
-				*offset +=
-				    sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": sampler;\n\n", set_index, binding, set_index, g->var_index);
-				binding += 1;
-			}
-			else if (base_type == tex2d_type_id) {
-				if (t->array_size == UINT32_MAX) {
-					assert(false);
-				}
-				else if (writable) {
-					*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_storage_2d<rgba32float, write>;\n\n",
-					                   set_index, binding, set_index, g->var_index);
-				}
 				else {
-					*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_2d<f32>;\n\n", set_index, binding, set_index,
-					                   g->var_index);
+					// TODO
+					assert(false);
 				}
-				binding += 1;
-			}
-			else if (g->type == tex2darray_type_id) {
-				*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_2d_array<f32>;\n\n", set_index, binding,
-				                   set_index, g->var_index);
-				binding += 1;
-			}
-			else if (g->type == texcube_type_id) {
-				*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var _set%zu_%" PRIu64 ": texture_cube<f32>;\n\n", set_index, binding, set_index,
-				                   g->var_index);
-				binding += 1;
 			}
 			else if (base_type == bvh_type_id) {
 				assert(false);
@@ -394,16 +393,18 @@ static void write_globals(char *wgsl, size_t *offset) {
 					binding += 1;
 				}
 				else {
-					type *t = get_type(g->type);
-					char  type_name[256];
-					if (t->name != NO_NAME) {
-						strcpy(type_name, get_name(t->name));
-					}
-					else {
-						sprintf(type_name, "_%" PRIu64 "_type", g->var_index);
+					if (referenced) {
+						type *t = get_type(g->type);
+						char  type_name[256];
+						if (t->name != NO_NAME) {
+							strcpy(type_name, get_name(t->name));
+						}
+						else {
+							sprintf(type_name, "_%" PRIu64 "_type", g->var_index);
+						}
+						*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var<uniform> _set%zu_%" PRIu64 ": %s;\n\n", set_index, binding, set_index,
+						                   g->var_index, type_name);
 					}
-					*offset += sprintf(&wgsl[*offset], "@group(%zu) @binding(%u) var<uniform> _set%zu_%" PRIu64 ": %s;\n\n", set_index, binding, set_index,
-					                   g->var_index, type_name);
 
 					binding += 1;
 				}
@@ -457,20 +458,8 @@ static small_string get_var(variable var, function *f) {
 		descriptor_set *set = group->values[set_index];
 		for (size_t global_index = 0; global_index < set->globals.size; ++global_index) {
 			if (var.index == get_global(set->globals.globals[global_index])->var_index) {
-				bool found = false;
-
-				size_t global_set_index = 0;
-				for (; global_set_index < get_sets_count(); ++global_set_index) {
-					if (set == get_set(global_set_index)) {
-						found = true;
-						break;
-					}
-				}
-
-				assert(found);
-
 				small_string name;
-				sprintf(name.str, "_set%zu_%" PRIu64, global_set_index, var.index);
+				sprintf(name.str, "_set%zu_%" PRIu64, set_index, var.index);
 				return name;
 			}
 		}
@@ -481,13 +470,21 @@ static small_string get_var(variable var, function *f) {
 	return name;
 }
 
-static void write_functions(char *code, size_t *offset) {
-	for (function_id i = 0; get_function(i) != NULL; ++i) {
-		function *f = get_function(i);
+static void write_functions(char *code, size_t *offset, shader_stage stage, function *main) {
+	function *functions[256];
+	size_t    functions_size = 0;
 
-		if (f->block == NULL) {
-			continue;
-		}
+	functions[functions_size] = main;
+	functions_size += 1;
+
+	find_referenced_functions(main, functions, &functions_size);
+
+	for (size_t i = 0; i < functions_size; ++i) {
+		function *f = functions[i];
+		assert(f != NULL);
+
+		debug_context context = {0};
+		check(f->block != NULL, context, "Function block missing");
 
 		uint8_t *data = f->code.o;
 		size_t   size = f->code.size;
@@ -502,36 +499,13 @@ static void write_functions(char *code, size_t *offset) {
 			}
 		}
 
-		debug_context context = {0};
 		for (uint8_t parameter_index = 0; parameter_index < f->parameters_size; ++parameter_index) {
 			check(parameter_ids[parameter_index] != 0, context, "Parameter not found");
 		}
 
-		if (is_vertex_function(i)) {
-			*offset += sprintf(&code[*offset], "@vertex fn %s(", get_name(f->name));
-			for (uint8_t parameter_index = 0; parameter_index < f->parameters_size; ++parameter_index) {
-				if (parameter_index == 0) {
-					*offset +=
-					    sprintf(&code[*offset], "_%" PRIu64 ": %s", parameter_ids[parameter_index], type_string(f->parameter_types[parameter_index].type));
-				}
-				else {
-					*offset +=
-					    sprintf(&code[*offset], ", _%" PRIu64 ": %s", parameter_ids[parameter_index], type_string(f->parameter_types[parameter_index].type));
-				}
-			}
-			*offset += sprintf(&code[*offset], ") -> %s {\n", type_string(f->return_type.type));
-		}
-		else if (is_fragment_function(i)) {
-			if (get_type(f->return_type.type)->array_size > 0) {
-				type_id base_type = get_type(f->return_type.type)->base;
-
-				*offset += sprintf(&code[*offset], "struct _kong_colors_out {\n");
-				for (uint32_t j = 0; j < get_type(f->return_type.type)->array_size; ++j) {
-					*offset += sprintf(&code[*offset], "\t@location(%u) _%i: %s,\n", j, j, type_string(base_type));
-				}
-				*offset += sprintf(&code[*offset], "}\n\n");
-
-				*offset += sprintf(&code[*offset], "@fragment fn %s(", get_name(f->name));
+		if (f == main) {
+			if (stage == SHADER_STAGE_VERTEX) {
+				*offset += sprintf(&code[*offset], "@vertex fn main(");
 				for (uint8_t parameter_index = 0; parameter_index < f->parameters_size; ++parameter_index) {
 					if (parameter_index == 0) {
 						*offset +=
@@ -542,35 +516,59 @@ static void write_functions(char *code, size_t *offset) {
 						                   type_string(f->parameter_types[parameter_index].type));
 					}
 				}
-				*offset += sprintf(&code[*offset], ") -> _kong_colors_out {\n");
+				*offset += sprintf(&code[*offset], ") -> %s {\n", type_string(f->return_type.type));
 			}
-			else {
-				*offset += sprintf(&code[*offset], "@fragment fn %s(", get_name(f->name));
-				for (uint8_t parameter_index = 0; parameter_index < f->parameters_size; ++parameter_index) {
-					if (parameter_index == 0) {
-						*offset +=
-						    sprintf(&code[*offset], "_%" PRIu64 ": %s", parameter_ids[parameter_index], type_string(f->parameter_types[parameter_index].type));
+			else if (stage == SHADER_STAGE_FRAGMENT) {
+				if (get_type(f->return_type.type)->array_size > 0) {
+					type_id base_type = get_type(f->return_type.type)->base;
+
+					*offset += sprintf(&code[*offset], "struct _kong_colors_out {\n");
+					for (uint32_t j = 0; j < get_type(f->return_type.type)->array_size; ++j) {
+						*offset += sprintf(&code[*offset], "\t@location(%u) _%i: %s,\n", j, j, type_string(base_type));
 					}
-					else {
-						*offset += sprintf(&code[*offset], ", _%" PRIu64 ": %s", parameter_ids[parameter_index],
-						                   type_string(f->parameter_types[parameter_index].type));
+					*offset += sprintf(&code[*offset], "}\n\n");
+
+					*offset += sprintf(&code[*offset], "@fragment fn main(");
+					for (uint8_t parameter_index = 0; parameter_index < f->parameters_size; ++parameter_index) {
+						if (parameter_index == 0) {
+							*offset += sprintf(&code[*offset], "_%" PRIu64 ": %s", parameter_ids[parameter_index],
+							                   type_string(f->parameter_types[parameter_index].type));
+						}
+						else {
+							*offset += sprintf(&code[*offset], ", _%" PRIu64 ": %s", parameter_ids[parameter_index],
+							                   type_string(f->parameter_types[parameter_index].type));
+						}
+					}
+					*offset += sprintf(&code[*offset], ") -> _kong_colors_out {\n");
+				}
+				else {
+					*offset += sprintf(&code[*offset], "@fragment fn main(");
+					for (uint8_t parameter_index = 0; parameter_index < f->parameters_size; ++parameter_index) {
+						if (parameter_index == 0) {
+							*offset += sprintf(&code[*offset], "_%" PRIu64 ": %s", parameter_ids[parameter_index],
+							                   type_string(f->parameter_types[parameter_index].type));
+						}
+						else {
+							*offset += sprintf(&code[*offset], ", _%" PRIu64 ": %s", parameter_ids[parameter_index],
+							                   type_string(f->parameter_types[parameter_index].type));
+						}
 					}
+					*offset += sprintf(&code[*offset], ") -> @location(0) %s {\n", type_string(f->return_type.type));
 				}
-				*offset += sprintf(&code[*offset], ") -> @location(0) %s {\n", type_string(f->return_type.type));
 			}
-		}
-		else if (is_compute_function(i)) {
-			assert(f->parameters_size == 0);
-			assert(f->return_type.type == void_id);
-
-			attribute *threads = find_attribute(&f->attributes, add_name("threads"));
-			assert(threads != NULL && threads->paramters_count == 3);
-
-			*offset += sprintf(&code[*offset],
-			                   "@compute @workgroup_size(%u, %u, %u) fn %s(@builtin(local_invocation_id) _kong_group_thread_id: vec3<u32>, "
-			                   "@builtin(workgroup_id) _kong_group_id: vec3<u32>, @builtin(global_invocation_id) _kong_dispatch_thread_id: vec3<u32>, "
-			                   "@builtin(num_workgroups) _kong_threads_count: vec3<u32>, @builtin(local_invocation_index) _kong_group_index: u32) {\n",
-			                   (uint32_t)threads->parameters[0], (uint32_t)threads->parameters[1], (uint32_t)threads->parameters[2], get_name(f->name));
+			else if (stage == SHADER_STAGE_COMPUTE) {
+				assert(f->parameters_size == 0);
+				assert(f->return_type.type == void_id);
+
+				attribute *threads = find_attribute(&f->attributes, add_name("threads"));
+				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>, "
+				                   "@builtin(workgroup_id) _kong_group_id: vec3<u32>, @builtin(global_invocation_id) _kong_dispatch_thread_id: vec3<u32>, "
+				                   "@builtin(num_workgroups) _kong_threads_count: vec3<u32>, @builtin(local_invocation_index) _kong_group_index: u32) {\n",
+				                   (uint32_t)threads->parameters[0], (uint32_t)threads->parameters[1], (uint32_t)threads->parameters[2]);
+			}
 		}
 		else {
 			*offset += sprintf(&code[*offset], "fn %s(", get_name(f->name));
@@ -609,16 +607,15 @@ static void write_functions(char *code, size_t *offset) {
 				type_id from_type = o->op_load_access_list.from.type.type;
 
 				if (is_texture(from_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);
+					assert(o->op_load_access_list.access_list_size == 1);
+					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).str, type_string(o->op_load_access_list.to.type.type));
 
-					*offset += sprintf(&code[*offset], "textureLoad(%s, vec2<u32>(u32(%s.x), u32(%s.y)), 0);\n", get_var(o->op_store_access_list.from, f).str,
-					                   get_var(o->op_store_access_list.access_list[0].access_element.index, f).str,
-					                   get_var(o->op_store_access_list.access_list[0].access_element.index, f).str);
+					*offset += sprintf(&code[*offset], "textureLoad(%s, vec2<u32>(u32(%s.x), u32(%s.y)), 0);\n", get_var(o->op_load_access_list.from, f).str,
+					                   get_var(o->op_load_access_list.access_list[0].access_element.index, f).str,
+					                   get_var(o->op_load_access_list.access_list[0].access_element.index, f).str);
 				}
 				else {
 					*offset += sprintf(&code[*offset], "var %s: %s = %s", get_var(o->op_load_access_list.to, f).str,
@@ -730,7 +727,7 @@ static void write_functions(char *code, size_t *offset) {
 			}
 			case OPCODE_RETURN: {
 				if (o->size > offsetof(opcode, op_return)) {
-					if (is_fragment_function(i) && get_type(f->return_type.type)->array_size > 0) {
+					if (f == main && stage == SHADER_STAGE_FRAGMENT && get_type(f->return_type.type)->array_size > 0) {
 						indent(code, offset, indentation);
 						*offset += sprintf(&code[*offset], "{\n");
 						indent(code, offset, indentation + 1);
@@ -779,6 +776,90 @@ static void write_functions(char *code, size_t *offset) {
 				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
 				break;
 			}
+			case OPCODE_MOD: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " %% _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_EQUALS: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " == _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_NOT_EQUALS: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " != _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_GREATER: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " > _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_GREATER_EQUAL: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " >= _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_LESS: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " < _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_LESS_EQUAL: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " <= _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_AND: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " && _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_OR: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " || _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_BITWISE_XOR: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " ^ _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_BITWISE_AND: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " & _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_BITWISE_OR: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " | _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_LEFT_SHIFT: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " << _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
+			case OPCODE_RIGHT_SHIFT: {
+				indent(code, offset, indentation);
+				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = _%" PRIu64 " >> _%" PRIu64 ";\n", o->op_binary.result.index,
+				                   type_string(o->op_binary.result.type.type), o->op_binary.left.index, o->op_binary.right.index);
+				break;
+			}
 			case OPCODE_LOAD_FLOAT_CONSTANT:
 				indent(code, offset, indentation);
 				*offset += sprintf(&code[*offset], "var _%" PRIu64 ": %s = %f;\n", o->op_load_float_constant.to.index,
@@ -804,7 +885,7 @@ static void write_functions(char *code, size_t *offset) {
 					variable sampler = o->op_call.parameters[1];
 					variable coord   = o->op_call.parameters[2];
 
-					if (tex.type.type == tex2darray_type_id) {
+					if (get_type(tex.type.type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
 						*offset += sprintf(&code[*offset], "var %s: %s = textureSample(%s, %s, %s.xy, u32(%s.z));\n", get_var(o->op_call.var, f).str,
 						                   type_string(o->op_call.var.type.type), get_var(tex, f).str, get_var(sampler, f).str, get_var(coord, f).str,
 						                   get_var(coord, f).str);
@@ -909,19 +990,92 @@ static void write_functions(char *code, size_t *offset) {
 	}
 }
 
-static void wgsl_export_everything(char *directory) {
+static void wgsl_export_vertex(char *directory, function *main) {
+	char         *wgsl    = (char *)calloc(1024 * 1024, 1);
+	debug_context context = {0};
+	check(wgsl != NULL, context, "Could not allocate the wgsl string");
+
+	size_t offset = 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;
+	}
+	type_id vertex_output = main->return_type.type;
+
+	check(vertex_output != NO_TYPE, context, "vertex output missing");
+
+	write_types(wgsl, &offset, SHADER_STAGE_VERTEX, vertex_inputs, main->parameters_size, vertex_output, main);
+
+	bool framebuffer_texture_format = false;
+	write_globals(wgsl, &offset, main, &framebuffer_texture_format);
+
+	write_functions(wgsl, &offset, SHADER_STAGE_VERTEX, main);
+
+	char *name = get_name(main->name);
+
+	char filename[512];
+	sprintf(filename, "kong_%s", name);
+
+	char var_name[256];
+	sprintf(var_name, "%s_code", name);
+
+	write_code(wgsl, directory, filename, var_name, framebuffer_texture_format);
+}
+
+static void wgsl_export_fragment(char *directory, function *main) {
+	char         *wgsl    = (char *)calloc(1024 * 1024, 1);
+	debug_context context = {0};
+	check(wgsl != NULL, context, "Could not allocate the wgsl string");
+
+	size_t offset = 0;
+
+	assert(main->parameters_size > 0);
+	type_id pixel_input = main->parameter_types[0].type;
+
+	check(pixel_input != NO_TYPE, context, "fragment input missing");
+
+	write_types(wgsl, &offset, SHADER_STAGE_FRAGMENT, &pixel_input, 1, NO_TYPE, main);
+
+	bool framebuffer_texture_format = false;
+	write_globals(wgsl, &offset, main, &framebuffer_texture_format);
+
+	write_functions(wgsl, &offset, SHADER_STAGE_FRAGMENT, main);
+
+	char *name = get_name(main->name);
+
+	char filename[512];
+	sprintf(filename, "kong_%s", name);
+
+	char var_name[256];
+	sprintf(var_name, "%s_code", name);
+
+	write_code(wgsl, directory, filename, var_name, framebuffer_texture_format);
+}
+
+static void wgsl_export_compute(char *directory, function *main) {
 	char         *wgsl    = (char *)calloc(1024 * 1024, 1);
 	debug_context context = {0};
 	check(wgsl != NULL, context, "Could not allocate the wgsl string");
+
 	size_t offset = 0;
 
-	write_types(wgsl, &offset);
+	write_types(wgsl, &offset, SHADER_STAGE_COMPUTE, NULL, 0, NO_TYPE, main);
+
+	bool framebuffer_texture_format = false;
+	write_globals(wgsl, &offset, main, &framebuffer_texture_format);
+
+	write_functions(wgsl, &offset, SHADER_STAGE_COMPUTE, main);
+
+	char *name = get_name(main->name);
 
-	write_globals(wgsl, &offset);
+	char filename[512];
+	sprintf(filename, "kong_%s", name);
 
-	write_functions(wgsl, &offset);
+	char var_name[256];
+	sprintf(var_name, "%s_code", name);
 
-	write_code(wgsl, directory, "wgsl");
+	write_code(wgsl, directory, filename, var_name, framebuffer_texture_format);
 }
 
 void wgsl_export(char *directory) {
@@ -980,5 +1134,15 @@ void wgsl_export(char *directory) {
 		}
 	}
 
-	wgsl_export_everything(directory);
+	for (size_t i = 0; i < vertex_functions_size; ++i) {
+		wgsl_export_vertex(directory, get_function(vertex_functions[i]));
+	}
+
+	for (size_t i = 0; i < fragment_functions_size; ++i) {
+		wgsl_export_fragment(directory, get_function(fragment_functions[i]));
+	}
+
+	for (size_t i = 0; i < compute_functions_size; ++i) {
+		wgsl_export_compute(directory, get_function(compute_functions[i]));
+	}
 }

base/sources/libs/kong/globals.c

@@ -151,13 +151,13 @@ void globals_init(void) {
 	uint_value.value.uints[0] = 37;
 	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH16_UNORM"), uint_value);
 	uint_value.value.uints[0] = 38;
-	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH24PLUS_NOTHING8"), uint_value);
+	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH24_NOTHING8"), uint_value);
 	uint_value.value.uints[0] = 39;
-	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH24PLUS_STENCIL8"), uint_value);
+	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH24_STENCIL8"), uint_value);
 	uint_value.value.uints[0] = 40;
-	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH32FLOAT"), uint_value);
+	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH32_FLOAT"), uint_value);
 	uint_value.value.uints[0] = 41;
-	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH32FLOAT_STENCIL8_NOTHING24"), uint_value);
+	add_global_with_value(uint_id, attributes, add_name("TEXTURE_FORMAT_DEPTH32_FLOAT_STENCIL8_NOTHING24"), uint_value);
 }
 
 global_id add_global(type_id type, attribute_list attributes, name_id name) {

base/sources/libs/kong/globals.h

@@ -37,6 +37,7 @@ typedef enum global_usage {
 	GLOBAL_USAGE_TEXTURE_READ   = 0x00000002,
 	GLOBAL_USAGE_TEXTURE_WRITE  = 0x00000004,
 	GLOBAL_USAGE_BUFFER_WRITE   = 0x00000008,
+	GLOBAL_USAGE_SAMPLE_DEPTH   = 0x00000010,
 } global_usage;
 
 typedef struct global {

base/sources/libs/kong/integrations/kore3.c

@@ -256,13 +256,13 @@ static const char *convert_texture_format(int format) {
 	case 37:
 		return "KORE_GPU_TEXTURE_FORMAT_DEPTH16_UNORM";
 	case 38:
-		return "KORE_GPU_TEXTURE_FORMAT_DEPTH24PLUS_NOTHING8";
+		return "KORE_GPU_TEXTURE_FORMAT_DEPTH24_NOTHING8";
 	case 39:
-		return "KORE_GPU_TEXTURE_FORMAT_DEPTH24PLUS_STENCIL8";
+		return "KORE_GPU_TEXTURE_FORMAT_DEPTH24_STENCIL8";
 	case 40:
-		return "KORE_GPU_TEXTURE_FORMAT_DEPTH32FLOAT";
+		return "KORE_GPU_TEXTURE_FORMAT_DEPTH32_FLOAT";
 	case 41:
-		return "KORE_GPU_TEXTURE_FORMAT_DEPTH32FLOAT_STENCIL8_NOTHING24";
+		return "KORE_GPU_TEXTURE_FORMAT_DEPTH32_FLOAT_STENCIL8_NOTHING24";
 	default: {
 		debug_context context = {0};
 		error(context, "Unknown texture format");
@@ -552,6 +552,27 @@ static void to_upper(char *from, char *to) {
 	to[from_size] = 0;
 }
 
+static void format_to_string_wgsl(texture_format format, char *str) {
+	switch (format) {
+	case TEXTURE_FORMAT_FRAMEBUFFER:
+		strcpy(str, "RGBA8Unorm");
+		break;
+	case TEXTURE_FORMAT_RGBA32_FLOAT:
+		strcpy(str, "RGBA32Float");
+		break;
+	case TEXTURE_FORMAT_RGBA8_UNORM:
+	case TEXTURE_FORMAT_UNDEFINED:
+		strcpy(str, "RGBA8Unorm");
+		break;
+	case TEXTURE_FORMAT_DEPTH:
+		strcpy(str, "Depth");
+		break;
+	default:
+		assert(false);
+		break;
+	}
+}
+
 static int global_register_indices[512];
 
 void kore3_export(char *directory, api_kind api) {
@@ -1007,9 +1028,6 @@ void kore3_export(char *directory, api_kind api) {
 		if (api == API_METAL) {
 			// Code is added directly to the Xcode project instead
 		}
-		else if (api == API_WEBGPU) {
-			fprintf(output, "#include \"wgsl.h\"\n");
-		}
 		else {
 			for (type_id i = 0; get_type(i) != NULL; ++i) {
 				type *t = get_type(i);
@@ -1090,23 +1108,25 @@ void kore3_export(char *directory, api_kind api) {
 			fprintf(output, "}\n\n");
 		}
 
-		if (api != API_WEBGPU) {
-			for (size_t set_index = 0; set_index < sets_count; ++set_index) {
-				descriptor_set *set = sets[set_index];
-				if (api == API_METAL) {
-					fprintf(output, "static uint32_t %s_vertex_table_index = UINT32_MAX;\n\n", get_name(set->name));
-					fprintf(output, "static uint32_t %s_fragment_table_index = UINT32_MAX;\n\n", get_name(set->name));
-					fprintf(output, "static uint32_t %s_compute_table_index = UINT32_MAX;\n\n", get_name(set->name));
-				}
-				else if (api == API_VULKAN) {
-					if (set->name != add_name("root_constants")) {
-						fprintf(output, "static uint32_t %s_table_index = UINT32_MAX;\n\n", get_name(set->name));
-					}
-				}
-				else {
+		if (api == API_WEBGPU) {
+			fprintf(output, "static uint32_t root_constants_table_index = UINT32_MAX;\n\n");
+		}
+
+		for (size_t set_index = 0; set_index < sets_count; ++set_index) {
+			descriptor_set *set = sets[set_index];
+			if (api == API_METAL) {
+				fprintf(output, "static uint32_t %s_vertex_table_index = UINT32_MAX;\n\n", get_name(set->name));
+				fprintf(output, "static uint32_t %s_fragment_table_index = UINT32_MAX;\n\n", get_name(set->name));
+				fprintf(output, "static uint32_t %s_compute_table_index = UINT32_MAX;\n\n", get_name(set->name));
+			}
+			else if (api == API_VULKAN || api == API_WEBGPU) {
+				if (set->name != add_name("root_constants")) {
 					fprintf(output, "static uint32_t %s_table_index = UINT32_MAX;\n\n", get_name(set->name));
 				}
 			}
+			else {
+				fprintf(output, "static uint32_t %s_table_index = UINT32_MAX;\n\n", get_name(set->name));
+			}
 		}
 
 		for (type_id i = 0; get_type(i) != NULL; ++i) {
@@ -1203,28 +1223,26 @@ void kore3_export(char *directory, api_kind api) {
 					}
 				}
 
-				if (api != API_WEBGPU) {
-					descriptor_set_group *group = find_descriptor_set_group_for_pipe_type(t);
+				descriptor_set_group *group = find_descriptor_set_group_for_pipe_type(t);
 
-					if (api == API_VULKAN) {
-						size_t index = 0;
-						for (size_t group_index = 0; group_index < group->size; ++group_index) {
-							if (group->values[group_index]->name != add_name("root_constants")) {
-								fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), index);
-								index += 1;
-							}
+				if (api == API_VULKAN) {
+					size_t index = 0;
+					for (size_t group_index = 0; group_index < group->size; ++group_index) {
+						if (group->values[group_index]->name != add_name("root_constants")) {
+							fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), index);
+							index += 1;
 						}
 					}
-					else {
-						for (size_t group_index = 0; group_index < group->size; ++group_index) {
-							if (api == API_METAL) {
-								fprintf(output, "\t%s_vertex_table_index = %zu;\n", get_name(group->values[group_index]->name),
-								        group_index + vertex_function->parameters_size);
-								fprintf(output, "\t%s_fragment_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index + 1);
-							}
-							else {
-								fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index);
-							}
+				}
+				else {
+					for (size_t group_index = 0; group_index < group->size; ++group_index) {
+						if (api == API_METAL) {
+							fprintf(output, "\t%s_vertex_table_index = %zu;\n", get_name(group->values[group_index]->name),
+							        group_index + vertex_function->parameters_size);
+							fprintf(output, "\t%s_fragment_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index + 1);
+						}
+						else {
+							fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index);
 						}
 					}
 				}
@@ -1273,7 +1291,7 @@ void kore3_export(char *directory, api_kind api) {
 						fprintf(output, "\tusage |= KORE_D3D12_TEXTURE_USAGE_UAV;\n");
 					}
 				}
-				else if (api == API_VULKAN || api == API_WEBGPU || api == API_OPENGL) {
+				else if (api == API_VULKAN || api == API_OPENGL) {
 					if (global_has_usage(i, GLOBAL_USAGE_TEXTURE_SAMPLE)) {
 						fprintf(output, "\tusage |= KORE_%s_TEXTURE_USAGE_SAMPLED;\n", api_caps);
 					}
@@ -1281,6 +1299,14 @@ void kore3_export(char *directory, api_kind api) {
 						fprintf(output, "\tusage |= KORE_%s_TEXTURE_USAGE_STORAGE;\n", api_caps);
 					}
 				}
+				else if (api == API_WEBGPU) {
+					if (global_has_usage(i, GLOBAL_USAGE_TEXTURE_SAMPLE) || global_has_usage(i, GLOBAL_USAGE_TEXTURE_READ)) {
+						fprintf(output, "\tusage |= KORE_%s_TEXTURE_USAGE_SAMPLED;\n", api_caps);
+					}
+					if (global_has_usage(i, GLOBAL_USAGE_TEXTURE_WRITE)) {
+						fprintf(output, "\tusage |= KORE_%s_TEXTURE_USAGE_STORAGE;\n", api_caps);
+					}
+				}
 				else {
 					if (global_has_usage(i, GLOBAL_USAGE_TEXTURE_SAMPLE)) {
 						fprintf(output, "\tusage |= KORE_%s_TEXTURE_USAGE_SAMPLE;\n", api_caps);
@@ -1428,6 +1454,10 @@ void kore3_export(char *directory, api_kind api) {
 			}
 		}
 
+		if (api == API_WEBGPU) {
+			fprintf(output, "extern WGPUBindGroupLayout root_constants_set_layout;\n\n");
+		}
+
 		for (size_t set_index = 0; set_index < sets_count; ++set_index) {
 			descriptor_set *set = sets[set_index];
 
@@ -1528,61 +1558,68 @@ void kore3_export(char *directory, api_kind api) {
 						fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
 						other_index += 1;
 					}
-					else if (base_type_id == tex2d_type_id) {
-						type *t = get_type(g->type);
-						if (t->array_size == UINT32_MAX) {
-							fprintf(output, "\tset->%s = (kore_gpu_texture_view *)malloc(sizeof(kore_gpu_texture_view) * parameters->textures_count);\n",
-							        get_name(g->name));
-							fprintf(output, "\tassert(set->%s != NULL);\n", get_name(g->name));
-							fprintf(output, "\tfor (size_t index = 0; index < parameters->textures_count; ++index) {\n");
-							fprintf(output,
-							        "\t\tkore_%s_descriptor_set_set_texture_view_srv(device, set->set.bindless_descriptor_allocation.offset + (uint32_t)index, "
-							        "&parameters->%s[index]);\n",
-							        api_short, get_name(g->name));
-							fprintf(output, "\t\tset->%s[index] = parameters->%s[index];\n", get_name(g->name), get_name(g->name));
-							fprintf(output, "\t}\n");
+					else if (get_type(base_type_id)->tex_kind != TEXTURE_KIND_NONE) {
+						if (get_type(base_type_id)->tex_kind == TEXTURE_KIND_2D) {
+							type *t = get_type(g->type);
+							if (t->array_size == UINT32_MAX) {
+								fprintf(output, "\tset->%s = (kore_gpu_texture_view *)malloc(sizeof(kore_gpu_texture_view) * parameters->textures_count);\n",
+								        get_name(g->name));
+								fprintf(output, "\tassert(set->%s != NULL);\n", get_name(g->name));
+								fprintf(output, "\tfor (size_t index = 0; index < parameters->textures_count; ++index) {\n");
+								fprintf(
+								    output,
+								    "\t\tkore_%s_descriptor_set_set_texture_view_srv(device, set->set.bindless_descriptor_allocation.offset + (uint32_t)index, "
+								    "&parameters->%s[index]);\n",
+								    api_short, get_name(g->name));
+								fprintf(output, "\t\tset->%s[index] = parameters->%s[index];\n", get_name(g->name), get_name(g->name));
+								fprintf(output, "\t}\n");
 
-							fprintf(output, "\tset->%s_count = parameters->%s_count;\n", get_name(g->name), get_name(g->name));
-						}
-						else {
-							if (writable) {
-								fprintf(output, "\tkore_%s_descriptor_set_set_texture_view_uav(device, &set->set, &parameters->%s, %zu);\n", api_short,
-								        get_name(g->name), other_index);
+								fprintf(output, "\tset->%s_count = parameters->%s_count;\n", get_name(g->name), get_name(g->name));
 							}
 							else {
-								fprintf(output,
-								        "\tkore_%s_descriptor_set_set_texture_view_srv(device, set->set.descriptor_allocation.offset + %zu, "
-								        "&parameters->%s);\n",
-								        api_short, other_index, get_name(g->name));
-							}
+								if (writable) {
+									fprintf(output, "\tkore_%s_descriptor_set_set_texture_view_uav(device, &set->set, &parameters->%s, %zu);\n", api_short,
+									        get_name(g->name), other_index);
+								}
+								else {
+									fprintf(output,
+									        "\tkore_%s_descriptor_set_set_texture_view_srv(device, set->set.descriptor_allocation.offset + %zu, "
+									        "&parameters->%s);\n",
+									        api_short, other_index, get_name(g->name));
+								}
 
-							fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
+								fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
 
-							other_index += 1;
-						}
-					}
-					else if (base_type_id == tex2darray_type_id) {
-						if (writable) {
-							debug_context context = {0};
-							error(context, "Texture arrays can not be writable");
+								other_index += 1;
+							}
 						}
+						else if (get_type(base_type_id)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+							if (writable) {
+								debug_context context = {0};
+								error(context, "Texture arrays can not be writable");
+							}
 
-						fprintf(output, "\tkore_%s_descriptor_set_set_texture_array_view_srv(device, &set->set, &parameters->%s, %zu);\n", api_short,
-						        get_name(g->name), other_index);
+							fprintf(output, "\tkore_%s_descriptor_set_set_texture_array_view_srv(device, &set->set, &parameters->%s, %zu);\n", api_short,
+							        get_name(g->name), other_index);
 
-						fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
-						other_index += 1;
-					}
-					else if (base_type_id == texcube_type_id) {
-						if (writable) {
-							debug_context context = {0};
-							error(context, "Cube maps can not be writable");
+							fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
+							other_index += 1;
 						}
-						fprintf(output, "\tkore_%s_descriptor_set_set_texture_cube_view_srv(device, &set->set, &parameters->%s, %zu);\n", api_short,
-						        get_name(g->name), other_index);
+						else if (get_type(base_type_id)->tex_kind == TEXTURE_KIND_CUBE) {
+							if (writable) {
+								debug_context context = {0};
+								error(context, "Cube maps can not be writable");
+							}
+							fprintf(output, "\tkore_%s_descriptor_set_set_texture_cube_view_srv(device, &set->set, &parameters->%s, %zu);\n", api_short,
+							        get_name(g->name), other_index);
 
-						fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
-						other_index += 1;
+							fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
+							other_index += 1;
+						}
+						else {
+							// TODO
+							assert(false);
+						}
 					}
 					else if (is_sampler(g->type)) {
 						fprintf(output, "\tkore_%s_descriptor_set_set_sampler(device, &set->set, parameters->%s, %zu);\n", api_short, get_name(g->name),
@@ -1666,26 +1703,33 @@ void kore3_export(char *directory, api_kind api) {
 					else if (is_texture(g->type)) {
 						fprintf(output, "\t\tid<MTLTexture> texture = (__bridge id<MTLTexture>)parameters->%s.texture->metal.texture;\n", get_name(g->name));
 
-						if (g->type == tex2d_type_id) {
-							fprintf(output,
-							        "\t\tid<MTLTexture> view = [texture newTextureViewWithPixelFormat:texture.pixelFormat textureType:MTLTextureType2D "
-							        "levels:NSMakeRange(parameters->%s.base_mip_level, parameters->%s.mip_level_count) "
-							        "slices:NSMakeRange(parameters->%s.base_array_layer, parameters->%s.array_layer_count)];\n",
-							        get_name(g->name), get_name(g->name), get_name(g->name), get_name(g->name));
-						}
-						else if (g->type == tex2darray_type_id) {
-							fprintf(output,
-							        "\t\tid<MTLTexture> view = [texture newTextureViewWithPixelFormat:texture.pixelFormat textureType:MTLTextureType2DArray "
-							        "levels:NSMakeRange(parameters->%s.base_mip_level, parameters->%s.mip_level_count) "
-							        "slices:NSMakeRange(parameters->%s.base_array_layer, parameters->%s.array_layer_count)];\n",
-							        get_name(g->name), get_name(g->name), get_name(g->name), get_name(g->name));
-						}
-						else if (g->type == texcube_type_id) {
-							fprintf(output,
-							        "\t\tid<MTLTexture> view = [texture newTextureViewWithPixelFormat:texture.pixelFormat textureType:MTLTextureTypeCube "
-							        "levels:NSMakeRange(parameters->%s.base_mip_level, parameters->%s.mip_level_count) "
-							        "slices:NSMakeRange(parameters->%s.base_array_layer, parameters->%s.array_layer_count)];\n",
-							        get_name(g->name), get_name(g->name), get_name(g->name), get_name(g->name));
+						if (get_type(g->type)->tex_kind != TEXTURE_KIND_NONE) {
+							if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D) {
+								fprintf(output,
+								        "\t\tid<MTLTexture> view = [texture newTextureViewWithPixelFormat:texture.pixelFormat textureType:MTLTextureType2D "
+								        "levels:NSMakeRange(parameters->%s.base_mip_level, parameters->%s.mip_level_count) "
+								        "slices:NSMakeRange(parameters->%s.base_array_layer, parameters->%s.array_layer_count)];\n",
+								        get_name(g->name), get_name(g->name), get_name(g->name), get_name(g->name));
+							}
+							else if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+								fprintf(
+								    output,
+								    "\t\tid<MTLTexture> view = [texture newTextureViewWithPixelFormat:texture.pixelFormat textureType:MTLTextureType2DArray "
+								    "levels:NSMakeRange(parameters->%s.base_mip_level, parameters->%s.mip_level_count) "
+								    "slices:NSMakeRange(parameters->%s.base_array_layer, parameters->%s.array_layer_count)];\n",
+								    get_name(g->name), get_name(g->name), get_name(g->name), get_name(g->name));
+							}
+							else if (get_type(g->type)->tex_kind == TEXTURE_KIND_CUBE) {
+								fprintf(output,
+								        "\t\tid<MTLTexture> view = [texture newTextureViewWithPixelFormat:texture.pixelFormat textureType:MTLTextureTypeCube "
+								        "levels:NSMakeRange(parameters->%s.base_mip_level, parameters->%s.mip_level_count) "
+								        "slices:NSMakeRange(parameters->%s.base_array_layer, parameters->%s.array_layer_count)];\n",
+								        get_name(g->name), get_name(g->name), get_name(g->name), get_name(g->name));
+							}
+							else {
+								// TODO
+								assert(false);
+							}
 						}
 
 						fprintf(output, "\t\t[argument_encoder setTexture: view atIndex: %zu];\n", index);
@@ -1771,59 +1815,62 @@ void kore3_export(char *directory, api_kind api) {
 						fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
 						other_index += 1;
 					}
-					else if (base_type_id == tex2d_type_id) {
-						type *t = get_type(g->type);
-						if (t->array_size == UINT32_MAX) {
-							fprintf(output, "\tset->%s = (kore_gpu_texture_view *)malloc(sizeof(kore_gpu_texture_view) * parameters->textures_count);\n",
-							        get_name(g->name));
-							fprintf(output, "\tassert(set->%s != NULL);\n", get_name(g->name));
-							fprintf(output, "\tfor (size_t index = 0; index < parameters->textures_count; ++index) {\n");
-							fprintf(output,
-							        "\t\tkore_%s_descriptor_set_set_texture_view_srv(device, set->set.bindless_descriptor_allocation.offset + (uint32_t)index, "
-							        "&parameters->%s[index]);\n",
-							        api_short, get_name(g->name));
-							fprintf(output, "\t\tset->%s[index] = parameters->%s[index];\n", get_name(g->name), get_name(g->name));
-							fprintf(output, "\t}\n");
+					else if (get_type(base_type_id)->tex_kind != TEXTURE_KIND_NONE) {
+						if (get_type(base_type_id)->tex_kind == TEXTURE_KIND_2D) {
+							type *t = get_type(g->type);
+							if (t->array_size == UINT32_MAX) {
+								fprintf(output, "\tset->%s = (kore_gpu_texture_view *)malloc(sizeof(kore_gpu_texture_view) * parameters->textures_count);\n",
+								        get_name(g->name));
+								fprintf(output, "\tassert(set->%s != NULL);\n", get_name(g->name));
+								fprintf(output, "\tfor (size_t index = 0; index < parameters->textures_count; ++index) {\n");
+								fprintf(
+								    output,
+								    "\t\tkore_%s_descriptor_set_set_texture_view_srv(device, set->set.bindless_descriptor_allocation.offset + (uint32_t)index, "
+								    "&parameters->%s[index]);\n",
+								    api_short, get_name(g->name));
+								fprintf(output, "\t\tset->%s[index] = parameters->%s[index];\n", get_name(g->name), get_name(g->name));
+								fprintf(output, "\t}\n");
 
-							fprintf(output, "\tset->%s_count = parameters->%s_count;\n", get_name(g->name), get_name(g->name));
-						}
-						else {
-							if (readable | writable) {
-								fprintf(output, "\tkore_vulkan_descriptor_set_set_storage_image_descriptor(device, &set->set, &parameters->%s, %zu);\n",
-								        get_name(g->name), other_index);
+								fprintf(output, "\tset->%s_count = parameters->%s_count;\n", get_name(g->name), get_name(g->name));
 							}
 							else {
-								fprintf(output, "\tkore_vulkan_descriptor_set_set_sampled_image_descriptor(device, &set->set, &parameters->%s, %zu);\n",
-								        get_name(g->name), other_index);
-							}
+								if (readable | writable) {
+									fprintf(output, "\tkore_vulkan_descriptor_set_set_storage_image_descriptor(device, &set->set, &parameters->%s, %zu);\n",
+									        get_name(g->name), other_index);
+								}
+								else {
+									fprintf(output, "\tkore_vulkan_descriptor_set_set_sampled_image_descriptor(device, &set->set, &parameters->%s, %zu);\n",
+									        get_name(g->name), other_index);
+								}
 
-							fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
+								fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
 
-							other_index += 1;
-						}
-					}
-					else if (base_type_id == tex2darray_type_id) {
-						if (writable) {
-							debug_context context = {0};
-							error(context, "Texture arrays can not be writable");
+								other_index += 1;
+							}
 						}
+						else if (get_type(base_type_id)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+							if (writable) {
+								debug_context context = {0};
+								error(context, "Texture arrays can not be writable");
+							}
 
-						fprintf(output, "\tkore_vulkan_descriptor_set_set_sampled_image_array_descriptor(device, &set->set, &parameters->%s, %zu);\n",
-						        get_name(g->name), other_index);
+							fprintf(output, "\tkore_vulkan_descriptor_set_set_sampled_image_array_descriptor(device, &set->set, &parameters->%s, %zu);\n",
+							        get_name(g->name), other_index);
 
-						fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
-						other_index += 1;
-					}
-					else if (base_type_id == texcube_type_id) {
-						if (writable) {
-							debug_context context = {0};
-							error(context, "Cube maps can not be writable");
+							fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
+							other_index += 1;
 						}
-						fprintf(output, "\tkore_vulkan_descriptor_set_set_sampled_cube_image_descriptor(device, &set->set, &parameters->%s, %zu);\n",
-						        get_name(g->name), other_index);
+						else if (get_type(base_type_id)->tex_kind == TEXTURE_KIND_CUBE) {
+							if (writable) {
+								debug_context context = {0};
+								error(context, "Cube maps can not be writable");
+							}
+							fprintf(output, "\tkore_vulkan_descriptor_set_set_sampled_cube_image_descriptor(device, &set->set, &parameters->%s, %zu);\n",
+							        get_name(g->name), other_index);
 
-						fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
-						other_index += 1;
+							fprintf(output, "\tset->%s = parameters->%s;\n", get_name(g->name), get_name(g->name));
+							other_index += 1;
+						}
 					}
 					else if (is_sampler(g->type)) {
 						fprintf(output, "\tkore_%s_descriptor_set_set_sampler(device, &set->set, parameters->%s, %zu);\n", api_short, get_name(g->name),
@@ -1851,15 +1898,22 @@ void kore3_export(char *directory, api_kind api) {
 
 					if (is_texture(g->type)) {
 						fprintf(output, "\tWGPUTextureViewDescriptor texture_view_descriptor%zu = {\n", global_index);
-						fprintf(output, "\t\t.format = kore_webgpu_convert_texture_format(parameters->%s.texture->webgpu.format),\n", get_name(g->name));
-						if (g->type == tex2darray_type_id) {
-							fprintf(output, "\t\t.dimension = WGPUTextureViewDimension_2DArray,\n");
-						}
-						else if (g->type == texcube_type_id) {
-							fprintf(output, "\t\t.dimension = WGPUTextureViewDimension_Cube,\n");
-						}
-						else {
-							fprintf(output, "\t\t.dimension = WGPUTextureViewDimension_2D,\n");
+						fprintf(output, "\t\t.format = kore_webgpu_convert_texture_format_to_webgpu(parameters->%s.texture->webgpu.format),\n",
+						        get_name(g->name));
+						if (get_type(g->type)->tex_kind != TEXTURE_KIND_NONE) {
+							if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D) {
+								fprintf(output, "\t\t.dimension = WGPUTextureViewDimension_2D,\n");
+							}
+							else if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+								fprintf(output, "\t\t.dimension = WGPUTextureViewDimension_2DArray,\n");
+							}
+							else if (get_type(g->type)->tex_kind == TEXTURE_KIND_CUBE) {
+								fprintf(output, "\t\t.dimension = WGPUTextureViewDimension_Cube,\n");
+							}
+							else {
+								// TODO
+								assert(false);
+							}
 						}
 						fprintf(output, "\t\t.baseArrayLayer  = parameters->%s.base_array_layer,\n", get_name(g->name));
 						fprintf(output, "\t\t.arrayLayerCount = parameters->%s.array_layer_count,\n", get_name(g->name));
@@ -2127,7 +2181,10 @@ void kore3_export(char *directory, api_kind api) {
 							}
 						}
 					}
-					else if (!is_sampler(g->type) && g->type != bvh_type_id) {
+					else if (is_sampler(g->type)) {
+						fprintf(output, "\tkore_opengl_command_list_set_sampler(list, set->%s);\n", get_name(g->name));
+					}
+					else if (g->type != bvh_type_id) {
 						if (has_attribute(&g->attributes, add_name("indexed"))) {
 							fprintf(output, "\tkore_vulkan_descriptor_set_prepare_buffer(list, set->%s);\n", get_name(g->name));
 						}
@@ -2209,19 +2266,7 @@ void kore3_export(char *directory, api_kind api) {
 					fprintf(output, ");\n");
 				}
 				else if (api == API_WEBGPU) {
-					bool   found            = false;
-					size_t global_set_index = 0;
-
-					for (; global_set_index < get_sets_count(); ++global_set_index) {
-						if (get_set(global_set_index) == set) {
-							found = true;
-							break;
-						}
-					}
-
-					assert(found);
-
-					fprintf(output, "\n\tkore_webgpu_command_list_set_bind_group(list, %zu, &set->set", global_set_index);
+					fprintf(output, "\n\tkore_webgpu_command_list_set_bind_group(list, %s_table_index, &set->set", get_name(set->name));
 					if (dynamic_count > 0) {
 						fprintf(output, ", %u, dynamic_offsets", dynamic_count);
 					}
@@ -2230,16 +2275,6 @@ void kore3_export(char *directory, api_kind api) {
 					}
 					fprintf(output, ");\n");
 				}
-				else if (api == API_OPENGL) {
-					fprintf(output, "\tkore_%s_command_list_set_descriptor_table(list, %s_table_index, &set->set", api_short, get_name(set->name));
-					if (dynamic_count > 0) {
-						fprintf(output, ", dynamic_buffers, dynamic_offsets, dynamic_sizes");
-					}
-					else {
-						fprintf(output, ", NULL, NULL, NULL");
-					}
-					fprintf(output, ");\n");
-				}
 				fprintf(output, "}\n\n");
 			}
 		}
@@ -2278,25 +2313,23 @@ void kore3_export(char *directory, api_kind api) {
 					fprintf(output, "\tkore_%s_command_list_set_compute_pipeline(list, &%s);\n", api_short, get_name(f->name));
 				}
 
-				if (api != API_WEBGPU) {
-					descriptor_set_group *group = find_descriptor_set_group_for_function(f);
-					if (api == API_VULKAN) {
-						size_t index = 0;
-						for (size_t group_index = 0; group_index < group->size; ++group_index) {
-							if (group->values[group_index]->name != add_name("root_constants")) {
-								fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), index);
-								++index;
-							}
+				descriptor_set_group *group = find_descriptor_set_group_for_function(f);
+				if (api == API_VULKAN) {
+					size_t index = 0;
+					for (size_t group_index = 0; group_index < group->size; ++group_index) {
+						if (group->values[group_index]->name != add_name("root_constants")) {
+							fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), index);
+							++index;
 						}
 					}
-					else {
-						for (size_t group_index = 0; group_index < group->size; ++group_index) {
-							if (api == API_METAL) {
-								fprintf(output, "\t%s_compute_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index);
-							}
-							else {
-								fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index);
-							}
+				}
+				else {
+					for (size_t group_index = 0; group_index < group->size; ++group_index) {
+						if (api == API_METAL) {
+							fprintf(output, "\t%s_compute_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index);
+						}
+						else {
+							fprintf(output, "\t%s_table_index = %zu;\n", get_name(group->values[group_index]->name), group_index);
 						}
 					}
 				}
@@ -2355,10 +2388,19 @@ void kore3_export(char *directory, api_kind api) {
 
 				for (size_t j = 0; j < t->members.size; ++j) {
 					if (t->members.m[j].name == add_name("vertex")) {
-						if (api == API_METAL || api == API_WEBGPU) {
+						if (api == API_METAL) {
 							fprintf(output, "\t%s_parameters.vertex.shader.function_name = \"%s\";\n", get_name(t->name),
 							        get_name(t->members.m[j].value.identifier));
 						}
+						else if (api == API_WEBGPU) {
+							fprintf(output,
+							        "\t%s_parameters.vertex.shader.data = kore_webgpu_prepare_shader(device, %s_code, %s_code_size, "
+							        "%s_code_uses_framebuffer_texture_format);\n",
+							        get_name(t->name), get_name(t->members.m[j].value.identifier), get_name(t->members.m[j].value.identifier),
+							        get_name(t->members.m[j].value.identifier));
+							fprintf(output, "\t%s_parameters.vertex.shader.size = %s_code_size;\n\n", get_name(t->name),
+							        get_name(t->members.m[j].value.identifier));
+						}
 						else {
 							fprintf(output, "\t%s_parameters.vertex.shader.data = %s_code;\n", get_name(t->name), get_name(t->members.m[j].value.identifier));
 							fprintf(output, "\t%s_parameters.vertex.shader.size = %s_code_size;\n\n", get_name(t->name),
@@ -2379,10 +2421,19 @@ void kore3_export(char *directory, api_kind api) {
 						mesh_shader_name = t->members.m[j].value.identifier;
 					}
 					else if (t->members.m[j].name == add_name("fragment")) {
-						if (api == API_METAL || api == API_WEBGPU) {
+						if (api == API_METAL) {
 							fprintf(output, "\t%s_parameters.fragment.shader.function_name = \"%s\";\n", get_name(t->name),
 							        get_name(t->members.m[j].value.identifier));
 						}
+						else if (api == API_WEBGPU) {
+							fprintf(output,
+							        "\t%s_parameters.fragment.shader.data = kore_webgpu_prepare_shader(device, %s_code, %s_code_size, "
+							        "%s_code_uses_framebuffer_texture_format);\n",
+							        get_name(t->name), get_name(t->members.m[j].value.identifier), get_name(t->members.m[j].value.identifier),
+							        get_name(t->members.m[j].value.identifier));
+							fprintf(output, "\t%s_parameters.fragment.shader.size = %s_code_size;\n\n", get_name(t->name),
+							        get_name(t->members.m[j].value.identifier));
+						}
 						else {
 							fprintf(output, "\t%s_parameters.fragment.shader.data = %s_code;\n", get_name(t->name), get_name(t->members.m[j].value.identifier));
 							fprintf(output, "\t%s_parameters.fragment.shader.size = %s_code_size;\n\n", get_name(t->name),
@@ -2723,7 +2774,6 @@ void kore3_export(char *directory, api_kind api) {
 					fprintf(output, "\t}\n");
 				}
 				else if (api == API_WEBGPU) {
-
 					descriptor_set_group *group = find_descriptor_set_group_for_pipe_type(t);
 
 					fprintf(output, "\t{\n");
@@ -2734,27 +2784,14 @@ void kore3_export(char *directory, api_kind api) {
 					}
 					else {
 
-						fprintf(output, "\t\tWGPUBindGroupLayout layouts[%zu];\n", get_sets_count());
+						fprintf(output, "\t\tWGPUBindGroupLayout layouts[%zu];\n", group->size);
 
-						for (size_t layout_index = 0; layout_index < get_sets_count(); ++layout_index) {
-							bool found = false;
-							for (size_t layout_in_group_index = 0; layout_in_group_index < group->size; ++layout_in_group_index) {
-								if (get_set(layout_index) == group->values[layout_in_group_index]) {
-									found = true;
-									break;
-								}
-							}
-
-							if (found) {
-								fprintf(output, "\t\tlayouts[%zu] = %s_set_layout;\n", layout_index, get_name(get_set(layout_index)->name));
-							}
-							else {
-								fprintf(output, "\t\tlayouts[%zu] = NULL;\n", layout_index);
-							}
+						for (size_t layout_index = 0; layout_index < group->size; ++layout_index) {
+							fprintf(output, "\t\tlayouts[%zu] = %s_set_layout;\n", layout_index, get_name(group->values[layout_index]->name));
 						}
 
 						fprintf(output, "\t\tkore_webgpu_render_pipeline_init(&device->webgpu, &%s, &%s_parameters, layouts, %zu);\n", get_name(t->name),
-						        get_name(t->name), get_sets_count());
+						        get_name(t->name), group->size);
 					}
 
 					fprintf(output, "\t}\n");
@@ -2791,9 +2828,16 @@ void kore3_export(char *directory, api_kind api) {
 			function *f = get_function(i);
 			if (has_attribute(&f->attributes, add_name("compute"))) {
 				fprintf(output, "\tkore_%s_compute_pipeline_parameters %s_parameters;\n", api_short, get_name(f->name));
-				if (api == API_METAL || api == API_WEBGPU) {
+				if (api == API_METAL) {
 					fprintf(output, "\t%s_parameters.shader.function_name = \"%s\";\n", get_name(f->name), get_name(f->name));
 				}
+				else if (api == API_WEBGPU) {
+					fprintf(output,
+					        "\t%s_parameters.shader.data = kore_webgpu_prepare_shader(device, %s_code, %s_code_size, "
+					        "%s_code_uses_framebuffer_texture_format);\n",
+					        get_name(f->name), get_name(f->name), get_name(f->name), get_name(f->name));
+					fprintf(output, "\t%s_parameters.shader.size = %s_code_size;\n", get_name(f->name), get_name(f->name));
+				}
 				else {
 					fprintf(output, "\t%s_parameters.shader.data = %s_code;\n", get_name(f->name), get_name(f->name));
 					fprintf(output, "\t%s_parameters.shader.size = %s_code_size;\n", get_name(f->name), get_name(f->name));
@@ -2850,27 +2894,14 @@ void kore3_export(char *directory, api_kind api) {
 						        get_name(f->name));
 					}
 					else {
-						fprintf(output, "\t\tWGPUBindGroupLayout layouts[%zu];\n", get_sets_count());
-
-						for (size_t layout_index = 0; layout_index < get_sets_count(); ++layout_index) {
-							bool found = false;
-							for (size_t layout_in_group_index = 0; layout_in_group_index < group->size; ++layout_in_group_index) {
-								if (get_set(layout_index) == group->values[layout_in_group_index]) {
-									found = true;
-									break;
-								}
-							}
+						fprintf(output, "\t\tWGPUBindGroupLayout layouts[%zu];\n", group->size);
 
-							if (found) {
-								fprintf(output, "\t\tlayouts[%zu] = %s_set_layout;\n", layout_index, get_name(get_set(layout_index)->name));
-							}
-							else {
-								fprintf(output, "\t\tlayouts[%zu] = NULL;\n", layout_index);
-							}
+						for (size_t layout_index = 0; layout_index < group->size; ++layout_index) {
+							fprintf(output, "\t\tlayouts[%zu] = %s_set_layout;\n", layout_index, get_name(group->values[layout_index]->name));
 						}
 
 						fprintf(output, "\t\tkore_webgpu_compute_pipeline_init(&device->webgpu, &%s, &%s_parameters, layouts, %zu);\n", get_name(f->name),
-						        get_name(f->name), get_sets_count());
+						        get_name(f->name), group->size);
 					}
 
 					fprintf(output, "\t}\n");
@@ -2990,47 +3021,53 @@ void kore3_export(char *directory, api_kind api) {
 				bool    readable = set->globals.readable[global_index];
 				bool    writable = set->globals.writable[global_index];
 
-				if (g->type == tex2d_type_id) {
-					fprintf(output, "\t\t\t{\n");
-					fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
-					if (readable | writable) {
-						fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,\n");
+				if (get_type(g->type)->tex_kind != TEXTURE_KIND_NONE) {
+					if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D) {
+						fprintf(output, "\t\t\t{\n");
+						fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
+						if (readable | writable) {
+							fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,\n");
+						}
+						else {
+							fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,\n");
+						}
+						fprintf(output, "\t\t\t\t.descriptorCount = 1,\n");
+						fprintf(output, "\t\t\t\t.stageFlags = VK_SHADER_STAGE_FRAGMENT_BIT | VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_COMPUTE_BIT,\n");
+						fprintf(output, "\t\t\t\t.pImmutableSamplers = NULL,\n");
+						fprintf(output, "\t\t\t},\n");
 					}
-					else {
-						fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,\n");
+					else if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+						fprintf(output, "\t\t\t{\n");
+						fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
+						if (readable | writable) {
+							fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,\n");
+						}
+						else {
+							fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,\n");
+						}
+						fprintf(output, "\t\t\t\t.descriptorCount = 1,\n");
+						fprintf(output, "\t\t\t\t.stageFlags = VK_SHADER_STAGE_FRAGMENT_BIT | VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_COMPUTE_BIT,\n");
+						fprintf(output, "\t\t\t\t.pImmutableSamplers = NULL,\n");
+						fprintf(output, "\t\t\t},\n");
 					}
-					fprintf(output, "\t\t\t\t.descriptorCount = 1,\n");
-					fprintf(output, "\t\t\t\t.stageFlags = VK_SHADER_STAGE_FRAGMENT_BIT | VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_COMPUTE_BIT,\n");
-					fprintf(output, "\t\t\t\t.pImmutableSamplers = NULL,\n");
-					fprintf(output, "\t\t\t},\n");
-				}
-				else if (g->type == tex2darray_type_id) {
-					fprintf(output, "\t\t\t{\n");
-					fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
-					if (readable | writable) {
-						fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,\n");
+					else if (get_type(g->type)->tex_kind == TEXTURE_KIND_CUBE) {
+						fprintf(output, "\t\t\t{\n");
+						fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
+						if (readable | writable) {
+							fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,\n");
+						}
+						else {
+							fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,\n");
+						}
+						fprintf(output, "\t\t\t\t.descriptorCount = 1,\n");
+						fprintf(output, "\t\t\t\t.stageFlags = VK_SHADER_STAGE_FRAGMENT_BIT | VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_COMPUTE_BIT,\n");
+						fprintf(output, "\t\t\t\t.pImmutableSamplers = NULL,\n");
+						fprintf(output, "\t\t\t},\n");
 					}
 					else {
-						fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,\n");
+						// TODO
+						assert(false);
 					}
-					fprintf(output, "\t\t\t\t.descriptorCount = 1,\n");
-					fprintf(output, "\t\t\t\t.stageFlags = VK_SHADER_STAGE_FRAGMENT_BIT | VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_COMPUTE_BIT,\n");
-					fprintf(output, "\t\t\t\t.pImmutableSamplers = NULL,\n");
-					fprintf(output, "\t\t\t},\n");
-				}
-				else if (g->type == texcube_type_id) {
-					fprintf(output, "\t\t\t{\n");
-					fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
-					if (readable | writable) {
-						fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,\n");
-					}
-					else {
-						fprintf(output, "\t\t\t\t.descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,\n");
-					}
-					fprintf(output, "\t\t\t\t.descriptorCount = 1,\n");
-					fprintf(output, "\t\t\t\t.stageFlags = VK_SHADER_STAGE_FRAGMENT_BIT | VK_SHADER_STAGE_VERTEX_BIT | VK_SHADER_STAGE_COMPUTE_BIT,\n");
-					fprintf(output, "\t\t\t\t.pImmutableSamplers = NULL,\n");
-					fprintf(output, "\t\t\t},\n");
 				}
 				else if (is_sampler(g->type)) {
 					fprintf(output, "\t\t\t{\n");
@@ -3099,9 +3136,15 @@ void kore3_export(char *directory, api_kind api) {
 
 		fprintf(output, "#include <assert.h>\n\n");
 
+		fprintf(output, "WGPUBindGroupLayout root_constants_set_layout;\n");
+
 		for (size_t set_index = 0; set_index < sets_count; ++set_index) {
 			descriptor_set *set = sets[set_index];
 
+			if (set->name == add_name("root_constants")) {
+				continue;
+			}
+
 			fprintf(output, "WGPUBindGroupLayout %s_set_layout;\n", get_name(set->name));
 		}
 
@@ -3109,9 +3152,29 @@ void kore3_export(char *directory, api_kind api) {
 
 		fprintf(output, "void create_bind_group_layouts(kore_gpu_device *device) {\n");
 
+		fprintf(output, "\t{\n");
+		fprintf(output, "\t\tWGPUBindGroupLayoutEntry layout_entries[1] = {\n");
+		fprintf(output, "\t\t\t{\n");
+		fprintf(output, "\t\t\t\t.binding    = 0,\n");
+		fprintf(output, "\t\t\t\t.buffer     = {.type = WGPUBufferBindingType_Uniform, .hasDynamicOffset = true},\n");
+		fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+		fprintf(output, "\t\t\t},\n");
+		fprintf(output, "\t\t};\n");
+		fprintf(output, "\n");
+		fprintf(output, "\t\tWGPUBindGroupLayoutDescriptor bind_group_layout_descriptor = {\n");
+		fprintf(output, "\t\t\t.entryCount = 1,\n");
+		fprintf(output, "\t\t\t.entries    = layout_entries,\n");
+		fprintf(output, "\t\t};\n");
+		fprintf(output, "\t\troot_constants_set_layout = wgpuDeviceCreateBindGroupLayout(device->webgpu.device, &bind_group_layout_descriptor);\n");
+		fprintf(output, "\t}\n");
+
 		for (size_t set_index = 0; set_index < sets_count; ++set_index) {
 			descriptor_set *set = sets[set_index];
 
+			if (set->name == add_name("root_constants")) {
+				continue;
+			}
+
 			fprintf(output, "\t{\n");
 
 			fprintf(output, "\t\tWGPUBindGroupLayoutEntry layout_entries[%zu] = {\n", set->globals.size);
@@ -3120,65 +3183,92 @@ void kore3_export(char *directory, api_kind api) {
 				global *g        = get_global(set->globals.globals[global_index]);
 				bool    writable = set->globals.writable[global_index];
 
-				if (g->type == tex2d_type_id) {
-					fprintf(output, "\t\t\t{\n");
-					fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
-					if (writable) {
-						fprintf(output, "\t\t\t\t.storageTexture = {.viewDimension = WGPUTextureViewDimension_2D, .format = WGPUTextureFormat_RGBA32Float, "
-						                ".access = WGPUStorageTextureAccess_WriteOnly},\n");
-					}
-					else {
-						fprintf(output, "\t\t\t\t.texture = {.sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_2D},\n");
-					}
-					if (writable) {
-						fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
-					}
-					else {
-						fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
-					}
-					fprintf(output, "\t\t\t},\n");
-				}
-				else if (g->type == tex2darray_type_id) {
-					fprintf(output, "\t\t\t{\n");
-					fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
-					if (writable) {
-						fprintf(output,
-						        "\t\t\t\t.storageTexture = {.viewDimension = WGPUTextureViewDimension_2DArray, .format = WGPUTextureFormat_RGBA32Float, "
-						        ".access = WGPUStorageTextureAccess_WriteOnly},\n");
+				if (get_type(g->type)->tex_kind != TEXTURE_KIND_NONE) {
+					char format[64];
+					format_to_string_wgsl(get_type(g->type)->tex_format, format);
+
+					if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D) {
+						fprintf(output, "\t\t\t{\n");
+						fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
+						if (writable) {
+							fprintf(output,
+							        "\t\t\t\t.storageTexture = {.viewDimension = WGPUTextureViewDimension_2D, .format = WGPUTextureFormat_%s, "
+							        ".access = WGPUStorageTextureAccess_WriteOnly},\n",
+							        format);
+						}
+						else {
+							if (is_depth(get_type(g->type)->tex_format)) {
+								fprintf(output,
+								        "\t\t\t\t.texture = {.sampleType = WGPUTextureSampleType_Depth, .viewDimension = WGPUTextureViewDimension_2D},\n");
+							}
+							else {
+								fprintf(output,
+								        "\t\t\t\t.texture = {.sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_2D},\n");
+							}
+						}
+						if (writable) {
+							fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+						}
+						else {
+							fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+						}
+						fprintf(output, "\t\t\t},\n");
 					}
-					else {
-						fprintf(output, "\t\t\t\t.texture = {.sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_2DArray},\n");
+					else if (get_type(g->type)->tex_kind == TEXTURE_KIND_2D_ARRAY) {
+						fprintf(output, "\t\t\t{\n");
+						fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
+						if (writable) {
+							fprintf(output,
+							        "\t\t\t\t.storageTexture = {.viewDimension = WGPUTextureViewDimension_2DArray, .format = WGPUTextureFormat_%s, "
+							        ".access = WGPUStorageTextureAccess_WriteOnly},\n",
+							        format);
+						}
+						else {
+							fprintf(output,
+							        "\t\t\t\t.texture = {.sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_2DArray},\n");
+						}
+						if (writable) {
+							fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+						}
+						else {
+							fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+						}
+						fprintf(output, "\t\t\t},\n");
 					}
-					if (writable) {
-						fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+					else if (get_type(g->type)->tex_kind == TEXTURE_KIND_CUBE) {
+						fprintf(output, "\t\t\t{\n");
+						fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
+						if (writable) {
+							fprintf(output,
+							        "\t\t\t\t.storageTexture = {.viewDimension = WGPUTextureViewDimension_Cube, .format = WGPUTextureFormat_%s, "
+							        ".access = WGPUStorageTextureAccess_WriteOnly},\n",
+							        format);
+						}
+						else {
+							fprintf(output,
+							        "\t\t\t\t.texture = {.sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_Cube},\n");
+						}
+						if (writable) {
+							fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+						}
+						else {
+							fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+						}
+						fprintf(output, "\t\t\t},\n");
 					}
-					else {
-						fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
+					else { // TODO
+						assert(false);
 					}
-					fprintf(output, "\t\t\t},\n");
 				}
-				else if (g->type == texcube_type_id) {
+				else if (is_sampler(g->type)) {
 					fprintf(output, "\t\t\t{\n");
 					fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
-					if (writable) {
-						fprintf(output, "\t\t\t\t.storageTexture = {.viewDimension = WGPUTextureViewDimension_Cube, .format = WGPUTextureFormat_RGBA32Float, "
-						                ".access = WGPUStorageTextureAccess_WriteOnly},\n");
+					if ((g->usage & GLOBAL_USAGE_SAMPLE_DEPTH) != 0) {
+						fprintf(output, "\t\t\t\t.sampler = {.type = WGPUSamplerBindingType_NonFiltering},\n");
 					}
 					else {
-						fprintf(output, "\t\t\t\t.texture = {.sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_Cube},\n");
+						fprintf(output, "\t\t\t\t.sampler = {.type = WGPUSamplerBindingType_Filtering},\n");
 					}
-					if (writable) {
-						fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
-					}
-					else {
-						fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
-					}
-					fprintf(output, "\t\t\t},\n");
-				}
-				else if (is_sampler(g->type)) {
-					fprintf(output, "\t\t\t{\n");
-					fprintf(output, "\t\t\t\t.binding = %zu,\n", global_index);
-					fprintf(output, "\t\t\t\t.sampler = {.type = WGPUSamplerBindingType_Filtering},\n");
 					fprintf(output, "\t\t\t\t.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute,\n");
 					fprintf(output, "\t\t\t},\n");
 				}

base/sources/libs/kong/kong.c

@@ -70,6 +70,7 @@ typedef enum integration_kind { INTEGRATION_KORE3 } integration_kind;
 ////
 /*
 ////
+
 int main(int argc, char **argv) {
 	arg_mode mode = MODE_MODECHECK;
 

base/sources/libs/kong/parser.c

@@ -1114,6 +1114,148 @@ static definition parse_function(state_t *state) {
 	return d;
 }
 
+static texture_format convert_texture_format(state_t *state, name_id format_name) {
+	if (format_name == NO_NAME) {
+		return TEXTURE_FORMAT_UNDEFINED;
+	}
+	else if (format_name == add_name("framebuffer_format")) {
+		return TEXTURE_FORMAT_FRAMEBUFFER;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_DEPTH")) {
+		return TEXTURE_FORMAT_DEPTH;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R8_UNORM")) {
+		return TEXTURE_FORMAT_R8_UNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R8_SNORM")) {
+		return TEXTURE_FORMAT_R8_SNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R8_UINT")) {
+		return TEXTURE_FORMAT_R8_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R8_SINT")) {
+		return TEXTURE_FORMAT_R8_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R16_UINT")) {
+		return TEXTURE_FORMAT_R16_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R16_SINT")) {
+		return TEXTURE_FORMAT_R16_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R16_FLOAT")) {
+		return TEXTURE_FORMAT_R16_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG8_UNORM")) {
+		return TEXTURE_FORMAT_RG8_UNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG8_SNORM")) {
+		return TEXTURE_FORMAT_RG8_SNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG8_UINT")) {
+		return TEXTURE_FORMAT_RG8_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG8_SINT")) {
+		return TEXTURE_FORMAT_RG8_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R32_UINT")) {
+		return TEXTURE_FORMAT_R32_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R32_SINT")) {
+		return TEXTURE_FORMAT_R32_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_R32_FLOAT")) {
+		return TEXTURE_FORMAT_R32_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG16_UINT")) {
+		return TEXTURE_FORMAT_RG16_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG16_SINT")) {
+		return TEXTURE_FORMAT_RG16_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG16_FLOAT")) {
+		return TEXTURE_FORMAT_RG16_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA8_UNORM")) {
+		return TEXTURE_FORMAT_RGBA8_UNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA8_UNORM_SRGB")) {
+		return TEXTURE_FORMAT_RGBA8_UNORM_SRGB;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA8_SNORM")) {
+		return TEXTURE_FORMAT_RGBA8_SNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA8_UINT")) {
+		return TEXTURE_FORMAT_RGBA8_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA8_SINT")) {
+		return TEXTURE_FORMAT_RGBA8_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_BGRA8_UNORM")) {
+		return TEXTURE_FORMAT_BGRA8_UNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_BGRA8_UNORM_SRGB")) {
+		return TEXTURE_FORMAT_BGRA8_UNORM_SRGB;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGB9E5U_FLOAT")) {
+		return TEXTURE_FORMAT_RGB9E5U_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGB10A2_UINT")) {
+		return TEXTURE_FORMAT_RGB10A2_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGB10A2_UNORM")) {
+		return TEXTURE_FORMAT_RGB10A2_UNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG11B10U_FLOAT")) {
+		return TEXTURE_FORMAT_RG11B10U_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG32_UINT")) {
+		return TEXTURE_FORMAT_RG32_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG32_SINT")) {
+		return TEXTURE_FORMAT_RG32_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RG32_FLOAT")) {
+		return TEXTURE_FORMAT_RG32_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA16_UINT")) {
+		return TEXTURE_FORMAT_RGBA16_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA16_SINT")) {
+		return TEXTURE_FORMAT_RGBA16_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA16_FLOAT")) {
+		return TEXTURE_FORMAT_RGBA16_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA32_UINT")) {
+		return TEXTURE_FORMAT_RGBA32_UINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA32_SINT")) {
+		return TEXTURE_FORMAT_RGBA32_SINT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_RGBA32_FLOAT")) {
+		return TEXTURE_FORMAT_RGBA32_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_DEPTH16_UNORM")) {
+		return TEXTURE_FORMAT_DEPTH16_UNORM;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_DEPTH24_NOTHING8")) {
+		return TEXTURE_FORMAT_DEPTH24_NOTHING8;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_DEPTH24_STENCIL8")) {
+		return TEXTURE_FORMAT_DEPTH24_STENCIL8;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_DEPTH32_FLOAT")) {
+		return TEXTURE_FORMAT_DEPTH32_FLOAT;
+	}
+	else if (format_name == add_name("TEXTURE_FORMAT_DEPTH32_FLOAT_STENCIL8_NOTHING24")) {
+		return TEXTURE_FORMAT_DEPTH32_FLOAT_STENCIL8_NOTHING24;
+	}
+	else {
+		error(state->context, "Unknown texture format %s", format_name);
+		return TEXTURE_FORMAT_UNDEFINED;
+	}
+}
+
 static definition parse_const(state_t *state, attribute_list attributes) {
 	advance_state(state);
 	match_token(state, TOKEN_IDENTIFIER, "Expected an identifier");
@@ -1123,8 +1265,9 @@ static definition parse_const(state_t *state, attribute_list attributes) {
 	match_token(state, TOKEN_COLON, "Expected a colon");
 	advance_state(state);
 
-	name_id type_name = NO_NAME;
-	type_id type      = NO_TYPE;
+	name_id type_name   = NO_NAME;
+	type_id type        = NO_TYPE;
+	name_id format_name = NO_NAME;
 
 	if (current(state).kind == TOKEN_LEFT_CURLY) {
 		type = parse_struct_inner(state, NO_NAME).type;
@@ -1157,39 +1300,99 @@ static definition parse_const(state_t *state, attribute_list attributes) {
 		value = parse_expression(state);
 	}
 
+	if (current(state).kind == TOKEN_OPERATOR && current(state).op == OPERATOR_LESS) {
+		advance_state(state);
+		match_token(state, TOKEN_IDENTIFIER, "Expected an identifier");
+		format_name = current(state).identifier;
+		advance_state(state);
+
+		if (current(state).kind == TOKEN_LEFT_PAREN) {
+			advance_state(state);
+			match_token(state, TOKEN_RIGHT_PAREN, "Expected a right paren");
+			advance_state(state);
+		}
+
+		if (current(state).kind != TOKEN_OPERATOR || current(state).op != OPERATOR_GREATER) {
+			error(state->context, "Expected a greater than");
+		}
+		advance_state(state);
+	}
+
 	match_token(state, TOKEN_SEMICOLON, "Expected a semicolon");
 	advance_state(state);
 
 	definition d = {0};
 
+	name_id tex1d_name        = add_name("tex1d");
+	name_id tex2d_name        = add_name("tex2d");
+	name_id tex3d_name        = add_name("tex3d");
+	name_id texcube_name      = add_name("texcube");
+	name_id tex1darray_name   = add_name("tex1darray");
+	name_id tex2darray_name   = add_name("tex2darray");
+	name_id texcubearray_name = add_name("texcubearray");
+
 	if (type_name == NO_NAME) {
 		debug_context context = {0};
 		check(type != NO_TYPE, context, "Const has no type");
 		d.kind   = DEFINITION_CONST_CUSTOM;
 		d.global = add_global(type, attributes, name.identifier);
 	}
-	else if (type_name == add_name("tex2d")) {
-		d.kind = DEFINITION_TEX2D;
+	else if (type_name == tex1d_name || type_name == tex2d_name || type_name == tex3d_name || type_name == texcube_name || type_name == tex1darray_name ||
+	         type_name == tex2darray_name || type_name == texcubearray_name) {
+		struct type tex_type;
+		tex_type.name                        = type_name;
+		tex_type.attributes.attributes_count = 0;
+		tex_type.members.size                = 0;
+		tex_type.built_in                    = true;
+		tex_type.array_size                  = 0;
+		tex_type.base                        = NO_TYPE;
+
+		if (type_name == tex1d_name) {
+			d.kind            = DEFINITION_TEX1D;
+			tex_type.tex_kind = TEXTURE_KIND_1D;
+		}
+		else if (type_name == tex2d_name) {
+			d.kind            = DEFINITION_TEX2D;
+			tex_type.tex_kind = TEXTURE_KIND_2D;
+		}
+		else if (type_name == tex3d_name) {
+			d.kind            = DEFINITION_TEX3D;
+			tex_type.tex_kind = TEXTURE_KIND_3D;
+		}
+		else if (type_name == texcube_name) {
+			d.kind            = DEFINITION_TEXCUBE;
+			tex_type.tex_kind = TEXTURE_KIND_CUBE;
+		}
+		else if (type_name == tex1darray_name) {
+			d.kind            = DEFINITION_TEX1DARRAY;
+			tex_type.tex_kind = TEXTURE_KIND_1D_ARRAY;
+		}
+		else if (type_name == tex2darray_name) {
+			d.kind            = DEFINITION_TEX2DARRAY;
+			tex_type.tex_kind = TEXTURE_KIND_2D_ARRAY;
+		}
+		else if (type_name == texcubearray_name) {
+			d.kind            = DEFINITION_TEXCUBEARRAY;
+			tex_type.tex_kind = TEXTURE_KIND_CUBE_ARRAY;
+		}
+		else {
+			assert(false);
+		}
+
+		tex_type.tex_format = convert_texture_format(state, format_name);
+
+		type_id t_id = add_full_type(&tex_type);
 
-		type_id t_id = tex2d_type_id;
 		if (array) {
-			type_id array_type_id               = add_type(get_type(t_id)->name);
+			type_id array_type_id               = add_type(type_name);
 			get_type(array_type_id)->base       = t_id;
-			get_type(array_type_id)->built_in   = get_type(t_id)->built_in;
+			get_type(array_type_id)->built_in   = true;
 			get_type(array_type_id)->array_size = array_size;
 			t_id                                = array_type_id;
 		}
 
 		d.global = add_global(t_id, attributes, name.identifier);
 	}
-	else if (type_name == add_name("tex2darray")) {
-		d.kind   = DEFINITION_TEX2DARRAY;
-		d.global = add_global(tex2darray_type_id, attributes, name.identifier);
-	}
-	else if (type_name == add_name("texcube")) {
-		d.kind   = DEFINITION_TEXCUBE;
-		d.global = add_global(texcube_type_id, attributes, name.identifier);
-	}
 	else if (type_name == add_name("sampler")) {
 		d.kind   = DEFINITION_SAMPLER;
 		d.global = add_global(sampler_type_id, attributes, name.identifier);

base/sources/libs/kong/parser.h

@@ -130,9 +130,13 @@ typedef struct definition {
 	enum {
 		DEFINITION_FUNCTION,
 		DEFINITION_STRUCT,
+		DEFINITION_TEX1D,
 		DEFINITION_TEX2D,
-		DEFINITION_TEX2DARRAY,
+		DEFINITION_TEX3D,
 		DEFINITION_TEXCUBE,
+		DEFINITION_TEX1DARRAY,
+		DEFINITION_TEX2DARRAY,
+		DEFINITION_TEXCUBEARRAY,
 		DEFINITION_SAMPLER,
 		DEFINITION_CONST_CUSTOM,
 		DEFINITION_CONST_BASIC,

base/sources/libs/kong/typer.c

@@ -43,14 +43,30 @@ static void resolve_types_in_element(statement *parent_block, expression *elemen
 
 	assert(of_type != NO_TYPE);
 
-	if (of_type == tex2d_type_id) {
-		element->type.type = float4_id;
-	}
-	else if (of_type == tex2darray_type_id) {
-		element->type.type = tex2d_type_id;
+	type *of = get_type(of_type);
+
+	if (of->tex_kind != TEXTURE_KIND_NONE) {
+		if (of->tex_format == TEXTURE_FORMAT_UNDEFINED) {
+			element->type.type = float4_id;
+		}
+		else if (of->tex_format == TEXTURE_FORMAT_FRAMEBUFFER) {
+			element->type.type = float4_id;
+		}
+		else if (of->tex_format == TEXTURE_FORMAT_RGBA32_FLOAT) {
+			element->type.type = float4_id;
+		}
+		else if (of->tex_format == TEXTURE_FORMAT_RGBA8_UNORM) {
+			element->type.type = float4_id;
+		}
+		else if (of->tex_format == TEXTURE_FORMAT_DEPTH) {
+			element->type.type = float_id;
+		}
+		else {
+			// TODO
+			assert(false);
+		}
 	}
 	else {
-		type *of = get_type(of_type);
 		if (of->array_size > 0) {
 			element->type.type = of->base;
 		}
@@ -561,13 +577,31 @@ void resolve_types_in_expression(statement *parent, expression *e) {
 		break;
 	}
 	case EXPRESSION_CALL: {
-		for (function_id i = 0; get_function(i) != NULL; ++i) {
-			function *f = get_function(i);
-			if (f->name == e->call.func_name) {
-				e->type = f->return_type;
-				break;
+		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);
+				if (is_depth(get_type(g->type)->tex_format)) {
+					e->type.type = float_id;
+				}
+				else {
+					e->type.type = float4_id;
+				}
+			}
+			else {
+				e->type.type = float4_id;
 			}
 		}
+		else {
+			for (function_id i = 0; get_function(i) != NULL; ++i) {
+				function *f = get_function(i);
+				if (f->name == e->call.func_name) {
+					e->type = f->return_type;
+					break;
+				}
+			}
+		}
+
 		for (size_t i = 0; i < e->call.parameters.size; ++i) {
 			resolve_types_in_expression(parent, e->call.parameters.e[i]);
 		}

base/sources/libs/kong/types.c

@@ -40,9 +40,6 @@ type_id bool2_id;
 type_id bool3_id;
 type_id bool4_id;
 type_id function_type_id;
-type_id tex2d_type_id;
-type_id tex2darray_type_id;
-type_id texcube_type_id;
 type_id sampler_type_id;
 type_id ray_type_id;
 type_id bvh_type_id;
@@ -68,14 +65,8 @@ void types_init(void) {
 	void_id                     = add_type(add_name("void"));
 	get_type(void_id)->built_in = true;
 
-	sampler_type_id                        = add_type(add_name("sampler"));
-	get_type(sampler_type_id)->built_in    = true;
-	tex2d_type_id                          = add_type(add_name("tex2d"));
-	get_type(tex2d_type_id)->built_in      = true;
-	tex2darray_type_id                     = add_type(add_name("tex2darray"));
-	get_type(tex2darray_type_id)->built_in = true;
-	texcube_type_id                        = add_type(add_name("texcube"));
-	get_type(texcube_type_id)->built_in    = true;
+	sampler_type_id                     = add_type(add_name("sampler"));
+	get_type(sampler_type_id)->built_in = true;
 
 	bool_id                      = add_type(add_name("bool"));
 	get_type(bool_id)->built_in  = true;
@@ -193,6 +184,11 @@ static void grow_if_needed(uint64_t size) {
 	}
 }
 
+static bool types_equal(type *a, type *b) {
+	return a->name == b->name && a->attributes.attributes_count == 0 && b->attributes.attributes_count == 0 && a->members.size == 0 && b->members.size == 0 &&
+	       a->built_in == b->built_in && a->array_size == b->array_size && a->base == b->base && a->tex_kind == b->tex_kind && a->tex_format == b->tex_format;
+}
+
 type_id add_type(name_id name) {
 	grow_if_needed(next_type_index + 1);
 
@@ -205,6 +201,25 @@ type_id add_type(name_id name) {
 	types[s].built_in                    = false;
 	types[s].array_size                  = 0;
 	types[s].base                        = NO_TYPE;
+	types[s].tex_kind                    = TEXTURE_KIND_NONE;
+	types[s].tex_format                  = TEXTURE_FORMAT_UNDEFINED;
+
+	return s;
+}
+
+type_id add_full_type(type *t) {
+	for (type_id type_index = 0; type_index < next_type_index; ++type_index) {
+		if (types_equal(&types[type_index], t)) {
+			return type_index;
+		}
+	}
+
+	grow_if_needed(next_type_index + 1);
+
+	type_id s = next_type_index;
+	++next_type_index;
+
+	types[s] = *t;
 
 	return s;
 }
@@ -401,3 +416,17 @@ type_id vector_to_size(type_id vector_type, uint32_t size) {
 		return float_id;
 	}
 }
+
+bool is_depth(texture_format format) {
+	switch (format) {
+	case TEXTURE_FORMAT_DEPTH16_UNORM:
+	case TEXTURE_FORMAT_DEPTH24_NOTHING8:
+	case TEXTURE_FORMAT_DEPTH24_STENCIL8:
+	case TEXTURE_FORMAT_DEPTH32_FLOAT:
+	case TEXTURE_FORMAT_DEPTH32_FLOAT_STENCIL8_NOTHING24:
+	case TEXTURE_FORMAT_DEPTH:
+		return true;
+	default:
+		return false;
+	}
+}

base/sources/libs/kong/types.h

@@ -51,6 +51,69 @@ bool has_attribute(attribute_list *attributes, name_id name);
 
 attribute *find_attribute(attribute_list *attributes, name_id name);
 
+typedef enum texture_kind {
+	TEXTURE_KIND_NONE,
+	TEXTURE_KIND_1D,
+	TEXTURE_KIND_2D,
+	TEXTURE_KIND_3D,
+	TEXTURE_KIND_1D_ARRAY,
+	TEXTURE_KIND_2D_ARRAY,
+	TEXTURE_KIND_CUBE,
+	TEXTURE_KIND_CUBE_ARRAY,
+} texture_kind;
+
+typedef enum texture_format {
+	TEXTURE_FORMAT_UNDEFINED,
+	TEXTURE_FORMAT_R8_UNORM,
+	TEXTURE_FORMAT_R8_SNORM,
+	TEXTURE_FORMAT_R8_UINT,
+	TEXTURE_FORMAT_R8_SINT,
+	TEXTURE_FORMAT_R16_UINT,
+	TEXTURE_FORMAT_R16_SINT,
+	TEXTURE_FORMAT_R16_FLOAT,
+	TEXTURE_FORMAT_RG8_UNORM,
+	TEXTURE_FORMAT_RG8_SNORM,
+	TEXTURE_FORMAT_RG8_UINT,
+	TEXTURE_FORMAT_RG8_SINT,
+	TEXTURE_FORMAT_R32_UINT,
+	TEXTURE_FORMAT_R32_SINT,
+	TEXTURE_FORMAT_R32_FLOAT,
+	TEXTURE_FORMAT_RG16_UINT,
+	TEXTURE_FORMAT_RG16_SINT,
+	TEXTURE_FORMAT_RG16_FLOAT,
+	TEXTURE_FORMAT_RGBA8_UNORM,
+	TEXTURE_FORMAT_RGBA8_UNORM_SRGB,
+	TEXTURE_FORMAT_RGBA8_SNORM,
+	TEXTURE_FORMAT_RGBA8_UINT,
+	TEXTURE_FORMAT_RGBA8_SINT,
+	TEXTURE_FORMAT_BGRA8_UNORM,
+	TEXTURE_FORMAT_BGRA8_UNORM_SRGB,
+	TEXTURE_FORMAT_RGB9E5U_FLOAT,
+	TEXTURE_FORMAT_RGB10A2_UINT,
+	TEXTURE_FORMAT_RGB10A2_UNORM,
+	TEXTURE_FORMAT_RG11B10U_FLOAT,
+	TEXTURE_FORMAT_RG32_UINT,
+	TEXTURE_FORMAT_RG32_SINT,
+	TEXTURE_FORMAT_RG32_FLOAT,
+	TEXTURE_FORMAT_RGBA16_UINT,
+	TEXTURE_FORMAT_RGBA16_SINT,
+	TEXTURE_FORMAT_RGBA16_FLOAT,
+	TEXTURE_FORMAT_RGBA32_UINT,
+	TEXTURE_FORMAT_RGBA32_SINT,
+	TEXTURE_FORMAT_RGBA32_FLOAT,
+	// TEXTURE_FORMAT_STENCIL8, // not available in d3d12
+	TEXTURE_FORMAT_DEPTH16_UNORM,
+	TEXTURE_FORMAT_DEPTH24_NOTHING8,
+	TEXTURE_FORMAT_DEPTH24_STENCIL8,
+	TEXTURE_FORMAT_DEPTH32_FLOAT,
+	TEXTURE_FORMAT_DEPTH32_FLOAT_STENCIL8_NOTHING24,
+
+	TEXTURE_FORMAT_DEPTH,
+	TEXTURE_FORMAT_FRAMEBUFFER,
+} texture_format;
+
+bool is_depth(texture_format format);
+
 typedef struct type {
 	attribute_list attributes;
 	name_id        name;
@@ -60,12 +123,17 @@ typedef struct type {
 
 	type_id  base;
 	uint32_t array_size;
+
+	texture_kind   tex_kind;
+	texture_format tex_format;
 } type;
 
 void types_init(void);
 
 type_id add_type(name_id name);
 
+type_id add_full_type(type *t);
+
 type_id find_type_by_name(name_id name);
 
 type_id find_type_by_ref(type_ref *t);
@@ -98,21 +166,22 @@ extern type_id bool_id;
 extern type_id bool2_id;
 extern type_id bool3_id;
 extern type_id bool4_id;
-extern type_id tex2d_type_id;
-extern type_id tex2darray_type_id;
-extern type_id texcube_type_id;
 extern type_id sampler_type_id;
 extern type_id ray_type_id;
 extern type_id bvh_type_id;
 
 static inline bool is_texture(type_id id) {
-	if (id == tex2d_type_id || id == tex2darray_type_id || id == texcube_type_id) {
-		return true;
-	}
+	while (id != NO_TYPE) {
+		type *t = get_type(id);
+
+		if (t->tex_kind != TEXTURE_KIND_NONE) {
+			return true;
+		}
 
-	type *t = get_type(id);
+		id = t->base;
+	}
 
-	return t->base == tex2d_type_id || t->base == tex2darray_type_id || t->base == texcube_type_id;
+	return false;
 }
 
 static inline bool is_cbv_srv_uav(type_id t) {