sharkk/gpu public

branch main
showing the latest snapshot on main
ir_builder.odin 24.3 KB · Plain text
package shader

import "core:fmt"

// IR Builder — lowers typed AST to structured IR

IR_Builder :: struct {
	module:        IR_Module,
	sema:          ^Sema,
	ast_mod:       ^Ast_Module,
	current_fn:    ^IR_Function,
	current_entry: ^Ast_Function,
	body_stack:    [dynamic]^[dynamic]IR_Stmt,
	diagnostics:   [dynamic]Diagnostic,
	var_map:       map[string]IR_Var_Id,
	next_var_id:   int,
}

ir_build_module :: proc(ast_mod: ^Ast_Module, sema: ^Sema, allocator := context.allocator) -> (IR_Module, []Diagnostic) {
	b := IR_Builder{
		module = IR_Module{
			functions      = make([dynamic]IR_Function, allocator),
			structs        = make([dynamic]IR_Struct, allocator),
			bindings       = make([dynamic]IR_Binding, allocator),
			spec_constants = make([dynamic]IR_Spec_Constant, allocator),
			shared_vars    = make([dynamic]IR_Shared_Var, allocator),
		},
		sema        = sema,
		ast_mod     = ast_mod,
		body_stack  = make([dynamic]^[dynamic]IR_Stmt, allocator),
		diagnostics = make([dynamic]Diagnostic, allocator),
		var_map     = make(map[string]IR_Var_Id, allocator = allocator),
	}

	build_structs(&b)
	build_bindings(&b)
	build_constants(&b)
	build_shared_vars(&b)
	build_functions(&b)

	return b.module, b.diagnostics[:]
}

// Append a statement to the current body target
@(private = "file")
emit_stmt :: proc(b: ^IR_Builder, stmt: IR_Stmt) {
	target := b.body_stack[len(b.body_stack) - 1]
	append(target, stmt)
}

@(private = "file")
push_body :: proc(b: ^IR_Builder, body: ^[dynamic]IR_Stmt) {
	append(&b.body_stack, body)
}

@(private = "file")
pop_body :: proc(b: ^IR_Builder) {
	pop(&b.body_stack)
}

@(private = "file")
alloc_var :: proc(b: ^IR_Builder, name: string, type: ^Resolved_Type) -> IR_Var_Id {
	id := IR_Var_Id(b.next_var_id)
	b.next_var_id += 1
	append(&b.current_fn.var_decls, IR_Var_Decl{id = id, name = name, type = type})
	b.var_map[name] = id
	return id
}

// -- Structs --

@(private = "file")
build_structs :: proc(b: ^IR_Builder) {
	for name, resolved in b.sema.structs {
		fields := make([dynamic]IR_Struct_Field)
		for f in resolved.fields {
			append(&fields, IR_Struct_Field{
				name       = f.name,
				type       = f.type,
				attributes = f.attributes,
			})
		}
		append(&b.module.structs, IR_Struct{
			name   = name,
			fields = fields[:],
		})
	}
}

// -- Bindings --

