sharkk/gpu public

branch main
showing the latest snapshot on main
emit_spirv.odin 69.7 KB · Plain text
package shader

import "core:fmt"

// SPIR-V 1.5 binary backend — emits from IR_Module
// Uses function-scoped OpVariable + OpStore/OpLoad instead of SSA phi nodes.
// This produces valid SPIR-V; GPU drivers handle mem2reg.

SPIRV_Loop_Context :: struct {
	continue_label: u32,
	merge_label:    u32,
}

SPIRV_Builder :: struct {
	next_id:    u32,
	module:     ^IR_Module,
	current_fn: ^IR_Function,

	// Section buffers — concatenated in spec order
	capabilities:  [dynamic]u32,
	extensions:    [dynamic]u32,
	ext_imports:   [dynamic]u32,
	mem_model:     [dynamic]u32,
	entry_points:  [dynamic]u32,
	exec_modes:    [dynamic]u32,
	debug_source:  [dynamic]u32, // OpString, OpSource (before OpName)
	debug_names:   [dynamic]u32, // OpName, OpMemberName
	debug_process: [dynamic]u32, // OpModuleProcessed (after OpName)
	annotations:   [dynamic]u32,
	type_section:  [dynamic]u32, // types + constants + global vars
	func_section:  [dynamic]u32,

	// Caches
	type_cache:    map[string]u32,  // type key string -> type ID
	const_cache:   map[string]u32,  // const key string -> const ID
	ptr_cache:     map[string]u32,  // "storage_class:type_id" -> pointer type ID
	wrapped_bindings: map[string]bool, // bindings wrapped in synthetic Block struct (non-struct uniforms)
	decorated_structs: map[u32]bool, // struct type IDs already decorated with offsets
	ssa_values: map[IR_Var_Id]u32, // let-bound variables kept as SSA values (not Function variables)

	// Tracking
	glsl_ext_id:  u32,
	void_type_id: u32,

	// Function ID tracking (pre-allocated for forward references)
	function_ids:  map[string]u32,   // function name -> function ID

	// Per-function state
	value_map:     map[IR_Var_Id]u32,   // variable id -> pointer ID (function-scoped OpVariable)
	binding_ids:   map[string]u32,   // binding name -> variable ID
	input_ids:     map[string]u32,   // input io name -> variable ID
	output_ids:    map[string]u32,   // output io name -> variable ID
	builtin_input_ids:  map[string]u32, // builtin name -> variable ID
	builtin_output_ids: map[string]u32, // builtin name -> variable ID
	shared_var_ids:     map[string]u32, // shared var name -> variable ID

	// Loop context stack for break/continue
	loop_stack: [dynamic]SPIRV_Loop_Context,

	// Interface variable IDs for OpEntryPoint
	interface_ids: [dynamic]u32,

	// Function-scoped OpVariable buffer (must be emitted first in entry block)
	var_buffer: [dynamic]u32,

	// Debug info
	debug:             bool,
	source_file:       string,
	source_file_id:    u32,
	last_emitted_line: int,

	diagnostics: [dynamic]Diagnostic,
}

emit_spirv :: proc(module: ^IR_Module, debug := false, source_file := "", allocator := context.allocator) -> ([]u8, []Diagnostic) {
	b := SPIRV_Builder{
		next_id      = 1,
		module       = module,
		capabilities  = make([dynamic]u32, allocator),
		extensions    = make([dynamic]u32, allocator),
		ext_imports   = make([dynamic]u32, allocator),
		mem_model     = make([dynamic]u32, allocator),
		entry_points  = make([dynamic]u32, allocator),
		exec_modes    = make([dynamic]u32, allocator),
		debug_source  = make([dynamic]u32, allocator),
		debug_names   = make([dynamic]u32, allocator),
		debug_process = make([dynamic]u32, allocator),
		annotations   = make([dynamic]u32, allocator),
		type_section  = make([dynamic]u32, allocator),
		func_section  = make([dynamic]u32, allocator),
		type_cache   = make(map[string]u32, allocator = allocator),
		const_cache  = make(map[string]u32, allocator = allocator),
		ptr_cache    = make(map[string]u32, allocator = allocator),
		wrapped_bindings = make(map[string]bool, allocator = allocator),
		function_ids = make(map[string]u32, allocator = allocator),
		binding_ids  = make(map[string]u32, allocator = allocator),
		input_ids    = make(map[string]u32, allocator = allocator),
		output_ids   = make(map[string]u32, allocator = allocator),
		builtin_input_ids  = make(map[string]u32, allocator = allocator),
		builtin_output_ids = make(map[string]u32, allocator = allocator),
		shared_var_ids     = make(map[string]u32, allocator = allocator),
		interface_ids = make([dynamic]u32, allocator),
		diagnostics  = make([dynamic]Diagnostic, allocator),
	}

	b.debug = debug
	b.source_file = source_file

	// Preamble
	spirv_emit_capability(&b, SpvCapability_Shader)
	b.glsl_ext_id = spirv_alloc_id(&b)
	spirv_emit_ext_inst_import(&b, b.glsl_ext_id, "GLSL.std.450")
	spirv_emit_memory_model(&b, SpvAddressingModel_Logical, SpvMemoryModel_GLSL450)

	// Debug info (if enabled)
	if b.debug {
		// OpString for source filename (goes before OpName)
		b.source_file_id = spirv_alloc_id(&b)
		file_str := b.source_file if b.source_file != "" else "<unknown>"
		spirv_encode_inst_str(&b.debug_source, SpvOp_String, {b.source_file_id}, file_str)

		// OpSource SourceLanguage_Unknown version file (goes before OpName)
		spirv_encode_inst(&b.debug_source, SpvOp_Source, SpvSourceLanguage_Unknown, 100, b.source_file_id)

		// OpModuleProcessed (goes after OpName/OpMemberName)
		spirv_encode_inst_str(&b.debug_process, SpvOp_ModuleProcessed, {}, "Luma Compiler v0.1")
	}

	// Pre-create void type
	b.void_type_id = spirv_get_or_create_type(&b, nil)

	// Emit specialization constants
	for &sc in module.spec_constants {
		spirv_emit_spec_constant(&b, &sc)
	}

	// Emit global bindings
	for &bind in module.bindings {
		spirv_emit_binding(&b, &bind)
	}

	// Emit shared variables (workgroup memory)
	for sv in module.shared_vars {
		spirv_emit_shared_var(&b, sv)
	}

	// Pre-allocate function IDs for forward references in function calls
	for &fn in module.functions {
		if !fn.is_entry {
			b.function_ids[fn.name] = spirv_alloc_id(&b)
		}
	}

	// Emit functions
	for &fn in module.functions {
		spirv_emit_function(&b, &fn)
	}

	// Assemble final binary
	result := spirv_assemble(&b)
	return result, b.diagnostics[:]
}

// -- ID allocation --

@(private = "file")
spirv_alloc_id :: proc(b: ^SPIRV_Builder) -> u32 {
	id := b.next_id
	b.next_id += 1
	return id
}

// -- Instruction encoding --

@(private = "file")
spirv_encode_inst :: proc(section: ^[dynamic]u32, opcode: u32, operands: ..u32) {
	word_count := u32(1 + len(operands))
	append(section, (word_count << 16) | opcode)
	for op in operands {
		append(section, op)
	}
}

@(private = "file")
spirv_encode_inst_str :: proc(section: ^[dynamic]u32, opcode: u32, pre_operands: []u32, str: string, post_operands: []u32 = {}) {
	// String is null-terminated and padded to 4-byte boundary
	str_words := (len(str) + 4) / 4 // includes null terminator
	word_count := u32(1 + len(pre_operands) + str_words + len(post_operands))
	append(section, (word_count << 16) | opcode)
	for op in pre_operands {
		append(section, op)
	}
	// Encode string as words
	spirv_encode_string(section, str)
	for op in post_operands {
		append(section, op)
	}
}

@(private = "file")
spirv_encode_string :: proc(section: ^[dynamic]u32, str: string) {
	bytes := transmute([]u8)str
	i := 0
	for i + 3 < len(bytes) {
		word := u32(bytes[i]) | (u32(bytes[i+1]) << 8) | (u32(bytes[i+2]) << 16) | (u32(bytes[i+3]) << 24)
		append(section, word)
		i += 4
	}
	// Remaining bytes + null terminator
	word: u32 = 0
	shift: uint = 0
	for i < len(bytes) {
		word |= u32(bytes[i]) << shift
		shift += 8
		i += 1
	}
	// null terminator is already 0 in remaining bits
	append(section, word)
}

// -- Preamble --

@(private = "file")
spirv_emit_capability :: proc(b: ^SPIRV_Builder, cap: u32) {
	spirv_encode_inst(&b.capabilities, SpvOp_Capability, cap)
}

@(private = "file")
spirv_emit_ext_inst_import :: proc(b: ^SPIRV_Builder, result_id: u32, name: string) {
	spirv_encode_inst_str(&b.ext_imports, SpvOp_ExtInstImport, {result_id}, name)
}

@(private = "file")
spirv_emit_memory_model :: proc(b: ^SPIRV_Builder, addressing: u32, memory: u32) {
	spirv_encode_inst(&b.mem_model, SpvOp_MemoryModel, addressing, memory)
}

// -- Type system --

