sharkk/corvette public

branch main
showing the latest snapshot on main
emit_msl.odin 22.3 KB · Plain text
package shader

import "core:fmt"
import "core:strings"

// MSL 2.4 backend — emits from IR_Module

Msl_Emitter :: struct {
	w:                Writer,
	module:           ^IR_Module,
	current_fn:       ^IR_Function,
	diagnostics:      [dynamic]Diagnostic,
	emitted_structs:  map[string]bool,       // track emitted I/O struct names
	struct_remap:     map[string]string,      // original name -> remapped name for duplicates
}

emit_msl :: proc(module: ^IR_Module, allocator := context.allocator) -> (string, []Diagnostic) {
	e := Msl_Emitter{
		w               = writer_init(allocator),
		module          = module,
		diagnostics     = make([dynamic]Diagnostic, allocator),
		emitted_structs = make(map[string]bool, allocator = allocator),
		struct_remap    = make(map[string]string, allocator = allocator),
	}

	write_line(&e.w, "#include <metal_stdlib>")
	write_line(&e.w, "using namespace metal;")
	write_line(&e.w, "")

	// Specialization constants (Metal function constants)
	for &sc in module.spec_constants {
		write_line(&e.w, "constant ", resolved_type_to_msl(sc.type), " ", sc.name, " [[function_constant(", sc.spec_id, ")]];")
	}
	if len(module.spec_constants) > 0 {
		write_line(&e.w, "")
	}

	// Structs used by bindings (uniform blocks)
	for &b in module.bindings {
		if b.struct_ref != nil && (b.kind == .Uniform || b.kind == .Buffer) {
			emit_msl_struct_def(&e, b.struct_ref)
			write_line(&e.w, "")
		}
	}

	// I/O structs for entry points (skip compute — no struct I/O)
	for &fn in module.functions {
		if !fn.is_entry do continue
		if fn.stage == .Compute do continue
		emit_msl_io_structs(&e, &fn)
	}

	// Functions
	for &fn in module.functions {
		emit_msl_function(&e, &fn)
		write_line(&e.w, "")
	}

	return writer_to_string(e.w), e.diagnostics[:]
}

// -- Struct definitions --

@(private = "file")
emit_msl_struct_def :: proc(e: ^Msl_Emitter, s: ^Type_Struct_Resolved) {
	write_line(&e.w, "struct ", s.name, " {")
	indent(&e.w)
	for f in s.fields {
		write_line(&e.w, resolved_type_to_msl(f.type), " ", f.name, ";")
	}
	dedent(&e.w)
	write_line(&e.w, "};")
}

// -- I/O Structs --

@(private = "file")
emit_msl_io_structs :: proc(e: ^Msl_Emitter, fn: ^IR_Function) {
	if len(fn.inputs) > 0 {
		input_name := len(fn.params) > 0 ? msl_struct_name(fn.params[0].type) : "VertexInput"
		// Deduplicate: if this struct name was already emitted (e.g. @varying used as
		// both vertex output and fragment input), suffix with _In
		if input_name in e.emitted_structs {
			remapped := fmt.aprintf("%s_In", input_name)
			e.struct_remap[input_name] = remapped
			input_name = remapped
		}
		e.emitted_structs[input_name] = true
		write_line(&e.w, "struct ", input_name, " {")
		indent(&e.w)
		for io in fn.inputs {
			attr := msl_input_attribute(io, fn.stage)
			write_line(&e.w, resolved_type_to_msl(io.type), " ", io.name, " ", attr, ";")
		}
		dedent(&e.w)
		write_line(&e.w, "};")
		write_line(&e.w, "")
	}

	if len(fn.outputs) > 0 {
		output_name := fn.return_type != nil ? msl_struct_name(fn.return_type) : "FragmentOutput"
		if output_name in e.emitted_structs {
			remapped := fmt.aprintf("%s_Out", output_name)
			e.struct_remap[output_name] = remapped
			output_name = remapped
		}
		e.emitted_structs[output_name] = true
		write_line(&e.w, "struct ", output_name, " {")
		indent(&e.w)
		for io in fn.outputs {
			attr := msl_output_attribute(io, fn.stage)
			write_line(&e.w, resolved_type_to_msl(io.type), " ", io.name, " ", attr, ";")
		}
		dedent(&e.w)
		write_line(&e.w, "};")
		write_line(&e.w, "")
	}
}