@(private = "file")
build_bindings :: proc(b: ^IR_Builder) {
	next_binding: map[int]int // group -> next auto-binding number

	for ast_b in b.ast_mod.bindings {
		// Check for @push_constant attribute first
		if has_attribute(ast_b.attributes, "push_constant") {
			sym := scope_lookup(b.sema.global_scope, ast_b.name)
			rt: ^Resolved_Type
			if sym != nil do rt = sym.type

			type_name := type_expr_name(ast_b.type_expr)
			struct_ref: ^Type_Struct_Resolved
			if resolved, ok := b.sema.structs[type_name]; ok {
				struct_ref = resolved
			}

			append(&b.module.bindings, IR_Binding{
				name          = ast_b.name,
				type          = rt,
				kind          = .Push_Constant,
				group         = -1,
				binding_num   = -1,
				struct_ref    = struct_ref,
				address_space = .Push_Constant,
			})
			continue
		}

		explicit_group, explicit_binding := get_group_binding(ast_b.attributes)

		// Missing @group defaults to 0
		group := explicit_group >= 0 ? explicit_group : 0

		// Missing @binding auto-increments within the group
		binding_num: int
		if explicit_binding >= 0 {
			binding_num = explicit_binding
			next_binding[group] = explicit_binding + 1
		} else {
			binding_num = next_binding[group] // 0 if not yet set
			next_binding[group] = binding_num + 1
		}

		sym := scope_lookup(b.sema.global_scope, ast_b.name)
		rt: ^Resolved_Type
		if sym != nil do rt = sym.type

		type_name := type_expr_name(ast_b.type_expr)

		if is_sampler_type_name(type_name) {
			// Split combined sampler into texture + sampler bindings
			append(&b.module.bindings, IR_Binding{
				name          = fmt.aprintf("%s_tex", ast_b.name),
				type          = rt,
				kind          = .Texture,
				group         = group,
				binding_num   = binding_num,
				combined_name = ast_b.name,
				address_space = .Uniform,
			})
			append(&b.module.bindings, IR_Binding{
				name          = fmt.aprintf("%s_samp", ast_b.name),
				type          = rt,
				kind          = .Sampler,
				group         = group,
				binding_num   = binding_num + 1,
				combined_name = ast_b.name,
				address_space = .Uniform,
			})
			// Consume two binding slots
			next_binding[group] = binding_num + 2
		} else {
			struct_ref: ^Type_Struct_Resolved
			if resolved, ok := b.sema.structs[type_name]; ok {
				struct_ref = resolved
			}

			ir_kind: IR_Binding_Kind
			addr_space: IR_Address_Space
			switch ast_b.kind {
			case .Uniform: ir_kind = .Uniform; addr_space = .Uniform
			case .Buffer:  ir_kind = .Buffer;  addr_space = .Storage
			}

			append(&b.module.bindings, IR_Binding{
				name          = ast_b.name,
				type          = rt,
				kind          = ir_kind,
				group         = group,
				binding_num   = binding_num,
				struct_ref    = struct_ref,
				address_space = addr_space,
			})
		}
	}
}

// -- Constants --

@(private = "file")
build_constants :: proc(b: ^IR_Builder) {
	for c in b.ast_mod.constants {
		// Check for @spec attribute
		spec_id := -1
		for attr in c.attributes {
			if attr.name == "spec" && len(attr.args) > 0 {
				spec_id = parse_int(attr.args[0])
			}
		}
		if spec_id < 0 do continue // regular const — folded by sema, skip IR

		// Resolve type from sema symbol table
		rt: ^Resolved_Type
		if sym := scope_lookup(b.sema.global_scope, c.name); sym != nil {
			rt = sym.type
		}

		// Extract default value
		default_val: IR_Const_Value
		if c.value != nil {
			#partial switch d in c.value.derived {
			case ^Ast_Literal:
				switch v in d.value {
				case i64:    default_val = v
				case f64:    default_val = v
				case bool:   default_val = v
				case string: // skip
				}
			case ^Ast_Unary:
				// Handle negative literals like -1
				if d.op == .Neg {
					if lit, ok := d.operand.derived.(^Ast_Literal); ok {
						switch v in lit.value {
						case i64:    default_val = -v
						case f64:    default_val = -v
						case bool:   // skip
						case string: // skip
						}
					}
				}
			}
		}

		append(&b.module.spec_constants, IR_Spec_Constant{
			name          = c.name,
			spec_id       = spec_id,
			type          = rt,
			default_value = default_val,
		})
	}
}

// -- Shared Variables --

@(private = "file")
build_shared_vars :: proc(b: ^IR_Builder) {
	for sv in b.ast_mod.shared_vars {
		sym := scope_lookup(b.sema.global_scope, sv.name)
		rt: ^Resolved_Type
		if sym != nil do rt = sym.type
		append(&b.module.shared_vars, IR_Shared_Var{
			name = sv.name,
			type = rt,
		})
	}
}

// -- Functions --