@(private = "file")
spirv_get_or_create_type :: proc(b: ^SPIRV_Builder, t: ^Resolved_Type) -> u32 {
	key := spirv_type_key(t)
	if id, ok := b.type_cache[key]; ok {
		return id
	}

	id := spirv_alloc_id(b)
	b.type_cache[key] = id

	if t == nil {
		spirv_encode_inst(&b.type_section, SpvOp_TypeVoid, id)
		return id
	}

	#partial switch v in t^ {
	case Type_Void:
		spirv_encode_inst(&b.type_section, SpvOp_TypeVoid, id)
	case Type_Scalar:
		switch v.kind {
		case .Bool:
			spirv_encode_inst(&b.type_section, SpvOp_TypeBool, id)
		case .Int:
			spirv_encode_inst(&b.type_section, SpvOp_TypeInt, id, 32, 1)
		case .Uint:
			spirv_encode_inst(&b.type_section, SpvOp_TypeInt, id, 32, 0)
		case .Float:
			spirv_encode_inst(&b.type_section, SpvOp_TypeFloat, id, 32)
		case .Half:
			spirv_encode_inst(&b.type_section, SpvOp_TypeFloat, id, 16)
		}
	case Type_Vector:
		elem_type := make_type(Type_Scalar{kind = v.elem})
		elem_id := spirv_get_or_create_type(b, elem_type)
		spirv_encode_inst(&b.type_section, SpvOp_TypeVector, id, elem_id, u32(v.size))
	case Type_Matrix:
		col_type := make_type(Type_Vector{elem = v.elem, size = v.rows})
		col_id := spirv_get_or_create_type(b, col_type)
		spirv_encode_inst(&b.type_section, SpvOp_TypeMatrix, id, col_id, u32(v.cols))
	case Type_Struct_Resolved:
		member_ids := make([dynamic]u32)
		append(&member_ids, id)
		for f in v.fields {
			append(&member_ids, spirv_get_or_create_type(b, f.type))
		}
		spirv_encode_inst(&b.type_section, SpvOp_TypeStruct, ..member_ids[:])
		// Debug names for struct members
		spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {id}, v.name)
		for f, i in v.fields {
			spirv_encode_inst_str(&b.debug_names, SpvOp_MemberName, {id, u32(i)}, f.name)
		}
	case Type_Array_Resolved:
		elem_id := spirv_get_or_create_type(b, v.elem)
		if v.size > 0 {
			// Fixed-size array: needs a constant for the length
			int_type := spirv_get_or_create_type(b, TYPE_UINT)
			length_id := spirv_get_or_create_const_int(b, int_type, u32(v.size))
			spirv_encode_inst(&b.type_section, SpvOp_TypeArray, id, elem_id, length_id)
		} else {
			spirv_encode_inst(&b.type_section, SpvOp_TypeRuntimeArray, id, elem_id)
		}
	case Type_Sampler:
		// Create image type, then sampled image type
		float_id := spirv_get_or_create_type(b, TYPE_FLOAT)
		image_id := spirv_alloc_id(b)
		dim := spirv_sampler_dim(v.kind)
		depth: u32 = v.kind == .Sampler2DShadow ? 1 : 0
		arrayed: u32 = v.kind == .Sampler2DArray ? 1 : 0
		// OpTypeImage: result_id, sampled_type, dim, depth, arrayed, multisampled, sampled, format
		spirv_encode_inst(&b.type_section, SpvOp_TypeImage, image_id, float_id, dim, depth, arrayed, 0, 1, SpvImageFormat_Unknown)
		spirv_encode_inst(&b.type_section, SpvOp_TypeSampledImage, id, image_id)
	}
	return id
}

@(private = "file")
spirv_type_key :: proc(t: ^Resolved_Type) -> string {
	if t == nil do return "void"
	#partial switch v in t^ {
	case Type_Void:              return "void"
	case Type_Scalar:
		switch v.kind {
		case .Bool:  return "bool"
		case .Int:   return "int32"
		case .Uint:  return "uint32"
		case .Float: return "float32"
		case .Half:  return "float16"
		}
	case Type_Vector:            return fmt.aprintf("vec_%v_%d", v.elem, v.size)
	case Type_Matrix:            return fmt.aprintf("mat_%v_%d_%d", v.elem, v.cols, v.rows)
	case Type_Struct_Resolved:   return fmt.aprintf("struct_%s", v.name)
	case Type_Array_Resolved:    return fmt.aprintf("arr_%s_%d", spirv_type_key(v.elem), v.size)
	case Type_Sampler:           return fmt.aprintf("sampler_%v", v.kind)
	}
	return "unknown"
}

@(private = "file")
spirv_get_ptr_type :: proc(b: ^SPIRV_Builder, storage_class: u32, pointee_type: u32) -> u32 {
	key := fmt.aprintf("%d:%d", storage_class, pointee_type)
	if id, ok := b.ptr_cache[key]; ok {
		return id
	}
	id := spirv_alloc_id(b)
	b.ptr_cache[key] = id
	spirv_encode_inst(&b.type_section, SpvOp_TypePointer, id, storage_class, pointee_type)
	return id
}

@(private = "file")
spirv_sampler_dim :: proc(kind: Sampler_Kind) -> u32 {
	switch kind {
	case .Sampler2D, .Sampler2DShadow, .Sampler2DArray: return SpvDim_2D
	case .Sampler3D:    return SpvDim_3D
	case .SamplerCube:  return SpvDim_Cube
	}
	return SpvDim_2D
}

// -- Constants --

@(private = "file")
spirv_get_or_create_const_float :: proc(b: ^SPIRV_Builder, type_id: u32, value: f64) -> u32 {
	key := fmt.aprintf("f:%d:%v", type_id, value)
	if id, ok := b.const_cache[key]; ok {
		return id
	}
	id := spirv_alloc_id(b)
	b.const_cache[key] = id
	bits := transmute(u32)f32(value)
	spirv_encode_inst(&b.type_section, SpvOp_Constant, type_id, id, bits)
	return id
}

@(private = "file")
spirv_get_or_create_const_int :: proc(b: ^SPIRV_Builder, type_id: u32, value: u32) -> u32 {
	key := fmt.aprintf("i:%d:%d", type_id, value)
	if id, ok := b.const_cache[key]; ok {
		return id
	}
	id := spirv_alloc_id(b)
	b.const_cache[key] = id
	spirv_encode_inst(&b.type_section, SpvOp_Constant, type_id, id, value)
	return id
}

@(private = "file")
spirv_get_or_create_const_bool :: proc(b: ^SPIRV_Builder, type_id: u32, value: bool) -> u32 {
	key := fmt.aprintf("b:%d:%v", type_id, value)
	if id, ok := b.const_cache[key]; ok {
		return id
	}
	id := spirv_alloc_id(b)
	b.const_cache[key] = id
	op := value ? u32(SpvOp_ConstantTrue) : u32(SpvOp_ConstantFalse)
	spirv_encode_inst(&b.type_section, op, type_id, id)
	return id
}

// -- Decorations --

@(private = "file")
spirv_decorate :: proc(b: ^SPIRV_Builder, target: u32, decoration: u32, operands: ..u32) {
	args := make([dynamic]u32)
	append(&args, target, decoration)
	for op in operands {
		append(&args, op)
	}
	spirv_encode_inst(&b.annotations, SpvOp_Decorate, ..args[:])
}

@(private = "file")
spirv_member_decorate :: proc(b: ^SPIRV_Builder, struct_id: u32, member: u32, decoration: u32, operands: ..u32) {
	args := make([dynamic]u32)
	append(&args, struct_id, member, decoration)
	for op in operands {
		append(&args, op)
	}
	spirv_encode_inst(&b.annotations, SpvOp_MemberDecorate, ..args[:])
}

// -- Specialization Constants --

@(private = "file")
spirv_emit_spec_constant :: proc(b: ^SPIRV_Builder, sc: ^IR_Spec_Constant) {
	type_id := spirv_get_or_create_type(b, sc.type)
	sc_id := spirv_alloc_id(b)

	switch v in sc.default_value {
	case bool:
		if v {
			spirv_encode_inst(&b.type_section, SpvOp_SpecConstantTrue, type_id, sc_id)
		} else {
			spirv_encode_inst(&b.type_section, SpvOp_SpecConstantFalse, type_id, sc_id)
		}
	case i64:
		spirv_encode_inst(&b.type_section, SpvOp_SpecConstant, type_id, sc_id, u32(v))
	case f64:
		bits := transmute(u32)f32(v)
		spirv_encode_inst(&b.type_section, SpvOp_SpecConstant, type_id, sc_id, bits)
	case:
		// Default to 0
		spirv_encode_inst(&b.type_section, SpvOp_SpecConstant, type_id, sc_id, 0)
	}

	spirv_decorate(b, sc_id, SpvDecoration_SpecId, u32(sc.spec_id))
	spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {sc_id}, sc.name)

	// Store ID so expressions can reference it
	b.const_cache[fmt.aprintf("spec_%s", sc.name)] = sc_id
}

// -- Bindings --

@(private = "file")
spirv_emit_binding :: proc(b: ^SPIRV_Builder, bind: ^IR_Binding) {
	switch bind.kind {
	case .Uniform:
		if bind.struct_ref != nil {
			// Struct-backed uniform block
			struct_type_id := spirv_get_or_create_type(b, make_type(bind.struct_ref^))
			spirv_decorate(b, struct_type_id, SpvDecoration_Block)
			spirv_decorate_struct_offsets(b, struct_type_id, bind.struct_ref)
			ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Uniform, struct_type_id)
			var_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Uniform)
			spirv_decorate(b, var_id, SpvDecoration_DescriptorSet, u32(bind.group))
			spirv_decorate(b, var_id, SpvDecoration_Binding, u32(bind.binding_num))
			spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, bind.name)
			b.binding_ids[bind.name] = var_id
		} else {
			// Non-struct uniform (e.g. uniform light_dir: vec3) — wrap in synthetic Block struct
			inner_type_id := spirv_get_or_create_type(b, bind.type)
			wrapper_key := fmt.aprintf("_wrap_%s_%d", bind.name, inner_type_id)
			wrapper_type_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.type_section, SpvOp_TypeStruct, wrapper_type_id, inner_type_id)
			spirv_decorate(b, wrapper_type_id, SpvDecoration_Block)
			spirv_member_decorate(b, wrapper_type_id, 0, SpvDecoration_Offset, 0)
			spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {wrapper_type_id}, wrapper_key)
			ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Uniform, wrapper_type_id)
			var_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Uniform)
			spirv_decorate(b, var_id, SpvDecoration_DescriptorSet, u32(bind.group))
			spirv_decorate(b, var_id, SpvDecoration_Binding, u32(bind.binding_num))
			spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, bind.name)
			b.binding_ids[bind.name] = var_id
			b.wrapped_bindings[bind.name] = true
		}
	case .Buffer:
		if bind.struct_ref == nil do return
		struct_type_id := spirv_get_or_create_type(b, make_type(bind.struct_ref^))
		spirv_decorate(b, struct_type_id, SpvDecoration_Block)
		spirv_decorate_struct_offsets(b, struct_type_id, bind.struct_ref)
		ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_StorageBuffer, struct_type_id)
		var_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_StorageBuffer)
		spirv_decorate(b, var_id, SpvDecoration_DescriptorSet, u32(bind.group))
		spirv_decorate(b, var_id, SpvDecoration_Binding, u32(bind.binding_num))
		spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, bind.name)
		b.binding_ids[bind.name] = var_id
	case .Texture:
		// Sampled image variable
		type_id := spirv_get_or_create_type(b, bind.type)
		ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_UniformConstant, type_id)
		var_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_UniformConstant)
		spirv_decorate(b, var_id, SpvDecoration_DescriptorSet, u32(bind.group))
		spirv_decorate(b, var_id, SpvDecoration_Binding, u32(bind.binding_num))
		spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, bind.name)
		b.binding_ids[bind.name] = var_id
	case .Sampler:
		// Skip — combined with texture in SPIR-V sampled image
		// The Texture binding's type is already SampledImage
		return
	case .Push_Constant:
		if bind.struct_ref == nil do return
		struct_type_id := spirv_get_or_create_type(b, make_type(bind.struct_ref^))
		spirv_decorate(b, struct_type_id, SpvDecoration_Block)
		spirv_decorate_struct_offsets(b, struct_type_id, bind.struct_ref)
		ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_PushConstant, struct_type_id)
		var_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_PushConstant)
		// Push constants have no descriptor set or binding decorations
		spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, bind.name)
		b.binding_ids[bind.name] = var_id
	}
}