// -- Functions --

@(private = "file")
emit_msl_function :: proc(e: ^Msl_Emitter, fn: ^IR_Function) {
	e.current_fn = fn

	if fn.is_entry {
		emit_msl_entry_point(e, fn)
	} else {
		emit_msl_helper_function(e, fn)
	}

	e.current_fn = nil
}

@(private = "file")
emit_msl_entry_point :: proc(e: ^Msl_Emitter, fn: ^IR_Function) {
	// Compute shader — kernel function with builtin params
	if fn.stage == .Compute {
		emit_msl_compute_entry(e, fn)
		return
	}

	input_name := len(fn.params) > 0 ? msl_struct_name(fn.params[0].type) : "VertexInput"
	if remapped, ok := e.struct_remap[input_name]; ok {
		input_name = remapped
	}
	output_name := fn.return_type != nil ? msl_struct_name(fn.return_type) : "FragmentOutput"
	if remapped, ok := e.struct_remap[output_name]; ok {
		output_name = remapped
	}

	// Stage qualifier
	stage_str: string
	#partial switch fn.stage {
	case .Vertex:   stage_str = "vertex"
	case .Fragment: stage_str = "fragment"
	case .None:     stage_str = "vertex"
	}

	// Collect bindings referenced by this entry point
	binding_refs := collect_binding_refs(e.module, fn)

	// Emit function signature
	for _ in 0 ..< e.w.indent do write(&e.w, "\t")
	write(&e.w, stage_str, " ", output_name, " ", fn.name, "(")

	// I/O parameter
	param_name := len(fn.params) > 0 ? fn.params[0].name : "input"
	write(&e.w, input_name, " ", param_name, " [[stage_in]]")

	// Binding parameters
	for &br in binding_refs {
		write(&e.w, ",\n")
		for _ in 0 ..< e.w.indent + 1 do write(&e.w, "\t")
		emit_msl_binding_param(e, &br)
	}

	write(&e.w, ") {\n")
	indent(&e.w)

	if len(fn.outputs) > 0 {
		write_line(&e.w, output_name, " __luma_output;")
	}
	emit_msl_stmts(e, fn.body[:])
	if len(fn.outputs) > 0 {
		write_line(&e.w, "return __luma_output;")
	}

	dedent(&e.w)
	write_line(&e.w, "}")
}

@(private = "file")
emit_msl_compute_entry :: proc(e: ^Msl_Emitter, fn: ^IR_Function) {
	// Collect bindings referenced by this entry point
	binding_refs := collect_binding_refs(e.module, fn)

	// Emit function signature
	write(&e.w, "kernel void ", fn.name, "(")
	first := true

	// Builtin parameters
	for io in fn.inputs {
		if io.builtin == "" do continue
		if !first do write(&e.w, ",\n\t")
		first = false
		write(&e.w, resolved_type_to_msl(io.type), " ", io.name, " ", msl_builtin_attribute(io.builtin))
	}

	// Binding parameters
	for &br in binding_refs {
		if !first do write(&e.w, ",\n\t")
		first = false
		emit_msl_binding_param(e, &br)
	}

	write(&e.w, ") {\n")
	indent(&e.w)

	// Emit threadgroup shared variables inside kernel body
	for sv in e.module.shared_vars {
		base, suffix := msl_type_and_array_suffix(sv.type)
		write_line(&e.w, "threadgroup ", base, " ", sv.name, suffix, ";")
	}
	if len(e.module.shared_vars) > 0 {
		write_line(&e.w, "")
	}

	emit_msl_stmts(e, fn.body[:])

	dedent(&e.w)
	write_line(&e.w, "}")
}

@(private = "file")
emit_msl_helper_function :: proc(e: ^Msl_Emitter, fn: ^IR_Function) {
	for _ in 0 ..< e.w.indent do write(&e.w, "\t")
	write(&e.w, resolved_type_to_msl(fn.return_type), " ", fn.name, "(")
	for p, i in fn.params {
		if i > 0 do write(&e.w, ", ")
		write(&e.w, resolved_type_to_msl(p.type), " ", p.name)
	}
	write(&e.w, ") {\n")
	indent(&e.w)

	emit_msl_stmts(e, fn.body[:])

	dedent(&e.w)
	write_line(&e.w, "}")
}