@(private = "file")
build_functions :: proc(b: ^IR_Builder) {
	for ast_fn in b.ast_mod.functions {
		ir_fn := IR_Function{
			name      = ast_fn.name,
			body      = make([dynamic]IR_Stmt),
			inputs    = make([dynamic]IR_IO_Var),
			outputs   = make([dynamic]IR_IO_Var),
			var_decls = make([dynamic]IR_Var_Decl),
		}

		// Params — resolve types from type expression
		params := make([dynamic]IR_Param)
		for p in ast_fn.params {
			type_name := type_expr_name(p.type)
			pt: ^Resolved_Type
			sym := scope_lookup(b.sema.global_scope, type_name)
			if sym != nil do pt = sym.type
			append(&params, IR_Param{name = p.name, type = pt})
		}
		ir_fn.params = params[:]

		// Return type
		if ast_fn.return_type != nil {
			name := type_expr_name(ast_fn.return_type^)
			sym := scope_lookup(b.sema.global_scope, name)
			if sym != nil {
				ir_fn.return_type = sym.type
			}
		} else {
			ir_fn.return_type = TYPE_VOID
		}

		// Entry point detection
		for attr in ast_fn.attributes {
			if attr.name == "entry" && len(attr.args) > 0 {
				ir_fn.is_entry = true
				switch attr.args[0] {
				case "vertex":   ir_fn.stage = .Vertex
				case "fragment": ir_fn.stage = .Fragment
				case "compute":  ir_fn.stage = .Compute
				}
			}
			if attr.name == "workgroup_size" {
				for i := 0; i < min(len(attr.args), 3); i += 1 {
					ir_fn.workgroup_size[i] = parse_int(attr.args[i])
				}
			}
		}

		// Flatten entry point I/O
		if ir_fn.is_entry {
			build_entry_io(b, ast_fn, &ir_fn)
		}

		b.current_fn = &ir_fn
		b.current_entry = ast_fn
		b.next_var_id = 1 // start at 1; 0 is invalid sentinel
		clear(&b.var_map)

		// Register parameters as var IDs so build_ident can resolve them
		for &p in ir_fn.params {
			p.id = alloc_var(b, p.name, p.type)
		}

		push_body(b, &ir_fn.body)

		for stmt in ast_fn.body {
			build_stmt(b, stmt)
		}

		pop_body(b)
		b.current_fn = nil
		b.current_entry = nil

		append(&b.module.functions, ir_fn)
	}
}

@(private = "file")
build_entry_io :: proc(b: ^IR_Builder, ast_fn: ^Ast_Function, ir_fn: ^IR_Function) {
	if len(ast_fn.params) > 0 {
		param_type_name := type_expr_name(ast_fn.params[0].type)
		if resolved, ok := b.sema.structs[param_type_name]; ok {
			// Check if the struct has @varying — fragment inputs skip @builtin fields
			is_varying := false
			for ast_s in b.ast_mod.structs {
				if ast_s.name == param_type_name {
					is_varying = has_attribute(ast_s.attributes, "varying")
					break
				}
			}
			strip_builtins := is_varying && ir_fn.stage == .Fragment

			loc_counter := 0
			for f in resolved.fields {
				builtin := get_builtin_name(f.attributes)

				// @varying struct: strip @builtin fields from fragment inputs
				if strip_builtins && builtin != "" do continue

				explicit_loc := get_location(f.attributes)
				loc: int
				if builtin != "" {
					loc = -1
				} else if explicit_loc >= 0 {
					loc = explicit_loc
					loc_counter = explicit_loc + 1
				} else {
					loc = loc_counter
					loc_counter += 1
				}
				append(&ir_fn.inputs, IR_IO_Var{
					name     = f.name,
					type     = f.type,
					location = loc,
					builtin  = builtin,
				})
			}
		}
	}

	if ast_fn.return_type != nil {
		ret_type_name := type_expr_name(ast_fn.return_type^)
		if resolved, ok := b.sema.structs[ret_type_name]; ok {
			loc_counter := 0
			for f in resolved.fields {
				builtin := get_builtin_name(f.attributes)
				explicit_loc := get_location(f.attributes)
				loc: int
				if builtin != "" {
					loc = -1
				} else if explicit_loc >= 0 {
					loc = explicit_loc
					loc_counter = explicit_loc + 1
				} else {
					loc = loc_counter
					loc_counter += 1
				}
				append(&ir_fn.outputs, IR_IO_Var{
					name     = f.name,
					type     = f.type,
					location = loc,
					builtin  = builtin,
				})
			}
		}
	}
}

// -- Statements --