@(private = "file")
spirv_emit_shared_var :: proc(b: ^SPIRV_Builder, sv: IR_Shared_Var) {
	type_id := spirv_get_or_create_type(b, sv.type)
	ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Workgroup, type_id)
	var_id := spirv_alloc_id(b)
	spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Workgroup)
	b.shared_var_ids[sv.name] = var_id
	spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, sv.name)
}

@(private = "file")
spirv_emit_constant_u32 :: proc(b: ^SPIRV_Builder, value: u32) -> u32 {
	key := fmt.aprintf("u32_%d", value)
	if cached, ok := b.const_cache[key]; ok {
		return cached
	}
	uint_type_id := spirv_get_or_create_type(b, TYPE_UINT)
	const_id := spirv_alloc_id(b)
	spirv_encode_inst(&b.type_section, SpvOp_Constant, uint_type_id, const_id, value)
	b.const_cache[key] = const_id
	return const_id
}

@(private = "file")
spirv_decorate_struct_offsets :: proc(b: ^SPIRV_Builder, struct_type_id: u32, s: ^Type_Struct_Resolved) {
	if struct_type_id in b.decorated_structs do return
	b.decorated_structs[struct_type_id] = true
	offset: u32 = 0
	for f, i in s.fields {
		align := spirv_type_alignment(f.type)
		offset = (offset + align - 1) & ~(align - 1)
		spirv_member_decorate(b, struct_type_id, u32(i), SpvDecoration_Offset, offset)
		// Matrix decorations
		if f.type != nil {
			if m, ok := f.type^.(Type_Matrix); ok {
				spirv_member_decorate(b, struct_type_id, u32(i), SpvDecoration_ColMajor)
				stride := spirv_type_size_scalar(make_type(Type_Vector{elem = m.elem, size = m.rows}))
				stride = (stride + 15) & ~u32(15) // round up to vec4 alignment for std140
				spirv_member_decorate(b, struct_type_id, u32(i), SpvDecoration_MatrixStride, stride)
			}
			// Nested struct: recursively decorate member offsets
			if nested, ok := f.type^.(Type_Struct_Resolved); ok {
				nested_type_id := spirv_get_or_create_type(b, f.type)
				spirv_decorate_struct_offsets(b, nested_type_id, &nested)
			}
			// Array stride decoration (required for arrays in Block-decorated structs)
			if arr, ok := f.type^.(Type_Array_Resolved); ok {
				arr_type_id := spirv_get_or_create_type(b, f.type)
				elem_stride := spirv_type_size(arr.elem)
				// Align to element alignment (std140: round up to 16 for vec/mat, 4 for scalar)
				elem_align := spirv_type_alignment(arr.elem)
				elem_stride = (elem_stride + elem_align - 1) & ~(elem_align - 1)
				spirv_decorate(b, arr_type_id, SpvDecoration_ArrayStride, elem_stride)
				// Decorate element struct members if array of structs
				if nested, ok2 := arr.elem^.(Type_Struct_Resolved); ok2 {
					nested_type_id := spirv_get_or_create_type(b, arr.elem)
					spirv_decorate_struct_offsets(b, nested_type_id, &nested)
				}
				// Matrix decorations for arrays of matrices (e.g. [4]mat4)
				if m, ok2 := arr.elem^.(Type_Matrix); ok2 {
					spirv_member_decorate(b, struct_type_id, u32(i), SpvDecoration_ColMajor)
					stride := spirv_type_size_scalar(make_type(Type_Vector{elem = m.elem, size = m.rows}))
					stride = (stride + 15) & ~u32(15)
					spirv_member_decorate(b, struct_type_id, u32(i), SpvDecoration_MatrixStride, stride)
				}
			}
		}
		offset += spirv_type_size(f.type)
	}
}

@(private = "file")
spirv_type_alignment :: proc(t: ^Resolved_Type) -> u32 {
	if t == nil do return 4
	#partial switch v in t^ {
	case Type_Scalar: return 4
	case Type_Vector:
		switch v.size {
		case 2: return 8
		case 3, 4: return 16
		}
	case Type_Matrix: return 16
	case Type_Struct_Resolved: return 16
	case Type_Array_Resolved: return 16
	case: return 4
	}
	return 4
}

@(private = "file")
spirv_type_size :: proc(t: ^Resolved_Type) -> u32 {
	if t == nil do return 0
	#partial switch v in t^ {
	case Type_Scalar:
		if v.kind == .Half do return 2
		return 4
	case Type_Vector:
		elem_size: u32 = v.elem == .Half ? 2 : 4
		return elem_size * u32(v.size)
	case Type_Matrix:
		col_size := spirv_type_size(make_type(Type_Vector{elem = v.elem, size = v.rows}))
		col_stride := (col_size + 15) & ~u32(15) // std140 column stride
		return col_stride * u32(v.cols)
	case Type_Struct_Resolved:
		total: u32 = 0
		for f in v.fields {
			align := spirv_type_alignment(f.type)
			total = (total + align - 1) & ~(align - 1)
			total += spirv_type_size(f.type)
		}
		return total
	case Type_Array_Resolved:
		elem_sz := spirv_type_size(v.elem)
		elem_align := spirv_type_alignment(v.elem)
		elem_stride := (elem_sz + elem_align - 1) & ~(elem_align - 1)
		return elem_stride * u32(v.size)
	case: return 4
	}
	return 0
}

@(private = "file")
spirv_type_size_scalar :: proc(t: ^Resolved_Type) -> u32 {
	if t == nil do return 0
	#partial switch v in t^ {
	case Type_Scalar:
		if v.kind == .Half do return 2
		return 4
	case Type_Vector:
		return (v.elem == .Half ? 2 : 4) * u32(v.size)
	case: return spirv_type_size(t)
	}
	return 0
}

// -- Functions --

@(private = "file")
spirv_emit_function :: proc(b: ^SPIRV_Builder, fn: ^IR_Function) {
	b.current_fn = fn
	b.value_map = make(map[IR_Var_Id]u32)

	if fn.is_entry {
		spirv_emit_entry_function(b, fn)
	} else {
		spirv_emit_helper_function(b, fn)
	}

	b.current_fn = nil
}

@(private = "file")
spirv_emit_entry_function :: proc(b: ^SPIRV_Builder, fn: ^IR_Function) {
	// Create I/O variables
	b.input_ids = make(map[string]u32)
	b.output_ids = make(map[string]u32)
	b.builtin_input_ids = make(map[string]u32)
	b.builtin_output_ids = make(map[string]u32)
	b.interface_ids = make([dynamic]u32)

	for io in fn.inputs {
		type_id := spirv_get_or_create_type(b, io.type)
		ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Input, type_id)
		var_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Input)
		spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, io.name)

		if io.builtin != "" {
			spirv_decorate(b, var_id, SpvDecoration_BuiltIn, spirv_builtin_id(io.builtin, fn.stage))
			b.builtin_input_ids[io.builtin] = var_id
		} else {
			spirv_decorate(b, var_id, SpvDecoration_Location, u32(io.location))
			b.input_ids[io.name] = var_id
		}
		append(&b.interface_ids, var_id)
	}

	for io in fn.outputs {
		type_id := spirv_get_or_create_type(b, io.type)
		ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Output, type_id)
		var_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.type_section, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Output)
		spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {var_id}, io.name)

		if io.builtin != "" {
			spirv_decorate(b, var_id, SpvDecoration_BuiltIn, spirv_builtin_id(io.builtin, fn.stage))
			b.builtin_output_ids[io.builtin] = var_id
		} else {
			spirv_decorate(b, var_id, SpvDecoration_Location, u32(io.location))
			b.output_ids[io.name] = var_id
		}
		append(&b.interface_ids, var_id)
	}

	// Add binding variables to interface list (required by SPIR-V 1.4+)
	for name, var_id in b.binding_ids {
		append(&b.interface_ids, var_id)
	}
	// Add shared variables to interface list
	for name, var_id in b.shared_var_ids {
		append(&b.interface_ids, var_id)
	}

	// Function type: void(void)
	fn_type_id := spirv_get_fn_type(b, b.void_type_id, {})
	fn_id := spirv_alloc_id(b)
	spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {fn_id}, fn.name)

	// Entry point declaration
	exec_model := spirv_execution_model(fn.stage)
	spirv_emit_entry_point_inst(b, exec_model, fn_id, fn.name, b.interface_ids[:])

	// Execution mode
	if fn.stage == .Fragment {
		spirv_encode_inst(&b.exec_modes, SpvOp_ExecutionMode, fn_id, SpvExecutionMode_OriginUpperLeft)
	}
	if fn.stage == .Compute {
		ws := fn.workgroup_size
		spirv_encode_inst(&b.exec_modes, SpvOp_ExecutionMode, fn_id, SpvExecutionMode_LocalSize, u32(ws[0] > 0 ? ws[0] : 1), u32(ws[1] > 0 ? ws[1] : 1), u32(ws[2] > 0 ? ws[2] : 1))
	}

	// Function definition
	spirv_encode_inst(&b.func_section, SpvOp_Function, b.void_type_id, fn_id, SpvFunctionControl_None, fn_type_id)

	// Entry label
	entry_label := spirv_alloc_id(b)
	spirv_encode_inst(&b.func_section, SpvOp_Label, entry_label)

	// Reset var buffer for this function
	b.var_buffer = make([dynamic]u32)

	// Save position after label — var_buffer will be inserted here
	label_end := len(b.func_section)

	// Emit body
	spirv_emit_stmts(b, fn.body[:])

	// Splice var_buffer right after the label (before body instructions)
	if len(b.var_buffer) > 0 {
		spirv_splice_vars(b, label_end)
	}

	// Implicit return (skip if body already terminates, e.g. discard/OpKill)
	if !spirv_block_has_terminator(fn.body[:]) {
		spirv_encode_inst(&b.func_section, SpvOp_Return)
	}
	spirv_encode_inst(&b.func_section, SpvOp_FunctionEnd)
}