// -- Binding parameters --

Msl_Binding_Ref :: struct {
	binding: ^IR_Binding,
	index:   int, // contiguous index per resource class
}

@(private = "file")
collect_binding_refs :: proc(module: ^IR_Module, fn: ^IR_Function) -> [dynamic]Msl_Binding_Ref {
	// Walk function body to find all IR_Load_Binding references
	used_names := make(map[string]bool)
	collect_binding_names_from_stmts(fn.body[:], &used_names)

	// Also include texture/sampler pairs — if texture is used, sampler is too
	for &b in module.bindings {
		if b.combined_name != "" && used_names[b.combined_name] {
			used_names[b.name] = true
		}
	}

	// Build refs with contiguous per-class indices
	refs := make([dynamic]Msl_Binding_Ref)
	next_buffer := 0
	next_texture := 0
	next_sampler := 0

	for &b in module.bindings {
		if !used_names[b.name] && !used_names[b.combined_name] do continue

		ref := Msl_Binding_Ref{binding = &b}
		switch b.kind {
		case .Uniform, .Buffer, .Push_Constant:
			ref.index = next_buffer
			next_buffer += 1
		case .Texture:
			ref.index = next_texture
			next_texture += 1
		case .Sampler:
			ref.index = next_sampler
			next_sampler += 1
		}
		append(&refs, ref)
	}
	return refs
}

@(private = "file")
collect_binding_names_from_stmts :: proc(stmts: []IR_Stmt, names: ^map[string]bool) {
	for stmt in stmts {
		switch s in stmt {
		case ^IR_Let:
			collect_binding_names_from_expr(s.value, names)
		case ^IR_Assign:
			collect_binding_names_from_expr(s.target, names)
			collect_binding_names_from_expr(s.value, names)
		case ^IR_Return:
			collect_binding_names_from_expr(s.value, names)
		case ^IR_Store_Output:
			collect_binding_names_from_expr(s.value, names)
		case ^IR_If:
			collect_binding_names_from_expr(s.condition, names)
			collect_binding_names_from_stmts(s.then_body[:], names)
			for ei in s.elseif_clauses {
				collect_binding_names_from_expr(ei.condition, names)
				collect_binding_names_from_stmts(ei.body[:], names)
			}
			collect_binding_names_from_stmts(s.else_body[:], names)
		case ^IR_For:
			collect_binding_names_from_expr(s.start, names)
			collect_binding_names_from_expr(s.stop, names)
			collect_binding_names_from_expr(s.step, names)
			collect_binding_names_from_stmts(s.body[:], names)
		case ^IR_While:
			collect_binding_names_from_expr(s.condition, names)
			collect_binding_names_from_stmts(s.body[:], names)
		case ^IR_Expr_Stmt:
			collect_binding_names_from_expr(s.expr, names)
		case ^IR_Barrier:
			// no binding references
		case ^IR_Discard:
			// no binding references
		case ^IR_Break:
			// no binding references
		case ^IR_Continue:
			// no binding references
		}
	}
}

@(private = "file")
collect_binding_names_from_expr :: proc(expr: ^IR_Expr, names: ^map[string]bool) {
	if expr == nil do return
	switch d in expr.derived {
	case ^IR_Load_Binding:
		names[d.name] = true
	case ^IR_Binary:
		collect_binding_names_from_expr(d.left, names)
		collect_binding_names_from_expr(d.right, names)
	case ^IR_Unary:
		collect_binding_names_from_expr(d.operand, names)
	case ^IR_Call:
		for arg in d.args {
			collect_binding_names_from_expr(arg, names)
		}
	case ^IR_Field_Access:
		collect_binding_names_from_expr(d.object, names)
	case ^IR_Swizzle:
		collect_binding_names_from_expr(d.object, names)
	case ^IR_Composite_Extract:
		collect_binding_names_from_expr(d.object, names)
	case ^IR_Vector_Shuffle:
		collect_binding_names_from_expr(d.object, names)
	case ^IR_Index:
		collect_binding_names_from_expr(d.object, names)
		collect_binding_names_from_expr(d.index, names)
	case ^IR_Construct:
		for arg in d.args {
			collect_binding_names_from_expr(arg, names)
		}
	case ^IR_Type_Cast:
		collect_binding_names_from_expr(d.value, names)
	case ^IR_Select:
		collect_binding_names_from_expr(d.condition, names)
		collect_binding_names_from_expr(d.true_val, names)
		collect_binding_names_from_expr(d.false_val, names)
	case ^IR_Literal, ^IR_Var_Ref, ^IR_Input_Field, ^IR_Builtin_Var, ^IR_Shared_Ref:
		// No bindings
	}
}