@(private = "file")
build_stmt :: proc(b: ^IR_Builder, node: ^Ast_Node) {
	if node == nil do return

	#partial switch d in node.derived {
	case ^Ast_Let:
		build_let(b, d, node)
	case ^Ast_Assign:
		build_assign(b, d)
	case ^Ast_Output_Assign:
		build_output_assign(b, d)
	case ^Ast_Return:
		build_return(b, d)
	case ^Ast_If:
		build_if(b, d)
	case ^Ast_For:
		build_for(b, d)
	case ^Ast_While:
		build_while(b, d)
	case ^Ast_Discard:
		s := new(IR_Discard)
		s.span = d.span
		emit_stmt(b, s)
	case ^Ast_Break:
		s := new(IR_Break)
		s.span = d.span
		emit_stmt(b, s)
	case ^Ast_Continue:
		s := new(IR_Continue)
		s.span = d.span
		emit_stmt(b, s)
	case:
		// Check for barrier() call → emit IR_Barrier statement
		if call, ok := node.derived.(^Ast_Call); ok {
			if ident, ok2 := call.callee.derived.(^Ast_Ident); ok2 && ident.name == "barrier" {
				barrier := new(IR_Barrier)
				barrier.kind = .Workgroup
				barrier.span = node.span
				emit_stmt(b, barrier)
				return
			}
		}
		expr := build_expr(b, node)
		if expr != nil {
			s := new(IR_Expr_Stmt)
			s.expr = expr
			s.span = node.span
			emit_stmt(b, s)
		}
	}
}

@(private = "file")
build_output_assign :: proc(b: ^IR_Builder, assign: ^Ast_Output_Assign) {
	if assign.io_index < 0 || assign.io_index >= len(b.current_fn.outputs) do return
	s := new(IR_Store_Output)
	s.io_index = assign.io_index
	s.value = build_expr(b, assign.value)
	s.span = assign.span
	emit_stmt(b, s)
}

@(private = "file")
build_let :: proc(b: ^IR_Builder, let: ^Ast_Let, node: ^Ast_Node) {
	s := new(IR_Let)
	s.name = let.name
	s.value = build_expr(b, let.value)
	s.type = node.resolved_type
	s.id = alloc_var(b, let.name, node.resolved_type)
	s.span = let.span
	s.mutable = let.mutable
	emit_stmt(b, s)
}

@(private = "file")
build_assign :: proc(b: ^IR_Builder, assign: ^Ast_Assign) {
	s := new(IR_Assign)
	s.target = build_expr(b, assign.target)
	s.value = build_expr(b, assign.value)
	s.span = assign.span
	emit_stmt(b, s)
}

@(private = "file")
build_return :: proc(b: ^IR_Builder, ret: ^Ast_Return) {
	if ret.value == nil {
		r := new(IR_Return)
		r.span = ret.span
		emit_stmt(b, r)
		return
	}

	// Entry point: decompose return into store_output instructions
	if b.current_fn.is_entry && len(b.current_fn.outputs) > 0 {
		build_entry_return(b, ret)
		return
	}

	s := new(IR_Return)
	s.value = build_expr(b, ret.value)
	s.span = ret.span
	emit_stmt(b, s)
}

@(private = "file")
build_entry_return :: proc(b: ^IR_Builder, ret: ^Ast_Return) {
	fn := b.current_fn
	value := ret.value

	if value.kind == .Struct_Literal {
		sl := value.derived.(^Ast_Struct_Literal)
		for f in sl.fields {
			for io, idx in fn.outputs {
				if io.name == f.name {
					s := new(IR_Store_Output)
					s.io_index = idx
					s.value = build_expr(b, f.value)
					s.span = ret.span
					emit_stmt(b, s)
					break
				}
			}
		}
		return
	}

	result_expr := build_expr(b, value)

	// Emit a synthetic let to avoid sharing result_expr across multiple field accesses.
	temp_name := "__result"
	temp_let := new(IR_Let)
	temp_let.name = temp_name
	temp_let.value = result_expr
	temp_let.type = result_expr.type
	temp_let.id = alloc_var(b, temp_name, result_expr.type)
	temp_let.span = ret.span
	emit_stmt(b, temp_let)

	for io, idx in fn.outputs {
		s := new(IR_Store_Output)
		s.io_index = idx
		s.span = ret.span
		vr := new(IR_Var_Ref)
		vr.name = temp_name
		vr.id = temp_let.id
		fa := new(IR_Field_Access)
		fa.object = make_ir_expr(.Var_Ref, result_expr.type, vr)
		fa.field_name = io.name
		s.value = make_ir_expr(.Field_Access, io.type, fa)
		emit_stmt(b, s)
	}
}