@(private = "file")
spirv_emit_entry_point_inst :: proc(b: ^SPIRV_Builder, exec_model: u32, fn_id: u32, name: string, interfaces: []u32) {
	str_words := (len(name) + 4) / 4
	word_count := u32(3 + str_words + len(interfaces))
	append(&b.entry_points, (word_count << 16) | u32(SpvOp_EntryPoint))
	append(&b.entry_points, exec_model)
	append(&b.entry_points, fn_id)
	spirv_encode_string(&b.entry_points, name)
	for iface in interfaces {
		append(&b.entry_points, iface)
	}
}

@(private = "file")
spirv_emit_helper_function :: proc(b: ^SPIRV_Builder, fn: ^IR_Function) {
	// Build parameter type list
	param_type_ids := make([dynamic]u32)
	for p in fn.params {
		append(&param_type_ids, spirv_get_or_create_type(b, p.type))
	}

	ret_type_id := spirv_get_or_create_type(b, fn.return_type)
	fn_type_id := spirv_get_fn_type(b, ret_type_id, param_type_ids[:])
	// Use pre-allocated function ID for forward reference support
	fn_id := b.function_ids[fn.name]
	spirv_encode_inst_str(&b.debug_names, SpvOp_Name, {fn_id}, fn.name)

	spirv_encode_inst(&b.func_section, SpvOp_Function, ret_type_id, fn_id, SpvFunctionControl_None, fn_type_id)

	// Parameters — collect IDs, then copy into function-scoped variables after label
	Param_Info :: struct { id: u32, type_id: u32, var_id: IR_Var_Id }
	param_infos := make([dynamic]Param_Info)
	for p in fn.params {
		param_id := spirv_alloc_id(b)
		param_type_id := spirv_get_or_create_type(b, p.type)
		spirv_encode_inst(&b.func_section, SpvOp_FunctionParameter, param_type_id, param_id)
		append(&param_infos, Param_Info{id = param_id, type_id = param_type_id, var_id = p.id})
	}

	// Entry label
	entry_label := spirv_alloc_id(b)
	spirv_encode_inst(&b.func_section, SpvOp_Label, entry_label)

	b.var_buffer = make([dynamic]u32)
	label_end := len(b.func_section)

	// Copy parameters into function-scoped variables (SPIR-V requires OpLoad from pointers)
	for pi in param_infos {
		ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Function, pi.type_id)
		var_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.var_buffer, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Function)
		spirv_encode_inst(&b.func_section, SpvOp_Store, var_id, pi.id)
		b.value_map[pi.var_id] = var_id
	}

	spirv_emit_stmts(b, fn.body[:])

	if len(b.var_buffer) > 0 {
		spirv_splice_vars(b, label_end)
	}

	if fn.return_type == nil && !spirv_block_has_terminator(fn.body[:]) {
		spirv_encode_inst(&b.func_section, SpvOp_Return)
	}
	spirv_encode_inst(&b.func_section, SpvOp_FunctionEnd)
}

@(private = "file")
spirv_get_fn_type :: proc(b: ^SPIRV_Builder, ret_type: u32, param_types: []u32) -> u32 {
	key := fmt.aprintf("fn_%d", ret_type)
	for pt in param_types {
		key = fmt.aprintf("%s_%d", key, pt)
	}
	if id, ok := b.type_cache[key]; ok {
		return id
	}
	id := spirv_alloc_id(b)
	b.type_cache[key] = id
	args := make([dynamic]u32)
	append(&args, id, ret_type)
	for pt in param_types {
		append(&args, pt)
	}
	spirv_encode_inst(&b.type_section, SpvOp_TypeFunction, ..args[:])
	return id
}

// -- Statements --

@(private = "file")
spirv_emit_stmts :: proc(b: ^SPIRV_Builder, stmts: []IR_Stmt) {
	for stmt in stmts {
		spirv_emit_stmt(b, stmt)
	}
}

@(private = "file")
spirv_emit_stmt :: proc(b: ^SPIRV_Builder, stmt: IR_Stmt) {
	// Emit OpLine debug info if enabled
	if b.debug && b.source_file_id != 0 {
		span := ir_stmt_span(stmt)
		if span.line_start > 0 && span.line_start != b.last_emitted_line {
			spirv_encode_inst(&b.func_section, SpvOp_Line,
				b.source_file_id, u32(span.line_start), u32(max(span.col_start - 1, 0)))
			b.last_emitted_line = span.line_start
		}
	}

	switch s in stmt {
	case ^IR_Let:
		val_id := spirv_emit_expr(b, s.value)
		// Create function-scoped variable (deferred to entry block) and store
		type_id := spirv_get_or_create_type(b, s.type)
		if !s.mutable && type_id in b.decorated_structs {
			// Immutable struct types with Offset decorations can't be used in Function storage.
			// Keep as SSA value; field accesses will use OpCompositeExtract.
			b.ssa_values[s.id] = val_id
		} else {
			ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Function, type_id)
			var_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.var_buffer, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Function)
			spirv_encode_inst(&b.func_section, SpvOp_Store, var_id, val_id)
			b.value_map[s.id] = var_id
		}

	case ^IR_Assign:
		val_id := spirv_emit_expr(b, s.value)
		ptr_id := spirv_emit_lvalue_ptr(b, s.target)
		if ptr_id != 0 {
			spirv_encode_inst(&b.func_section, SpvOp_Store, ptr_id, val_id)
		}

	case ^IR_Return:
		if s.value != nil {
			val_id := spirv_emit_expr(b, s.value)
			spirv_encode_inst(&b.func_section, SpvOp_ReturnValue, val_id)
		} else {
			spirv_encode_inst(&b.func_section, SpvOp_Return)
		}

	case ^IR_Store_Output:
		val_id := spirv_emit_expr(b, s.value)
		fn := b.current_fn
		if fn != nil && s.io_index >= 0 && s.io_index < len(fn.outputs) {
			io := fn.outputs[s.io_index]
			var_id: u32
			if io.builtin != "" {
				var_id = b.builtin_output_ids[io.builtin]
			} else {
				var_id = b.output_ids[io.name]
			}
			if var_id != 0 {
				spirv_encode_inst(&b.func_section, SpvOp_Store, var_id, val_id)
			}
		}

	case ^IR_If:
		spirv_emit_if(b, s)

	case ^IR_For:
		spirv_emit_for(b, s)

	case ^IR_While:
		spirv_emit_while(b, s)

	case ^IR_Expr_Stmt:
		spirv_emit_expr(b, s.expr) // result unused

	case ^IR_Barrier:
		// OpControlBarrier execution=Workgroup, memory=Workgroup, semantics=WorkgroupMemory|AcquireRelease
		scope_wg_id := spirv_emit_constant_u32(b, u32(SpvScope_Workgroup))
		semantics_id := spirv_emit_constant_u32(b, u32(SpvMemorySemantics_WorkgroupMemory | SpvMemorySemantics_AcquireRelease))
		spirv_encode_inst(&b.func_section, SpvOp_ControlBarrier, scope_wg_id, scope_wg_id, semantics_id)

	case ^IR_Discard:
		spirv_encode_inst(&b.func_section, SpvOp_Kill)

	case ^IR_Break:
		ctx := b.loop_stack[len(b.loop_stack) - 1]
		spirv_encode_inst(&b.func_section, SpvOp_Branch, ctx.merge_label)

	case ^IR_Continue:
		ctx := b.loop_stack[len(b.loop_stack) - 1]
		spirv_encode_inst(&b.func_section, SpvOp_Branch, ctx.continue_label)
	}
}

// -- Control flow --

@(private = "file")
spirv_emit_if :: proc(b: ^SPIRV_Builder, s: ^IR_If) {
	cond_id := spirv_emit_expr(b, s.condition)

	then_label := spirv_alloc_id(b)
	else_label := spirv_alloc_id(b)
	merge_label := spirv_alloc_id(b)

	has_else := len(s.else_body) > 0 || len(s.elseif_clauses) > 0
	false_label := has_else ? else_label : merge_label

	spirv_encode_inst(&b.func_section, SpvOp_SelectionMerge, merge_label, SpvSelectionControl_None)
	spirv_encode_inst(&b.func_section, SpvOp_BranchConditional, cond_id, then_label, false_label)

	// Then block
	spirv_encode_inst(&b.func_section, SpvOp_Label, then_label)
	spirv_emit_stmts(b, s.then_body[:])
	if !spirv_block_has_terminator(s.then_body[:]) {
		spirv_encode_inst(&b.func_section, SpvOp_Branch, merge_label)
	}

	// Else block
	if has_else {
		spirv_encode_inst(&b.func_section, SpvOp_Label, else_label)
		// Handle elseif chains as nested ifs
		if len(s.elseif_clauses) > 0 {
			spirv_emit_elseif_chain(b, s.elseif_clauses[:], s.else_body[:], merge_label)
		} else {
			spirv_emit_stmts(b, s.else_body[:])
		}
		// For elseif chains, the chain handles its own branching; for plain else, check body
		if len(s.elseif_clauses) == 0 {
			if !spirv_block_has_terminator(s.else_body[:]) {
				spirv_encode_inst(&b.func_section, SpvOp_Branch, merge_label)
			}
		}
	}

	// Merge block
	spirv_encode_inst(&b.func_section, SpvOp_Label, merge_label)
}

@(private = "file")
spirv_emit_elseif_chain :: proc(b: ^SPIRV_Builder, clauses: []IR_Elseif, else_body: []IR_Stmt, outer_merge: u32) {
	if len(clauses) == 0 {
		spirv_emit_stmts(b, else_body[:])
		if !spirv_block_has_terminator(else_body[:]) {
			spirv_encode_inst(&b.func_section, SpvOp_Branch, outer_merge)
		}
		return
	}

	clause := clauses[0]
	cond_id := spirv_emit_expr(b, clause.condition)

	then_label := spirv_alloc_id(b)
	else_label := spirv_alloc_id(b)
	merge_label := spirv_alloc_id(b)

	has_more := len(clauses) > 1 || len(else_body) > 0

	spirv_encode_inst(&b.func_section, SpvOp_SelectionMerge, merge_label, SpvSelectionControl_None)
	spirv_encode_inst(&b.func_section, SpvOp_BranchConditional, cond_id, then_label, has_more ? else_label : merge_label)

	spirv_encode_inst(&b.func_section, SpvOp_Label, then_label)
	spirv_emit_stmts(b, clause.body[:])
	if !spirv_block_has_terminator(clause.body[:]) {
		spirv_encode_inst(&b.func_section, SpvOp_Branch, merge_label)
	}

	if has_more {
		spirv_encode_inst(&b.func_section, SpvOp_Label, else_label)
		spirv_emit_elseif_chain(b, clauses[1:], else_body, merge_label)
	}

	// Merge block — branch to the outer merge so control flow propagates up
	spirv_encode_inst(&b.func_section, SpvOp_Label, merge_label)
	spirv_encode_inst(&b.func_section, SpvOp_Branch, outer_merge)
}