@(private = "file")
emit_msl_binding_param :: proc(e: ^Msl_Emitter, ref: ^Msl_Binding_Ref) {
	b := ref.binding
	switch b.kind {
	case .Uniform:
		type_name := b.struct_ref != nil ? b.struct_ref.name : resolved_type_to_msl(b.type)
		write(&e.w, "constant ", type_name, "& ", b.name, " [[buffer(", ref.index, ")]]")
	case .Buffer:
		type_name := b.struct_ref != nil ? b.struct_ref.name : resolved_type_to_msl(b.type)
		write(&e.w, "device ", type_name, "* ", b.name, " [[buffer(", ref.index, ")]]")
	case .Texture:
		msl_tex := msl_texture_type(b.type)
		write(&e.w, msl_tex, " ", b.name, " [[texture(", ref.index, ")]]")
	case .Sampler:
		write(&e.w, "sampler ", b.name, " [[sampler(", ref.index, ")]]")
	case .Push_Constant:
		type_name := b.struct_ref != nil ? b.struct_ref.name : resolved_type_to_msl(b.type)
		write(&e.w, "constant ", type_name, "& ", b.name, " [[buffer(", ref.index, ")]]")
	}
}

// -- Statements --

@(private = "file")
emit_msl_stmts :: proc(e: ^Msl_Emitter, stmts: []IR_Stmt) {
	for stmt in stmts {
		emit_msl_stmt(e, stmt)
	}
}