@(private = "file")
build_if :: proc(b: ^IR_Builder, if_node: ^Ast_If) {
	s := new(IR_If)
	s.condition = build_expr(b, if_node.condition)
	s.then_body = make([dynamic]IR_Stmt)
	s.else_body = make([dynamic]IR_Stmt)
	s.span = if_node.span

	push_body(b, &s.then_body)
	for stmt in if_node.then_body {
		build_stmt(b, stmt)
	}
	pop_body(b)

	elseifs := make([dynamic]IR_Elseif)
	for ei in if_node.elseif_clauses {
		ir_ei := IR_Elseif{
			body = make([dynamic]IR_Stmt),
		}
		ir_ei.condition = build_expr(b, ei.condition)
		push_body(b, &ir_ei.body)
		for stmt in ei.body {
			build_stmt(b, stmt)
		}
		pop_body(b)
		append(&elseifs, ir_ei)
	}
	s.elseif_clauses = elseifs[:]

	if len(if_node.else_body) > 0 {
		push_body(b, &s.else_body)
		for stmt in if_node.else_body {
			build_stmt(b, stmt)
		}
		pop_body(b)
	}

	emit_stmt(b, s)
}

@(private = "file")
build_for :: proc(b: ^IR_Builder, for_node: ^Ast_For) {
	s := new(IR_For)
	s.var_name = for_node.var_name
	s.var_id = alloc_var(b, for_node.var_name, nil) // type resolved later
	s.start = build_expr(b, for_node.start)
	s.stop = build_expr(b, for_node.stop)
	if for_node.step != nil {
		s.step = build_expr(b, for_node.step)
	}
	s.body = make([dynamic]IR_Stmt)
	s.span = for_node.span

	push_body(b, &s.body)
	for stmt in for_node.body {
		build_stmt(b, stmt)
	}
	pop_body(b)

	emit_stmt(b, s)
}

@(private = "file")
build_while :: proc(b: ^IR_Builder, while_node: ^Ast_While) {
	s := new(IR_While)
	s.condition = build_expr(b, while_node.condition)
	s.body = make([dynamic]IR_Stmt)
	s.span = while_node.span

	push_body(b, &s.body)
	for stmt in while_node.body {
		build_stmt(b, stmt)
	}
	pop_body(b)

	emit_stmt(b, s)
}

// -- Expressions --

@(private = "file")
build_expr :: proc(b: ^IR_Builder, node: ^Ast_Node) -> ^IR_Expr {
	if node == nil do return nil

	#partial switch d in node.derived {
	case ^Ast_Literal:
		return build_literal(d, node)
	case ^Ast_Ident:
		return build_ident(b, d, node)
	case ^Ast_Binary:
		return build_binary(b, d, node)
	case ^Ast_Unary:
		return build_unary(b, d, node)
	case ^Ast_Call:
		return build_call(b, d, node)
	case ^Ast_Field_Access:
		return build_field_access(b, d, node)
	case ^Ast_Index:
		return build_index(b, d, node)
	case ^Ast_Struct_Literal:
		return build_struct_literal(b, d, node)
	}

	return nil
}

@(private = "file")
build_literal :: proc(lit: ^Ast_Literal, node: ^Ast_Node) -> ^IR_Expr {
	ir_lit := new(IR_Literal)
	switch v in lit.value {
	case i64:    ir_lit.value = v
	case f64:    ir_lit.value = v
	case bool:   ir_lit.value = v
	case string: return nil
	}
	return make_ir_expr(.Literal, node.resolved_type, ir_lit)
}

@(private = "file")
build_ident :: proc(b: ^IR_Builder, ident: ^Ast_Ident, node: ^Ast_Node) -> ^IR_Expr {
	if _, ok := b.sema.bindings[ident.name]; ok {
		lb := new(IR_Load_Binding)
		lb.name = ident.name
		return make_ir_expr(.Load_Binding, node.resolved_type, lb)
	}

	// Check for shared variable reference
	if sym := scope_lookup(b.sema.global_scope, ident.name); sym != nil && sym.kind == .Shared {
		sr := new(IR_Shared_Ref)
		sr.name = ident.name
		return make_ir_expr(.Shared_Ref, node.resolved_type, sr)
	}

	// Inline constant references — const values are folded at IR build time
	// Skip spec constants — they must remain as variable references for specialization
	if sym := scope_lookup(b.sema.global_scope, ident.name); sym != nil && sym.kind == .Constant {
		is_spec := false
		for c in b.ast_mod.constants {
			if c.name == ident.name {
				for attr in c.attributes {
					if attr.name == "spec" {
						is_spec = true
						break
					}
				}
				break
			}
		}
		if !is_spec {
			for c in b.ast_mod.constants {
				if c.name == ident.name && c.value != nil {
					return build_expr(b, c.value)
				}
			}
		}
	}

	vr := new(IR_Var_Ref)
	vr.name = ident.name
	if id, ok := b.var_map[ident.name]; ok {
		vr.id = id
	}
	return make_ir_expr(.Var_Ref, node.resolved_type, vr)
}