@(private = "file")
spirv_emit_for :: proc(b: ^SPIRV_Builder, s: ^IR_For) {
	// Initialize loop variable
	start_id := spirv_emit_expr(b, s.start)
	loop_type := s.start != nil ? s.start.type : TYPE_INT
	is_uint := spirv_is_uint_type(loop_type)
	int_type_id := spirv_get_or_create_type(b, loop_type if loop_type != nil else TYPE_INT)
	ptr_type_id := spirv_get_ptr_type(b, SpvStorageClass_Function, int_type_id)
	var_id := spirv_alloc_id(b)
	spirv_encode_inst(&b.var_buffer, SpvOp_Variable, ptr_type_id, var_id, SpvStorageClass_Function)
	spirv_encode_inst(&b.func_section, SpvOp_Store, var_id, start_id)
	b.value_map[s.var_id] = var_id

	header_label := spirv_alloc_id(b)
	cond_label := spirv_alloc_id(b)
	body_label := spirv_alloc_id(b)
	continue_label := spirv_alloc_id(b)
	merge_label := spirv_alloc_id(b)

	spirv_encode_inst(&b.func_section, SpvOp_Branch, header_label)

	// Header — OpLoopMerge must be second-to-last, followed only by a branch
	spirv_encode_inst(&b.func_section, SpvOp_Label, header_label)
	spirv_encode_inst(&b.func_section, SpvOp_LoopMerge, merge_label, continue_label, SpvLoopControl_None)
	spirv_encode_inst(&b.func_section, SpvOp_Branch, cond_label)

	// Condition: var <= stop
	spirv_encode_inst(&b.func_section, SpvOp_Label, cond_label)
	cur_val := spirv_alloc_id(b)
	spirv_encode_inst(&b.func_section, SpvOp_Load, int_type_id, cur_val, var_id)
	stop_id := spirv_emit_expr(b, s.stop)
	bool_type_id := spirv_get_or_create_type(b, TYPE_BOOL)
	cond_id := spirv_alloc_id(b)
	cmp_op := is_uint ? u32(SpvOp_ULessThanEqual) : u32(SpvOp_SLessThanEqual)
	spirv_encode_inst(&b.func_section, cmp_op, bool_type_id, cond_id, cur_val, stop_id)
	spirv_encode_inst(&b.func_section, SpvOp_BranchConditional, cond_id, body_label, merge_label)

	// Body
	spirv_encode_inst(&b.func_section, SpvOp_Label, body_label)
	append(&b.loop_stack, SPIRV_Loop_Context{continue_label, merge_label})
	spirv_emit_stmts(b, s.body[:])
	pop(&b.loop_stack)
	if !spirv_block_has_terminator(s.body[:]) {
		spirv_encode_inst(&b.func_section, SpvOp_Branch, continue_label)
	}

	// Continue: increment
	spirv_encode_inst(&b.func_section, SpvOp_Label, continue_label)
	cur_val2 := spirv_alloc_id(b)
	spirv_encode_inst(&b.func_section, SpvOp_Load, int_type_id, cur_val2, var_id)
	step_id: u32
	if s.step != nil {
		step_id = spirv_emit_expr(b, s.step)
	} else {
		step_id = spirv_get_or_create_const_int(b, int_type_id, 1)
	}
	next_val := spirv_alloc_id(b)
	spirv_encode_inst(&b.func_section, SpvOp_IAdd, int_type_id, next_val, cur_val2, step_id)
	spirv_encode_inst(&b.func_section, SpvOp_Store, var_id, next_val)
	spirv_encode_inst(&b.func_section, SpvOp_Branch, header_label)

	// Merge
	spirv_encode_inst(&b.func_section, SpvOp_Label, merge_label)
}

@(private = "file")
spirv_emit_while :: proc(b: ^SPIRV_Builder, s: ^IR_While) {
	header_label := spirv_alloc_id(b)
	cond_label := spirv_alloc_id(b)
	body_label := spirv_alloc_id(b)
	continue_label := spirv_alloc_id(b)
	merge_label := spirv_alloc_id(b)

	spirv_encode_inst(&b.func_section, SpvOp_Branch, header_label)

	// Header — OpLoopMerge must be second-to-last, followed only by a branch
	spirv_encode_inst(&b.func_section, SpvOp_Label, header_label)
	spirv_encode_inst(&b.func_section, SpvOp_LoopMerge, merge_label, continue_label, SpvLoopControl_None)
	spirv_encode_inst(&b.func_section, SpvOp_Branch, cond_label)

	// Condition
	spirv_encode_inst(&b.func_section, SpvOp_Label, cond_label)
	cond_id := spirv_emit_expr(b, s.condition)
	spirv_encode_inst(&b.func_section, SpvOp_BranchConditional, cond_id, body_label, merge_label)

	// Body
	spirv_encode_inst(&b.func_section, SpvOp_Label, body_label)
	append(&b.loop_stack, SPIRV_Loop_Context{continue_label, merge_label})
	spirv_emit_stmts(b, s.body[:])
	pop(&b.loop_stack)
	if !spirv_block_has_terminator(s.body[:]) {
		spirv_encode_inst(&b.func_section, SpvOp_Branch, continue_label)
	}

	// Continue
	spirv_encode_inst(&b.func_section, SpvOp_Label, continue_label)
	spirv_encode_inst(&b.func_section, SpvOp_Branch, header_label)

	// Merge
	spirv_encode_inst(&b.func_section, SpvOp_Label, merge_label)
}

// -- Expressions --