@(private = "file")
emit_msl_stmt :: proc(e: ^Msl_Emitter, stmt: IR_Stmt) {
	switch s in stmt {
	case ^IR_Let:
		for _ in 0 ..< e.w.indent do write(&e.w, "\t")
		write(&e.w, resolved_type_to_msl(s.type), " ", s.name, " = ")
		emit_msl_expr(e, s.value)
		write(&e.w, ";\n")

	case ^IR_Assign:
		for _ in 0 ..< e.w.indent do write(&e.w, "\t")
		emit_msl_expr(e, s.target)
		write(&e.w, " = ")
		emit_msl_expr(e, s.value)
		write(&e.w, ";\n")

	case ^IR_Return:
		if s.value != nil {
			for _ in 0 ..< e.w.indent do write(&e.w, "\t")
			write(&e.w, "return ")
			emit_msl_expr(e, s.value)
			write(&e.w, ";\n")
		} else {
			write_line(&e.w, "return;")
		}

	case ^IR_Store_Output:
		fn := e.current_fn
		if fn != nil && s.io_index >= 0 && s.io_index < len(fn.outputs) {
			io := fn.outputs[s.io_index]
			for _ in 0 ..< e.w.indent do write(&e.w, "\t")
			write(&e.w, "__luma_output.", io.name, " = ")
			emit_msl_expr(e, s.value)
			write(&e.w, ";\n")
		}

	case ^IR_If:
		for _ in 0 ..< e.w.indent do write(&e.w, "\t")
		write(&e.w, "if (")
		emit_msl_expr(e, s.condition)
		write(&e.w, ") {\n")
		indent(&e.w)
		emit_msl_stmts(e, s.then_body[:])
		dedent(&e.w)
		for ei in s.elseif_clauses {
			for _ in 0 ..< e.w.indent do write(&e.w, "\t")
			write(&e.w, "} else if (")
			emit_msl_expr(e, ei.condition)
			write(&e.w, ") {\n")
			indent(&e.w)
			emit_msl_stmts(e, ei.body[:])
			dedent(&e.w)
		}
		if len(s.else_body) > 0 {
			write_line(&e.w, "} else {")
			indent(&e.w)
			emit_msl_stmts(e, s.else_body[:])
			dedent(&e.w)
		}
		write_line(&e.w, "}")

	case ^IR_For:
		for _ in 0 ..< e.w.indent do write(&e.w, "\t")
		write(&e.w, "for (int ", s.var_name, " = ")
		emit_msl_expr(e, s.start)
		write(&e.w, "; ", s.var_name, " <= ")
		emit_msl_expr(e, s.stop)
		write(&e.w, "; ", s.var_name)
		if s.step != nil {
			write(&e.w, " += ")
			emit_msl_expr(e, s.step)
		} else {
			write(&e.w, "++")
		}
		write(&e.w, ") {\n")
		indent(&e.w)
		emit_msl_stmts(e, s.body[:])
		dedent(&e.w)
		write_line(&e.w, "}")

	case ^IR_While:
		for _ in 0 ..< e.w.indent do write(&e.w, "\t")
		write(&e.w, "while (")
		emit_msl_expr(e, s.condition)
		write(&e.w, ") {\n")
		indent(&e.w)
		emit_msl_stmts(e, s.body[:])
		dedent(&e.w)
		write_line(&e.w, "}")

	case ^IR_Expr_Stmt:
		for _ in 0 ..< e.w.indent do write(&e.w, "\t")
		emit_msl_expr(e, s.expr)
		write(&e.w, ";\n")

	case ^IR_Barrier:
		write_line(&e.w, "threadgroup_barrier(mem_flags::mem_threadgroup);")

	case ^IR_Discard:
		write_line(&e.w, "discard_fragment();")

	case ^IR_Break:
		write_line(&e.w, "break;")

	case ^IR_Continue:
		write_line(&e.w, "continue;")
	}
}

// -- Expressions --