@(private = "file")
build_binary :: proc(b: ^IR_Builder, bin: ^Ast_Binary, node: ^Ast_Node) -> ^IR_Expr {
	ir_bin := new(IR_Binary)
	ir_bin.op = ast_binop_to_ir(bin.op)
	ir_bin.left = build_expr(b, bin.left)
	ir_bin.right = build_expr(b, bin.right)
	return make_ir_expr(.Binary, node.resolved_type, ir_bin)
}

@(private = "file")
build_unary :: proc(b: ^IR_Builder, un: ^Ast_Unary, node: ^Ast_Node) -> ^IR_Expr {
	ir_un := new(IR_Unary)
	ir_un.op = un.op == .Neg ? IR_Op.Neg : IR_Op.Not
	ir_un.operand = build_expr(b, un.operand)
	return make_ir_expr(.Unary, node.resolved_type, ir_un)
}

@(private = "file")
build_call :: proc(b: ^IR_Builder, call: ^Ast_Call, node: ^Ast_Node) -> ^IR_Expr {
	name := ""
	if call.callee.kind == .Ident {
		ident := call.callee.derived.(^Ast_Ident)
		name = ident.name
	}

	sym := scope_lookup(b.sema.global_scope, name)

	// Type constructor or type cast
	if sym != nil && sym.kind == .Struct_Type {
		// Scalar type cast: float(x), int(x) where input is a different scalar type
		if len(call.args) == 1 {
			if _, is_scalar := node.resolved_type^.(Type_Scalar); is_scalar {
				arg_expr := build_expr(b, call.args[0])
				if arg_expr.type != nil {
					if _, arg_is_scalar := arg_expr.type^.(Type_Scalar); arg_is_scalar {
						if arg_expr.type != node.resolved_type {
							tc := new(IR_Type_Cast)
							tc.value = arg_expr
							return make_ir_expr(.Type_Cast, node.resolved_type, tc)
						}
					}
				}
			}
		}

		args := make([dynamic]^IR_Expr)
		for arg in call.args {
			append(&args, build_expr(b, arg))
		}

		// Scalar-to-vector splat: vec3(x) → vec3(x, x, x)
		if vec, is_vec := node.resolved_type^.(Type_Vector); is_vec && len(args) == 1 {
			if args[0] != nil && args[0].type != nil {
				if _, is_scalar := args[0].type^.(Type_Scalar); is_scalar {
					scalar := args[0]
					for _ in 1 ..< vec.size {
						append(&args, scalar)
					}
				}
			}
		}

		c := new(IR_Construct)
		c.type_name = name
		c.args = args[:]
		return make_ir_expr(.Construct, node.resolved_type, c)
	}

	// select(cond, true_val, false_val) → IR_Select
	if name == "select" && len(call.args) == 3 {
		sel := new(IR_Select)
		sel.condition = build_expr(b, call.args[0])
		sel.true_val = build_expr(b, call.args[1])
		sel.false_val = build_expr(b, call.args[2])
		return make_ir_expr(.Select, node.resolved_type, sel)
	}

	// Function/builtin call
	args := make([dynamic]^IR_Expr)
	for arg in call.args {
		append(&args, build_expr(b, arg))
	}

	ir_call := new(IR_Call)
	ir_call.name = name
	ir_call.args = args[:]
	ir_call.is_builtin = sym != nil && sym.kind == .Builtin_Function
	return make_ir_expr(.Call, node.resolved_type, ir_call)
}