@(private = "file")
spirv_emit_expr :: proc(b: ^SPIRV_Builder, expr: ^IR_Expr) -> u32 {
	if expr == nil do return 0

	result_type_id := spirv_get_or_create_type(b, expr.type)

	switch d in expr.derived {
	case ^IR_Literal:
		return spirv_emit_literal(b, d, result_type_id)

	case ^IR_Var_Ref:
		// Check for spec constant first
		if sc_id, ok := b.const_cache[fmt.aprintf("spec_%s", d.name)]; ok {
			return sc_id
		}
		// Check for SSA value (let-bound decorated structs)
		if val_id, ok := b.ssa_values[d.id]; ok {
			return val_id
		}
		if ptr_id, ok := b.value_map[d.id]; ok {
			// Load from function-scoped variable
			loaded_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, ptr_id)
			return loaded_id
		}
		return 0

	case ^IR_Binary:
		left_id := spirv_emit_expr(b, d.left)
		right_id := spirv_emit_expr(b, d.right)
		result_id := spirv_alloc_id(b)
		opcode := spirv_binary_opcode(d.op, expr.type, d.left.type, d.right.type)
		// SPIR-V VectorTimesScalar requires (vector, scalar) order
		l_id, r_id := left_id, right_id
		if opcode == SpvOp_VectorTimesScalar || opcode == SpvOp_MatrixTimesScalar {
			if _, is_scalar := d.left.type^.(Type_Scalar); is_scalar {
				l_id, r_id = right_id, left_id
			}
		}
		// SPIR-V requires matching types for FDiv/FAdd/FSub etc — splat scalar to vector
		// Skip for opcodes that natively handle scalar operands
		if opcode != SpvOp_VectorTimesScalar && opcode != SpvOp_MatrixTimesScalar {
			if result_vec, is_vec := expr.type^.(Type_Vector); is_vec {
				if is_scalar(d.left.type) {
					l_id = spirv_splat_scalar(b, l_id, result_type_id, result_vec.size)
				}
				if is_scalar(d.right.type) {
					r_id = spirv_splat_scalar(b, r_id, result_type_id, result_vec.size)
				}
			}
		}
		spirv_encode_inst(&b.func_section, opcode, result_type_id, result_id, l_id, r_id)
		return result_id

	case ^IR_Unary:
		operand_id := spirv_emit_expr(b, d.operand)
		result_id := spirv_alloc_id(b)
		opcode: u32
		if d.op == .Neg {
			opcode = spirv_is_float_type(expr.type) ? u32(SpvOp_FNegate) : u32(SpvOp_SNegate)
		} else {
			opcode = SpvOp_LogicalNot
		}
		spirv_encode_inst(&b.func_section, opcode, result_type_id, result_id, operand_id)
		return result_id

	case ^IR_Call:
		return spirv_emit_call(b, d, expr)

	case ^IR_Field_Access:
		// Memory path: AccessChain + Load (binding field access)
		obj_id := spirv_emit_expr_ptr(b, d.object)
		if obj_id == 0 {
			// Fallback: treat as composite extract
			obj_val := spirv_emit_expr(b, d.object)
			idx := resolve_field_index(d.object.type, d.field_name)
			result_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_CompositeExtract, result_type_id, result_id, obj_val, u32(idx))
			return result_id
		}
		idx := resolve_field_index(d.object.type, d.field_name)
		int_type_id := spirv_get_or_create_type(b, TYPE_UINT)
		idx_const := spirv_get_or_create_const_int(b, int_type_id, u32(idx))
		sc := spirv_storage_class_for_expr(b, d.object)
		ptr_type_id := spirv_get_ptr_type(b, sc, result_type_id)
		chain_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_AccessChain, ptr_type_id, chain_id, obj_id, idx_const)
		loaded_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, chain_id)
		return loaded_id

	case ^IR_Swizzle:
		obj_id := spirv_emit_expr(b, d.object)
		// Single component -> CompositeExtract
		if len(d.components) == 1 {
			idx := spirv_swizzle_index(d.components[0])
			result_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_CompositeExtract, result_type_id, result_id, obj_id, u32(idx))
			return result_id
		}
		// Multi-component -> VectorShuffle
		result_id := spirv_alloc_id(b)
		args := make([dynamic]u32)
		append(&args, result_type_id, result_id, obj_id, obj_id) // two input vectors (same)
		for ch in d.components {
			append(&args, u32(spirv_swizzle_index(u8(ch))))
		}
		spirv_encode_inst(&b.func_section, SpvOp_VectorShuffle, ..args[:])
		return result_id

	case ^IR_Composite_Extract:
		obj_id := spirv_emit_expr(b, d.object)
		result_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_CompositeExtract, result_type_id, result_id, obj_id, u32(d.index))
		return result_id

	case ^IR_Vector_Shuffle:
		obj_id := spirv_emit_expr(b, d.object)
		result_id := spirv_alloc_id(b)
		args := make([dynamic]u32)
		append(&args, result_type_id, result_id, obj_id, obj_id)
		for idx in d.components {
			append(&args, u32(idx))
		}
		spirv_encode_inst(&b.func_section, SpvOp_VectorShuffle, ..args[:])
		return result_id

	case ^IR_Index:
		// Check if the object is memory-backed (shared, binding, var) — use AccessChain + Load
		base_ptr := spirv_emit_expr_ptr(b, d.object)
		if base_ptr != 0 {
			idx_id := spirv_emit_expr(b, d.index)
			sc := spirv_storage_class_for_expr(b, d.object)
			ptr_type_id := spirv_get_ptr_type(b, sc, result_type_id)
			chain_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_AccessChain, ptr_type_id, chain_id, base_ptr, idx_id)
			loaded_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, chain_id)
			return loaded_id
		}
		// Fallback: value-based extraction
		obj_id := spirv_emit_expr(b, d.object)
		idx_id := spirv_emit_expr(b, d.index)
		result_id := spirv_alloc_id(b)
		// Use VectorExtractDynamic for vector types (supports runtime index)
		spirv_encode_inst(&b.func_section, SpvOp_VectorExtractDynamic, result_type_id, result_id, obj_id, idx_id)
		return result_id

	case ^IR_Construct:
		result_id := spirv_alloc_id(b)
		target_type := expr.type^

		// Special case: mat3(mat4) — extract first 3 columns and truncate each vec4 to vec3
		if target_mat, ok := target_type.(Type_Matrix); ok && len(d.args) == 1 {
			arg := d.args[0]
			if src_mat, ok2 := arg.type^.(Type_Matrix); ok2 && src_mat.cols >= target_mat.cols && src_mat.rows >= target_mat.rows && !type_equals(arg.type, expr.type) {
				src_id := spirv_emit_expr(b, arg)
				col_type_id := spirv_get_or_create_type(b, make_type(Type_Vector{target_mat.elem, target_mat.rows}))
				src_col_type_id := spirv_get_or_create_type(b, make_type(Type_Vector{src_mat.elem, src_mat.rows}))
				col_ids := make([dynamic]u32)
				for c in 0 ..< target_mat.cols {
					// Extract column (vec4) from source matrix
					ext_id := spirv_alloc_id(b)
					spirv_encode_inst(&b.func_section, SpvOp_CompositeExtract, src_col_type_id, ext_id, src_id, u32(c))
					if src_mat.rows > target_mat.rows {
						// Truncate vec4 to vec3 via VectorShuffle
						trunc_id := spirv_alloc_id(b)
						shuffle_args := make([dynamic]u32)
						append(&shuffle_args, col_type_id, trunc_id, ext_id, ext_id) // two operands (same)
						for i in 0 ..< target_mat.rows {
							append(&shuffle_args, u32(i))
						}
						spirv_encode_inst(&b.func_section, SpvOp_VectorShuffle, ..shuffle_args[:])
						append(&col_ids, trunc_id)
					} else {
						append(&col_ids, ext_id)
					}
				}
				construct_args := make([dynamic]u32)
				append(&construct_args, result_type_id, result_id)
				for id in col_ids {
					append(&construct_args, id)
				}
				spirv_encode_inst(&b.func_section, SpvOp_CompositeConstruct, ..construct_args[:])
				return result_id
			}
		}

		// Special case: vecN(scalar) — splat scalar to N components
		if target_vec, ok := target_type.(Type_Vector); ok && len(d.args) == 1 {
			arg := d.args[0]
			if is_scalar(arg.type) {
				scalar_id := spirv_emit_expr(b, arg)
				splat_args := make([dynamic]u32)
				append(&splat_args, result_type_id, result_id)
				for _ in 0 ..< target_vec.size {
					append(&splat_args, scalar_id)
				}
				spirv_encode_inst(&b.func_section, SpvOp_CompositeConstruct, ..splat_args[:])
				return result_id
			}
		}

		// Default: pass args directly to OpCompositeConstruct
		args := make([dynamic]u32)
		append(&args, result_type_id)
		append(&args, result_id)
		for arg in d.args {
			append(&args, spirv_emit_expr(b, arg))
		}
		spirv_encode_inst(&b.func_section, SpvOp_CompositeConstruct, ..args[:])
		return result_id

	case ^IR_Type_Cast:
		val_id := spirv_emit_expr(b, d.value)
		result_id := spirv_alloc_id(b)
		opcode := spirv_cast_opcode(d.value.type, expr.type)
		spirv_encode_inst(&b.func_section, opcode, result_type_id, result_id, val_id)
		return result_id

	case ^IR_Load_Binding:
		// Look up by name or combined_name (for split sampler bindings)
		var_id: u32
		found := false
		if id, ok := b.binding_ids[d.name]; ok {
			var_id = id
			found = true
		} else {
			// Search by combined_name (e.g., "material_tex" -> "material_tex_tex")
			for &bind in b.module.bindings {
				if bind.combined_name == d.name && bind.kind == .Texture {
					if id2, ok2 := b.binding_ids[bind.name]; ok2 {
						var_id = id2
						found = true
						break
					}
				}
			}
		}
		if found {
			// For uniform blocks, return the pointer (field access will use AccessChain)
			for &bind in b.module.bindings {
				if bind.name == d.name {
					if bind.kind == .Uniform || bind.kind == .Buffer || bind.kind == .Push_Constant {
						// Non-struct uniforms wrapped in synthetic struct: AccessChain to member 0, then Load
						if b.wrapped_bindings[d.name] {
							sc := bind.kind == .Buffer ? u32(SpvStorageClass_StorageBuffer) : u32(SpvStorageClass_Uniform)
							member_ptr_type_id := spirv_get_ptr_type(b, sc, result_type_id)
							zero_id := spirv_get_or_create_const_int(b, spirv_get_or_create_type(b, make_type(Type_Scalar{kind = .Uint})), 0)
							chain_id := spirv_alloc_id(b)
							spirv_encode_inst(&b.func_section, SpvOp_AccessChain, member_ptr_type_id, chain_id, var_id, zero_id)
							loaded_id := spirv_alloc_id(b)
							spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, chain_id)
							return loaded_id
						}
						return var_id // return pointer for AccessChain
					}
					break
				}
			}
			loaded_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, var_id)
			return loaded_id
		}
		return 0

	case ^IR_Input_Field:
		if var_id, ok := b.input_ids[d.field_name]; ok {
			loaded_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, var_id)
			return loaded_id
		}
		return 0

	case ^IR_Builtin_Var:
		ids := d.is_input ? b.builtin_input_ids : b.builtin_output_ids
		if var_id, ok := ids[d.name]; ok {
			if !d.is_input {
				return var_id // return pointer for stores
			}
			loaded_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, var_id)
			return loaded_id
		}
		return 0

	case ^IR_Shared_Ref:
		if var_id, ok := b.shared_var_ids[d.name]; ok {
			loaded_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Load, result_type_id, loaded_id, var_id)
			return loaded_id
		}
		return 0

	case ^IR_Select:
		cond_id := spirv_emit_expr(b, d.condition)
		true_id := spirv_emit_expr(b, d.true_val)
		false_id := spirv_emit_expr(b, d.false_val)
		result_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_Select, result_type_id, result_id, cond_id, true_id, false_id)
		return result_id
	}
	return 0
}

@(private = "file")
spirv_emit_expr_ptr :: proc(b: ^SPIRV_Builder, expr: ^IR_Expr) -> u32 {
	// Returns a pointer ID for expressions that are memory-backed
	if expr == nil do return 0
	#partial switch d in expr.derived {
	case ^IR_Load_Binding:
		if var_id, ok := b.binding_ids[d.name]; ok {
			return var_id
		}
	case ^IR_Var_Ref:
		if ptr_id, ok := b.value_map[d.id]; ok {
			return ptr_id
		}
	case ^IR_Shared_Ref:
		if var_id, ok := b.shared_var_ids[d.name]; ok {
			return var_id
		}
	case ^IR_Field_Access:
		base_ptr := spirv_emit_expr_ptr(b, d.object)
		if base_ptr == 0 do return 0
		idx := resolve_field_index(d.object.type, d.field_name)
		int_type_id := spirv_get_or_create_type(b, TYPE_UINT)
		idx_const := spirv_get_or_create_const_int(b, int_type_id, u32(idx))
		sc := spirv_storage_class_for_expr(b, d.object)
		field_type_id := spirv_get_or_create_type(b, expr.type)
		ptr_type_id := spirv_get_ptr_type(b, sc, field_type_id)
		chain_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_AccessChain, ptr_type_id, chain_id, base_ptr, idx_const)
		return chain_id
	}
	return 0
}

// Returns a pointer ID for an lvalue expression (assignment target).
// Handles var refs, shared refs, indexed shared/binding access, and field access on bindings.
@(private = "file")
spirv_emit_lvalue_ptr :: proc(b: ^SPIRV_Builder, expr: ^IR_Expr) -> u32 {
	if expr == nil do return 0
	#partial switch d in expr.derived {
	case ^IR_Var_Ref:
		if ptr_id, ok := b.value_map[d.id]; ok {
			return ptr_id
		}
	case ^IR_Shared_Ref:
		if var_id, ok := b.shared_var_ids[d.name]; ok {
			return var_id
		}
	case ^IR_Index:
		// array[idx] — AccessChain from base pointer
		base_ptr := spirv_emit_lvalue_ptr(b, d.object)
		if base_ptr == 0 do return 0
		idx_id := spirv_emit_expr(b, d.index)
		sc := spirv_storage_class_for_expr(b, d.object)
		elem_type_id := spirv_get_or_create_type(b, expr.type)
		ptr_type_id := spirv_get_ptr_type(b, sc, elem_type_id)
		chain_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_AccessChain, ptr_type_id, chain_id, base_ptr, idx_id)
		return chain_id
	case ^IR_Field_Access:
		base_ptr := spirv_emit_lvalue_ptr(b, d.object)
		if base_ptr == 0 do return 0
		idx := resolve_field_index(d.object.type, d.field_name)
		int_type_id := spirv_get_or_create_type(b, TYPE_UINT)
		idx_const := spirv_get_or_create_const_int(b, int_type_id, u32(idx))
		sc := spirv_storage_class_for_expr(b, d.object)
		field_type_id := spirv_get_or_create_type(b, expr.type)
		ptr_type_id := spirv_get_ptr_type(b, sc, field_type_id)
		chain_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_AccessChain, ptr_type_id, chain_id, base_ptr, idx_const)
		return chain_id
	case ^IR_Load_Binding:
		if var_id, ok := b.binding_ids[d.name]; ok {
			return var_id
		}
	}
	return 0
}