@(private = "file")
emit_msl_expr :: proc(e: ^Msl_Emitter, expr: ^IR_Expr) {
	if expr == nil {
		write(&e.w, "/* nil */")
		return
	}

	switch d in expr.derived {
	case ^IR_Literal:
		switch v in d.value {
		case i64:
			write(&e.w, v)
		case f64:
			s := fmt.aprintf("%v", v)
			if !strings.contains(s, ".") && !strings.contains(s, "e") {
				write(&e.w, s, ".0")
			} else {
				write(&e.w, s)
			}
		case bool:
			write(&e.w, v ? "true" : "false")
		}

	case ^IR_Var_Ref:
		write(&e.w, d.name)

	case ^IR_Binary:
		write(&e.w, "(")
		emit_msl_expr(e, d.left)
		write(&e.w, " ", ir_op_to_msl(d.op), " ")
		emit_msl_expr(e, d.right)
		write(&e.w, ")")

	case ^IR_Unary:
		if d.op == .Neg {
			write(&e.w, "(-")
		} else {
			write(&e.w, "(!")
		}
		emit_msl_expr(e, d.operand)
		write(&e.w, ")")

	case ^IR_Call:
		msl_name := d.is_builtin ? builtin_to_msl(d.name) : d.name
		// Handle texture sampling — MSL uses tex.method(sampler, uv, ...) syntax
		if d.is_builtin && (d.name == "sample" || d.name == "sample_level" || d.name == "sample_shadow") && len(d.args) >= 2 {
			tex_name, samp_name: string
			if lb, ok := d.args[0].derived.(^IR_Load_Binding); ok {
				tex_name, samp_name = find_split_bindings(e.module, lb.name)
			} else {
				tex_name, samp_name = "tex", "default_sampler"
			}
			if d.name == "sample" {
				write(&e.w, tex_name, ".sample(", samp_name, ", ")
				emit_msl_expr(e, d.args[1])
				write(&e.w, ")")
			} else if d.name == "sample_level" && len(d.args) >= 3 {
				write(&e.w, tex_name, ".sample(", samp_name, ", ")
				emit_msl_expr(e, d.args[1])
				write(&e.w, ", level(")
				emit_msl_expr(e, d.args[2])
				write(&e.w, "))")
			} else if d.name == "sample_shadow" && len(d.args) >= 3 {
				write(&e.w, tex_name, ".sample_compare(", samp_name, ", ")
				emit_msl_expr(e, d.args[1])
				write(&e.w, ", ")
				emit_msl_expr(e, d.args[2])
				write(&e.w, ")")
			}
			return
		}
		write(&e.w, msl_name, "(")
		for arg, i in d.args {
			if i > 0 do write(&e.w, ", ")
			emit_msl_expr(e, arg)
		}
		write(&e.w, ")")

	case ^IR_Field_Access:
		emit_msl_expr(e, d.object)
		write(&e.w, ".", d.field_name)

	case ^IR_Swizzle:
		emit_msl_expr(e, d.object)
		write(&e.w, ".", d.components)

	case ^IR_Composite_Extract:
		emit_msl_expr(e, d.object)
		write(&e.w, ".", d.field_name)

	case ^IR_Vector_Shuffle:
		emit_msl_expr(e, d.object)
		write(&e.w, ".", swizzle_indices_to_string(d.components))

	case ^IR_Index:
		emit_msl_expr(e, d.object)
		write(&e.w, "[")
		emit_msl_expr(e, d.index)
		write(&e.w, "]")

	case ^IR_Construct:
		msl_name := resolved_type_to_msl(expr.type)
		write(&e.w, msl_name, "(")
		for arg, i in d.args {
			if i > 0 do write(&e.w, ", ")
			emit_msl_expr(e, arg)
		}
		write(&e.w, ")")

	case ^IR_Type_Cast:
		write(&e.w, "static_cast<", resolved_type_to_msl(expr.type), ">(")
		emit_msl_expr(e, d.value)
		write(&e.w, ")")

	case ^IR_Load_Binding:
		write(&e.w, d.name)

	case ^IR_Input_Field:
		// MSL entry points pass struct via [[stage_in]]
		write(&e.w, d.param_name, ".", d.field_name)

	case ^IR_Builtin_Var:
		fn := e.current_fn
		if fn != nil {
			io_list := d.is_input ? fn.inputs[:] : fn.outputs[:]
			for io in io_list {
				if io.builtin == d.name {
					// Compute shaders use inline params, not struct access
					if fn.stage == .Compute {
						write(&e.w, io.name)
					} else if d.is_input && len(fn.params) > 0 {
						write(&e.w, fn.params[0].name, ".", io.name)
					} else {
						write(&e.w, "output.", io.name)
					}
					return
				}
			}
		}
		write(&e.w, d.name)

	case ^IR_Shared_Ref:
		write(&e.w, d.name)

	case ^IR_Select:
		write(&e.w, "select(")
		emit_msl_expr(e, d.false_val)
		write(&e.w, ", ")
		emit_msl_expr(e, d.true_val)
		write(&e.w, ", ")
		emit_msl_expr(e, d.condition)
		write(&e.w, ")")
	}
}

// -- Helpers --

@(private = "file")
ir_op_to_msl :: proc(op: IR_Op) -> string {
	switch op {
	case .Add: return "+"
	case .Sub: return "-"
	case .Mul: return "*"
	case .Div: return "/"
	case .Mod: return "%"
	case .Eq:  return "=="
	case .Neq: return "!="
	case .Lt:  return "<"
	case .Gt:  return ">"
	case .Lte: return "<="
	case .Gte: return ">="
	case .And: return "&&"
	case .Or:  return "||"
	case .Neg: return "-"
	case .Not: return "!"
	}
	return "?"
}

@(private = "file")
builtin_to_msl :: proc(name: string) -> string {
	switch name {
	case "sample":          return "sample" // handled specially
	case "sample_level":    return "sample" // would use level() arg
	case "mix":             return "mix"
	case "fract":           return "fract"
	case "mod":             return "fmod"
	case "inversesqrt":     return "rsqrt"
	case "dfdx":            return "dfdx"
	case "dfdy":            return "dfdy"
	case "atan2":           return "atan2"
	case "texture_size":    return "get_width" // MSL uses method calls
	}
	return name
}