@(private = "file")
build_field_access :: proc(b: ^IR_Builder, fa: ^Ast_Field_Access, node: ^Ast_Node) -> ^IR_Expr {
	// Entry point input field access (e.g. input.position)
	if b.current_fn != nil && b.current_fn.is_entry && fa.object.kind == .Ident {
		ident := fa.object.derived.(^Ast_Ident)
		if b.current_entry != nil && len(b.current_entry.params) > 0 && ident.name == b.current_entry.params[0].name {
			for io in b.current_fn.inputs {
				if io.name == fa.field {
					if io.builtin != "" {
						bv := new(IR_Builtin_Var)
						bv.name = io.builtin
						bv.stage = b.current_fn.stage
						bv.is_input = true
						return make_ir_expr(.Builtin_Var, node.resolved_type, bv)
					}
					inf := new(IR_Input_Field)
					inf.param_name = ident.name
					inf.field_name = fa.field
					return make_ir_expr(.Input_Field, node.resolved_type, inf)
				}
			}
		}
	}

	// Swizzle on vector types
	obj_type := fa.object.resolved_type
	if obj_type != nil {
		if v, ok := obj_type^.(Type_Vector); ok {
			if is_valid_swizzle(fa.field, v.size) {
				sw := new(IR_Swizzle)
				sw.object = build_expr(b, fa.object)
				sw.components = fa.field
				return make_ir_expr(.Swizzle, node.resolved_type, sw)
			}
		}
	}

	// Build the object expression first to determine value vs memory semantics
	obj_expr := build_expr(b, fa.object)

	// If the object is a binding load, use Field_Access (memory semantics)
	if _, is_binding := obj_expr.derived.(^IR_Load_Binding); is_binding {
		ir_fa := new(IR_Field_Access)
		ir_fa.object = obj_expr
		ir_fa.field_name = fa.field
		return make_ir_expr(.Field_Access, node.resolved_type, ir_fa)
	}

	// Otherwise use Composite_Extract (value semantics on SSA values)
	ce := new(IR_Composite_Extract)
	ce.object = obj_expr
	ce.field_name = fa.field
	ce.index = resolve_field_index(obj_type, fa.field)
	return make_ir_expr(.Composite_Extract, node.resolved_type, ce)
}

resolve_field_index :: proc(t: ^Resolved_Type, field_name: string) -> int {
	if t == nil do return 0
	#partial switch v in t^ {
	case Type_Struct_Resolved:
		for f, i in v.fields {
			if f.name == field_name do return i
		}
	case Type_Vector:
		// Map vector component names to indices
		if len(field_name) == 1 {
			ch := field_name[0]
			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
}

build_index :: proc(b: ^IR_Builder, idx: ^Ast_Index, node: ^Ast_Node) -> ^IR_Expr {
	ir_idx := new(IR_Index)
	ir_idx.object = build_expr(b, idx.object)
	ir_idx.index = build_expr(b, idx.index)
	return make_ir_expr(.Index, node.resolved_type, ir_idx)
}

@(private = "file")
build_struct_literal :: proc(b: ^IR_Builder, sl: ^Ast_Struct_Literal, node: ^Ast_Node) -> ^IR_Expr {
	resolved, ok := b.sema.structs[sl.type_name]

	if !ok {
		args := make([dynamic]^IR_Expr)
		for f in sl.fields {
			append(&args, build_expr(b, f.value))
		}
		c := new(IR_Construct)
		c.type_name = sl.type_name
		c.args = args[:]
		return make_ir_expr(.Construct, node.resolved_type, c)
	}

	// Emit args in declaration order
	args := make([dynamic]^IR_Expr)
	for rf in resolved.fields {
		found := false
		for f in sl.fields {
			if f.name == rf.name {
				append(&args, build_expr(b, f.value))
				found = true
				break
			}
		}
		if !found {
			append(&args, nil)
		}
	}

	c := new(IR_Construct)
	c.type_name = sl.type_name
	c.args = args[:]
	return make_ir_expr(.Construct, node.resolved_type, c)
}

// -- Helpers --

@(private = "file")
make_ir_expr :: proc(kind: IR_Expr_Kind, type: ^Resolved_Type, derived: IR_Expr_Derived) -> ^IR_Expr {
	e := new(IR_Expr)
	e.kind = kind
	e.type = type
	e.derived = derived
	return e
}

@(private = "file")
ast_binop_to_ir :: proc(op: Binary_Op) -> IR_Op {
	switch op {
	case .Add: return .Add
	case .Sub: return .Sub
	case .Mul: return .Mul
	case .Div: return .Div
	case .Mod: return .Mod
	case .Eq:  return .Eq
	case .Neq: return .Neq
	case .Lt:  return .Lt
	case .Gt:  return .Gt
	case .Lte: return .Lte
	case .Gte: return .Gte
	case .And: return .And
	case .Or:  return .Or
	}
	return .Add
}