// Determine the SPIR-V storage class for a given expression's base.
@(private = "file")
spirv_storage_class_for_expr :: proc(b: ^SPIRV_Builder, expr: ^IR_Expr) -> u32 {
	if expr == nil do return SpvStorageClass_Function
	#partial switch d in expr.derived {
	case ^IR_Shared_Ref:
		return SpvStorageClass_Workgroup
	case ^IR_Load_Binding:
		for &bind in b.module.bindings {
			if bind.name == d.name {
				if bind.kind == .Buffer {
					return SpvStorageClass_StorageBuffer
				}
				if bind.kind == .Push_Constant {
					return SpvStorageClass_PushConstant
				}
				return SpvStorageClass_Uniform
			}
		}
		return SpvStorageClass_Uniform
	case ^IR_Var_Ref:
		return SpvStorageClass_Function
	case ^IR_Field_Access:
		return spirv_storage_class_for_expr(b, d.object)
	case ^IR_Index:
		return spirv_storage_class_for_expr(b, d.object)
	}
	return SpvStorageClass_Function
}

@(private = "file")
spirv_emit_literal :: proc(b: ^SPIRV_Builder, lit: ^IR_Literal, type_id: u32) -> u32 {
	switch v in lit.value {
	case f64:
		return spirv_get_or_create_const_float(b, type_id, v)
	case i64:
		return spirv_get_or_create_const_int(b, type_id, u32(v))
	case bool:
		return spirv_get_or_create_const_bool(b, type_id, v)
	}
	return 0
}

@(private = "file")
spirv_emit_call :: proc(b: ^SPIRV_Builder, call: ^IR_Call, expr: ^IR_Expr) -> u32 {
	result_type_id := spirv_get_or_create_type(b, expr.type)

	// Handle texture sampling
	if call.is_builtin && call.name == "sample" && len(call.args) >= 2 {
		// Load the sampled image (combined texture+sampler in SPIR-V)
		sampled_image_id := spirv_emit_expr(b, call.args[0])
		coord_id := spirv_emit_expr(b, call.args[1])
		result_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_ImageSampleImplicitLod, result_type_id, result_id, sampled_image_id, coord_id)
		return result_id
	}

	// Handle sample_level — OpImageSampleExplicitLod with Lod operand
	if call.is_builtin && call.name == "sample_level" && len(call.args) >= 3 {
		sampled_image_id := spirv_emit_expr(b, call.args[0])
		coord_id := spirv_emit_expr(b, call.args[1])
		lod_id := spirv_emit_expr(b, call.args[2])
		result_id := spirv_alloc_id(b)
		// ImageOperandsMask: Lod = 0x2
		spirv_encode_inst(&b.func_section, SpvOp_ImageSampleExplicitLod, result_type_id, result_id, sampled_image_id, coord_id, 0x2, lod_id)
		return result_id
	}

	// Handle shadow texture sampling (depth comparison)
	if call.is_builtin && call.name == "sample_shadow" && len(call.args) >= 3 {
		sampled_image_id := spirv_emit_expr(b, call.args[0])
		coord_id := spirv_emit_expr(b, call.args[1])
		dref_id := spirv_emit_expr(b, call.args[2])
		result_id := spirv_alloc_id(b)
		spirv_encode_inst(&b.func_section, SpvOp_ImageSampleDrefImplicitLod, result_type_id, result_id, sampled_image_id, coord_id, dref_id)
		return result_id
	}

	// GLSL.std.450 extended instructions
	if call.is_builtin {
		if glsl_inst, ok := spirv_glsl_ext_inst(call.name); ok {
			// Select integer variants for integer types
			if len(call.args) > 0 && call.args[0].type != nil {
				is_uint := spirv_is_uint_type(call.args[0].type)
				is_int := spirv_is_int_type(call.args[0].type) && !is_uint
				switch call.name {
				case "min":   if is_uint { glsl_inst = GLSLstd450_UMin } else if is_int { glsl_inst = GLSLstd450_SMin }
				case "max":   if is_uint { glsl_inst = GLSLstd450_UMax } else if is_int { glsl_inst = GLSLstd450_SMax }
				case "clamp": if is_uint { glsl_inst = GLSLstd450_UClamp } else if is_int { glsl_inst = GLSLstd450_SClamp }
				case "abs":   if is_int { glsl_inst = GLSLstd450_SAbs }
				case "sign":  if is_int { glsl_inst = GLSLstd450_SSign }
				}
			}
			arg_ids := make([dynamic]u32)
			for arg in call.args {
				append(&arg_ids, spirv_emit_expr(b, arg))
			}
			// GLSL.std.450 requires all operands to match result type.
			// Splat scalar args to vectors when result is a vector.
			if vec, is_vec := expr.type^.(Type_Vector); is_vec {
				for arg, i in call.args {
					if arg.type != nil {
						if _, is_scalar := arg.type^.(Type_Scalar); is_scalar {
							// Splat: construct vector from repeated scalar
							splat_args := make([dynamic]u32)
							append(&splat_args, result_type_id, spirv_alloc_id(b))
							for _ in 0 ..< vec.size {
								append(&splat_args, arg_ids[i])
							}
							splat_id := splat_args[1]
							spirv_encode_inst(&b.func_section, SpvOp_CompositeConstruct, ..splat_args[:])
							arg_ids[i] = splat_id
						}
					}
				}
			}
			result_id := spirv_alloc_id(b)
			// OpExtInst: result_type, result_id, ext_set, instruction, ...args
			full_args := make([dynamic]u32)
			append(&full_args, result_type_id, result_id, b.glsl_ext_id, glsl_inst)
			for aid in arg_ids {
				append(&full_args, aid)
			}
			spirv_encode_inst(&b.func_section, SpvOp_ExtInst, ..full_args[:])
			return result_id
		}

		// Handle dot product specially — it's a dedicated opcode, not GLSL.std.450
		if call.name == "dot" && len(call.args) >= 2 {
			left_id := spirv_emit_expr(b, call.args[0])
			right_id := spirv_emit_expr(b, call.args[1])
			result_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Dot, result_type_id, result_id, left_id, right_id)
			return result_id
		}

		// Handle transpose — native SPIR-V opcode, not GLSL.std.450
		if call.name == "transpose" && len(call.args) >= 1 {
			mat_id := spirv_emit_expr(b, call.args[0])
			result_id := spirv_alloc_id(b)
			spirv_encode_inst(&b.func_section, SpvOp_Transpose, result_type_id, result_id, mat_id)
			return result_id
		}

		// Handle derivative ops — native SPIR-V opcodes, not GLSL.std.450
		if (call.name == "dfdx" || call.name == "dfdy" || call.name == "fwidth") && len(call.args) >= 1 {
			arg_id := spirv_emit_expr(b, call.args[0])
			result_id := spirv_alloc_id(b)
			opcode: u32
			switch call.name {
			case "dfdx":   opcode = SpvOp_DPdx
			case "dfdy":   opcode = SpvOp_DPdy
			case "fwidth": opcode = SpvOp_Fwidth
			}
			spirv_encode_inst(&b.func_section, opcode, result_type_id, result_id, arg_id)
			return result_id
		}
	}

	// Regular function call
	arg_ids := make([dynamic]u32)
	for arg in call.args {
		append(&arg_ids, spirv_emit_expr(b, arg))
	}
	result_id := spirv_alloc_id(b)
	full_args := make([dynamic]u32)
	append(&full_args, result_type_id, result_id)
	// Look up pre-allocated function ID
	fn_id := b.function_ids[call.name] or_else 0
	append(&full_args, fn_id)
	for aid in arg_ids {
		append(&full_args, aid)
	}
	spirv_encode_inst(&b.func_section, SpvOp_FunctionCall, ..full_args[:])
	return result_id
}

// -- Helpers --

@(private = "file")
spirv_binary_opcode :: proc(op: IR_Op, result_type: ^Resolved_Type, left_type: ^Resolved_Type, right_type: ^Resolved_Type) -> u32 {
	is_float := spirv_is_float_type(left_type)
	is_uint := spirv_is_uint_type(left_type)

	switch op {
	case .Add: return is_float ? u32(SpvOp_FAdd) : u32(SpvOp_IAdd)
	case .Sub: return is_float ? u32(SpvOp_FSub) : u32(SpvOp_ISub)
	case .Mul:
		if left_type != nil && right_type != nil {
			_, l_is_mat := left_type^.(Type_Matrix)
			_, r_is_mat := right_type^.(Type_Matrix)
			_, l_is_vec := left_type^.(Type_Vector)
			_, r_is_vec := right_type^.(Type_Vector)
			_, r_is_scalar := right_type^.(Type_Scalar)
			_, l_is_scalar := left_type^.(Type_Scalar)

			if l_is_mat && r_is_vec { return SpvOp_MatrixTimesVector }
			if l_is_vec && r_is_mat { return SpvOp_VectorTimesMatrix }
			if l_is_mat && r_is_mat { return SpvOp_MatrixTimesMatrix }
			if l_is_mat && r_is_scalar { return SpvOp_MatrixTimesScalar }
			if l_is_scalar && r_is_mat { return SpvOp_MatrixTimesScalar }
			if l_is_vec && r_is_scalar { return SpvOp_VectorTimesScalar }
			if l_is_scalar && r_is_vec { return SpvOp_VectorTimesScalar }
		}
		return is_float ? u32(SpvOp_FMul) : u32(SpvOp_IMul)
	case .Div: return is_float ? u32(SpvOp_FDiv) : (is_uint ? u32(SpvOp_UDiv) : u32(SpvOp_SDiv))
	case .Mod: return is_float ? u32(SpvOp_FMod) : (is_uint ? u32(SpvOp_UMod) : u32(SpvOp_SMod))
	case .Eq:  return is_float ? u32(SpvOp_FOrdEqual) : u32(SpvOp_IEqual)
	case .Neq: return is_float ? u32(SpvOp_FOrdNotEqual) : u32(SpvOp_INotEqual)
	case .Lt:  return is_float ? u32(SpvOp_FOrdLessThan) : (is_uint ? u32(SpvOp_ULessThan) : u32(SpvOp_SLessThan))
	case .Gt:  return is_float ? u32(SpvOp_FOrdGreaterThan) : (is_uint ? u32(SpvOp_UGreaterThan) : u32(SpvOp_SGreaterThan))
	case .Lte: return is_float ? u32(SpvOp_FOrdLessThanEqual) : (is_uint ? u32(SpvOp_ULessThanEqual) : u32(SpvOp_SLessThanEqual))
	case .Gte: return is_float ? u32(SpvOp_FOrdGreaterThanEqual) : (is_uint ? u32(SpvOp_UGreaterThanEqual) : u32(SpvOp_SGreaterThanEqual))
	case .And: return SpvOp_LogicalAnd
	case .Or:  return SpvOp_LogicalOr
	case .Neg: return is_float ? u32(SpvOp_FNegate) : u32(SpvOp_SNegate)
	case .Not: return SpvOp_LogicalNot
	}
	return SpvOp_Nop
}