@(private = "file")
msl_input_attribute :: proc(io: IR_IO_Var, stage: Shader_Stage) -> string {
	if io.builtin != "" {
		return msl_builtin_attribute(io.builtin)
	}
	if stage == .Vertex {
		return fmt.aprintf("[[attribute(%d)]]", io.location)
	}
	// Inter-stage (fragment inputs)
	return fmt.aprintf("[[user(locn%d)]]", io.location)
}

@(private = "file")
msl_output_attribute :: proc(io: IR_IO_Var, stage: Shader_Stage) -> string {
	if io.builtin != "" {
		return msl_builtin_attribute(io.builtin)
	}
	// Inter-stage outputs
	return fmt.aprintf("[[user(locn%d)]]", io.location)
}

@(private = "file")
msl_builtin_attribute :: proc(name: string) -> string {
	switch name {
	case "position":              return "[[position]]"
	case "vertex_id":             return "[[vertex_id]]"
	case "instance_id":           return "[[instance_id]]"
	case "frag_coord":            return "[[position]]"
	case "front_facing":          return "[[front_facing]]"
	case "local_invocation_id":   return "[[thread_position_in_threadgroup]]"
	case "local_invocation_index": return "[[thread_index_in_threadgroup]]"
	case "global_invocation_id":  return "[[thread_position_in_grid]]"
	case "workgroup_id":          return "[[threadgroup_position_in_grid]]"
	}
	return fmt.aprintf("[[user(%s)]]", name)
}

@(private = "file")
msl_texture_type :: proc(t: ^Resolved_Type) -> string {
	if t == nil do return "texture2d<float>"
	#partial switch v in t^ {
	case Type_Sampler:
		switch v.kind {
		case .Sampler2D:       return "texture2d<float>"
		case .Sampler3D:       return "texture3d<float>"
		case .SamplerCube:     return "texturecube<float>"
		case .Sampler2DArray:  return "texture2d_array<float>"
		case .Sampler2DShadow: return "depth2d<float>"
		}
	}
	return "texture2d<float>"
}

resolved_type_to_msl :: proc(t: ^Resolved_Type) -> string {
	if t == nil do return "void"
	switch v in t^ {
	case Type_Scalar:
		switch v.kind {
		case .Bool:  return "bool"
		case .Int:   return "int"
		case .Uint:  return "uint"
		case .Float: return "float"
		case .Half:  return "half"
		}
	case Type_Vector:
		elem: string
		switch v.elem {
		case .Float: elem = "float"
		case .Int:   elem = "int"
		case .Uint:  elem = "uint"
		case .Bool:  elem = "bool"
		case .Half:  elem = "half"
		}
		return fmt.aprintf("%s%d", elem, v.size)
	case Type_Matrix:
		elem: string
		#partial switch v.elem {
		case .Float: elem = "float"
		case .Half:  elem = "half"
		case:        elem = "float"
		}
		return fmt.aprintf("%s%dx%d", elem, v.cols, v.rows)
	case Type_Struct_Resolved:
		return v.name
	case Type_Array_Resolved:
		elem_str := resolved_type_to_msl(v.elem)
		if v.size == 0 {
			return fmt.aprintf("device %s*", elem_str) // unsized arrays are pointers in MSL
		}
		return fmt.aprintf("array<%s, %d>", elem_str, v.size)
	case Type_Sampler:
		return msl_texture_type(t)
	case Type_Void:
		return "void"
	}
	return "void"
}

@(private = "file")
msl_struct_name :: proc(t: ^Resolved_Type) -> string {
	if t == nil do return "Unknown"
	#partial switch v in t^ {
	case Type_Struct_Resolved:
		return v.name
	}
	return type_to_string(t)
}

@(private = "file")
msl_type_and_array_suffix :: proc(t: ^Resolved_Type) -> (base: string, suffix: string) {
	if t == nil do return "void", ""
	#partial switch v in t^ {
	case Type_Array_Resolved:
		elem_str := resolved_type_to_msl(v.elem)
		if v.size == 0 {
			return elem_str, "[]"
		}
		return elem_str, fmt.aprintf("[%d]", v.size)
	}
	return resolved_type_to_msl(t), ""
}