@(private = "file")
spirv_is_float_type :: proc(t: ^Resolved_Type) -> bool {
	if t == nil do return false
	#partial switch v in t^ {
	case Type_Scalar: return v.kind == .Float || v.kind == .Half
	case Type_Vector: return v.elem == .Float || v.elem == .Half
	case Type_Matrix: return true // matrices are always float
	case: return false
	}
	return false
}

@(private = "file")
spirv_splat_scalar :: proc(b: ^SPIRV_Builder, scalar_id: u32, vec_type_id: u32, size: int) -> u32 {
	splat_args := make([dynamic]u32)
	splat_id := spirv_alloc_id(b)
	append(&splat_args, vec_type_id, splat_id)
	for _ in 0 ..< size {
		append(&splat_args, scalar_id)
	}
	spirv_encode_inst(&b.func_section, SpvOp_CompositeConstruct, ..splat_args[:])
	return splat_id
}

@(private = "file")
spirv_is_int_type :: proc(t: ^Resolved_Type) -> bool {
	if t == nil do return false
	#partial switch v in t^ {
	case Type_Scalar: return v.kind == .Int || v.kind == .Uint
	case Type_Vector: return v.elem == .Int || v.elem == .Uint
	case: return false
	}
	return false
}

@(private = "file")
spirv_is_uint_type :: proc(t: ^Resolved_Type) -> bool {
	if t == nil do return false
	#partial switch v in t^ {
	case Type_Scalar: return v.kind == .Uint
	case Type_Vector: return v.elem == .Uint
	case: return false
	}
	return false
}

@(private = "file")
spirv_cast_opcode :: proc(from: ^Resolved_Type, to: ^Resolved_Type) -> u32 {
	from_float := spirv_is_float_type(from)
	to_float := spirv_is_float_type(to)

	if from_float && !to_float {
		// Check if target is signed (scalar or vector of signed int)
		if to != nil {
			#partial switch v in to^ {
			case Type_Scalar:
				if v.kind == .Int do return SpvOp_ConvertFToS
			case Type_Vector:
				if v.elem == .Int do return SpvOp_ConvertFToS
			}
		}
		return SpvOp_ConvertFToU
	}
	if !from_float && to_float {
		if spirv_is_uint_type(from) {
			return SpvOp_ConvertUToF
		}
		return SpvOp_ConvertSToF
	}
	if from_float && to_float {
		return SpvOp_FConvert
	}
	return SpvOp_Bitcast
}

@(private = "file")
spirv_swizzle_index :: proc(ch: u8) -> int {
	switch ch {
	case 'x', 'r', 's': return 0
	case 'y', 'g', 't': return 1
	case 'z', 'b', 'p': return 2
	case 'w', 'a', 'q': return 3
	}
	return 0
}

@(private = "file")
spirv_execution_model :: proc(stage: Shader_Stage) -> u32 {
	#partial switch stage {
	case .Vertex:   return SpvExecutionModel_Vertex
	case .Fragment: return SpvExecutionModel_Fragment
	case .Compute:  return SpvExecutionModel_GLCompute
	}
	return SpvExecutionModel_Vertex
}

@(private = "file")
spirv_builtin_id :: proc(name: string, stage: Shader_Stage = .Vertex) -> u32 {
	switch name {
	case "position":
		return stage == .Fragment ? SpvBuiltIn_FragCoord : SpvBuiltIn_Position
	case "vertex_id":             return SpvBuiltIn_VertexIndex
	case "instance_id":           return SpvBuiltIn_InstanceIndex
	case "frag_coord":            return SpvBuiltIn_FragCoord
	case "front_facing":          return SpvBuiltIn_FrontFacing
	case "frag_depth":            return SpvBuiltIn_FragDepth
	case "local_invocation_id":   return SpvBuiltIn_LocalInvocationId
	case "local_invocation_index": return SpvBuiltIn_LocalInvocationIndex
	case "global_invocation_id":  return SpvBuiltIn_GlobalInvocationId
	case "workgroup_id":          return SpvBuiltIn_WorkgroupId
	}
	return 0
}

@(private = "file")
spirv_glsl_ext_inst :: proc(name: string) -> (u32, bool) {
	switch name {
	case "round":       return GLSLstd450_Round, true
	case "floor":       return GLSLstd450_Floor, true
	case "ceil":        return GLSLstd450_Ceil, true
	case "fract":       return GLSLstd450_Fract, true
	case "abs":         return GLSLstd450_FAbs, true
	case "sign":        return GLSLstd450_FSign, true
	case "sin":         return GLSLstd450_Sin, true
	case "cos":         return GLSLstd450_Cos, true
	case "tan":         return GLSLstd450_Tan, true
	case "asin":        return GLSLstd450_Asin, true
	case "acos":        return GLSLstd450_Acos, true
	case "atan":        return GLSLstd450_Atan, true
	case "atan2":       return GLSLstd450_Atan2, true
	case "pow":         return GLSLstd450_Pow, true
	case "exp":         return GLSLstd450_Exp, true
	case "log":         return GLSLstd450_Log, true
	case "exp2":        return GLSLstd450_Exp2, true
	case "log2":        return GLSLstd450_Log2, true
	case "sqrt":        return GLSLstd450_Sqrt, true
	case "inversesqrt": return GLSLstd450_InverseSqrt, true
	case "inverse":     return GLSLstd450_MatrixInverse, true
	case "determinant": return GLSLstd450_Determinant, true
	case "min":         return GLSLstd450_FMin, true
	case "max":         return GLSLstd450_FMax, true
	case "clamp":       return GLSLstd450_FClamp, true
	case "mix":         return GLSLstd450_FMix, true
	case "step":        return GLSLstd450_Step, true
	case "smoothstep":  return GLSLstd450_SmoothStep, true
	case "length":      return GLSLstd450_Length, true
	case "distance":    return GLSLstd450_Distance, true
	case "cross":       return GLSLstd450_Cross, true
	case "normalize":   return GLSLstd450_Normalize, true
	case "reflect":     return GLSLstd450_Reflect, true
	case "refract":     return GLSLstd450_Refract, true
	}
	return 0, false
}

// -- Final assembly --

@(private = "file")
spirv_splice_vars :: proc(b: ^SPIRV_Builder, insert_pos: int) {
	// Insert var_buffer contents at insert_pos in func_section
	// (right after OpLabel, before any other instructions)
	old_len := len(b.func_section)
	var_len := len(b.var_buffer)

	// Extend func_section by var_len
	resize(&b.func_section, old_len + var_len)

	// Shift existing instructions after insert_pos to make room
	copy(b.func_section[insert_pos + var_len:], b.func_section[insert_pos:old_len])

	// Copy var_buffer into the gap
	copy(b.func_section[insert_pos:], b.var_buffer[:])
}

// Check if a statement list ends with a block-terminating statement (discard/return).
// These produce SPIR-V terminators (OpKill, OpReturn), so we must not emit OpBranch after them.
@(private = "file")
spirv_block_has_terminator :: proc(stmts: []IR_Stmt) -> bool {
	if len(stmts) == 0 do return false
	last := stmts[len(stmts) - 1]
	#partial switch s in last {
	case ^IR_Discard:  return true
	case ^IR_Return:   return true
	case ^IR_Break:    return true
	case ^IR_Continue: return true
	}
	return false
}

@(private = "file")
ir_stmt_span :: proc(stmt: IR_Stmt) -> Source_Span {
	switch s in stmt {
	case ^IR_Let:          return s.span
	case ^IR_Assign:       return s.span
	case ^IR_Return:       return s.span
	case ^IR_If:           return s.span
	case ^IR_For:          return s.span
	case ^IR_While:        return s.span
	case ^IR_Store_Output: return s.span
	case ^IR_Expr_Stmt:    return s.span
	case ^IR_Barrier:      return s.span
	case ^IR_Discard:      return s.span
	case ^IR_Break:        return s.span
	case ^IR_Continue:     return s.span
	}
	return {}
}

@(private = "file")
spirv_assemble :: proc(b: ^SPIRV_Builder) -> []u8 {
	words := make([dynamic]u32)

	// Header
	append(&words, SPIRV_MAGIC)
	append(&words, SPIRV_VERSION)
	append(&words, SPIRV_GENERATOR)
	append(&words, b.next_id) // bound
	append(&words, 0)         // schema

	// Sections in order
	for w in b.capabilities  { append(&words, w) }
	for w in b.extensions    { append(&words, w) }
	for w in b.ext_imports   { append(&words, w) }
	for w in b.mem_model     { append(&words, w) }
	for w in b.entry_points  { append(&words, w) }
	for w in b.exec_modes    { append(&words, w) }
	for w in b.debug_source  { append(&words, w) } // OpString, OpSource
	for w in b.debug_names   { append(&words, w) } // OpName, OpMemberName
	for w in b.debug_process { append(&words, w) } // OpModuleProcessed
	for w in b.annotations   { append(&words, w) }
	for w in b.type_section  { append(&words, w) }
	for w in b.func_section  { append(&words, w) }

	// Convert to bytes
	byte_len := len(words) * 4
	bytes := make([]u8, byte_len)
	for w, i in words {
		bytes[i*4 + 0] = u8(w)
		bytes[i*4 + 1] = u8(w >> 8)
		bytes[i*4 + 2] = u8(w >> 16)
		bytes[i*4 + 3] = u8(w >> 24)
	}
	return bytes
}