sharkk/corvette public

branch main
showing the latest snapshot on main
emit_hlsl.odin 21.0 KB · Plain text
package shader

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

// HLSL SM6.0 backend — emits from IR_Module

Hlsl_Emitter :: struct {
	w:                Writer,
	module:           ^IR_Module,
	options:          ^Compile_Options,
	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_hlsl :: proc(module: ^IR_Module, options: ^Compile_Options = nil, allocator := context.allocator) -> (string, []Diagnostic) {
	e := Hlsl_Emitter{
		w               = writer_init(allocator),
		module          = module,
		options         = options,
		diagnostics     = make([dynamic]Diagnostic, allocator),
		emitted_structs = make(map[string]bool, allocator = allocator),
		struct_remap    = make(map[string]string, allocator = allocator),
	}

	// Specialization constants (Vulkan HLSL extension)
	for &sc in module.spec_constants {
		write_line(&e.w, "[[vk::constant_id(", sc.spec_id, ")]] const ", resolved_type_to_hlsl(sc.type), " ", sc.name, " = ", ir_const_value_to_string(sc.default_value), ";")
	}
	if len(module.spec_constants) > 0 {
		write_line(&e.w, "")
	}

	// Shared variables (groupshared)
	for sv in module.shared_vars {
		base, suffix := hlsl_type_and_array_suffix(sv.type)
		write_line(&e.w, "groupshared ", base, " ", sv.name, suffix, ";")

	}
	if len(module.shared_vars) > 0 {
		write_line(&e.w, "")
	}

	// Bindings (cbuffers, textures, samplers)
	for &b in module.bindings {
		emit_hlsl_binding(&e, &b)
	}
	if len(module.bindings) > 0 {
		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_hlsl_io_structs(&e, &fn)
	}

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

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

// -- Bindings --

@(private = "file")
hlsl_register :: proc(e: ^Hlsl_Emitter, prefix: string, num: int, group: int) {
	omit_spaces := e.options != nil && e.options.hlsl_omit_register_spaces
	if omit_spaces {
		write(&e.w, "register(", prefix, num, ")")
	} else {
		write(&e.w, "register(", prefix, num, ", space", group, ")")
	}
}

@(private = "file")
emit_hlsl_binding :: proc(e: ^Hlsl_Emitter, b: ^IR_Binding) {
	switch b.kind {
	case .Texture:
		write(&e.w, "Texture2D ", b.name, " : ")
		hlsl_register(e, "t", b.binding_num, b.group)
		write(&e.w, ";\n")
	case .Sampler:
		write(&e.w, "SamplerState ", b.name, " : ")
		hlsl_register(e, "s", b.binding_num, b.group)
		write(&e.w, ";\n")
	case .Uniform:
		if b.struct_ref != nil {
			// Emit struct definition + cbuffer containing it, so binding.field works
			write_line(&e.w, "struct ", b.struct_ref.name, " {")
			indent(&e.w)
			for f in b.struct_ref.fields {
				base, suffix := hlsl_type_and_array_suffix(f.type)
				write_line(&e.w, base, " ", f.name, suffix, ";")
			}
			dedent(&e.w)
			write_line(&e.w, "};")
			write(&e.w, "cbuffer ", b.name, "_CB : ")
			hlsl_register(e, "b", b.binding_num, b.group)
			write(&e.w, " {\n")
			indent(&e.w)
			write_line(&e.w, b.struct_ref.name, " ", b.name, ";")
			dedent(&e.w)
			write_line(&e.w, "};")
		} else {
			// Bare uniform (e.g. uniform light_dir: vec3) — wrap as single field in cbuffer
			write(&e.w, "cbuffer ", b.name, "_CB : ")
			hlsl_register(e, "b", b.binding_num, b.group)
			write(&e.w, " {\n")
			indent(&e.w)
			base, suffix := hlsl_type_and_array_suffix(b.type)
			write_line(&e.w, base, " ", b.name, suffix, ";")
			dedent(&e.w)
			write_line(&e.w, "};")
		}
	case .Buffer:
		// Storage buffer — emit struct definition if needed
		if b.struct_ref != nil {
			if !(b.struct_ref.name in e.emitted_structs) {
				e.emitted_structs[b.struct_ref.name] = true
				write_line(&e.w, "struct ", b.struct_ref.name, " {")
				indent(&e.w)
				for f in b.struct_ref.fields {
					base, suffix := hlsl_type_and_array_suffix(f.type)
					write_line(&e.w, base, " ", f.name, suffix, ";")
				}
				dedent(&e.w)
				write_line(&e.w, "};")
			}
			write(&e.w, "RWStructuredBuffer<", b.struct_ref.name, "> ", b.name, " : ")
			hlsl_register(e, "u", b.binding_num, b.group)
			write(&e.w, ";\n")
		} else {
			write(&e.w, "RWByteAddressBuffer ", b.name, " : ")
			hlsl_register(e, "u", b.binding_num, b.group)
			write(&e.w, ";\n")
		}
	case .Push_Constant:
		if b.struct_ref != nil {
			if !(b.struct_ref.name in e.emitted_structs) {
				e.emitted_structs[b.struct_ref.name] = true
				write_line(&e.w, "struct ", b.struct_ref.name, " {")
				indent(&e.w)
				for f in b.struct_ref.fields {
					base, suffix := hlsl_type_and_array_suffix(f.type)
					write_line(&e.w, base, " ", f.name, suffix, ";")
				}
				dedent(&e.w)
				write_line(&e.w, "};")
			}
			use_cbuffer := e.options != nil && e.options.hlsl_cbuffer_push_constants
			if use_cbuffer {
				slot := e.options.hlsl_push_constant_slot
				write(&e.w, "cbuffer ", b.name, "_CB : ")
				hlsl_register(e, "b", slot, 0)
				write(&e.w, " {\n")
				indent(&e.w)
				write_line(&e.w, b.struct_ref.name, " ", b.name, ";")
				dedent(&e.w)
				write_line(&e.w, "};")
			} else {
				write_line(&e.w, "[[vk::push_constant]] ", b.struct_ref.name, " ", b.name, ";")
			}
		}
	}
}

// -- I/O Structs --

@(private = "file")
emit_hlsl_io_structs :: proc(e: ^Hlsl_Emitter, fn: ^IR_Function) {
	if len(fn.inputs) > 0 {
		input_name := len(fn.params) > 0 ? hlsl_struct_name(fn.params[0].type) : "VS_INPUT"
		// Deduplicate: if this struct name was already emitted (e.g. @varying used as
		// both vertex output and fragment input), the FS reuses the VS output struct.
		// HLSL allows PS input to have fields the shader doesn't read, so the VS output
		// struct (which includes SV_Position) works directly as the PS input struct.
		if input_name in e.emitted_structs {
			if fn.stage == .Fragment {
				// Skip emitting a new struct — FS reuses the VS output struct as-is.
				// Don't add to struct_remap so VS keeps using the original name too.
			} else {
				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 {
					semantic := hlsl_input_semantic(io, fn.stage)
					write_line(&e.w, resolved_type_to_hlsl(io.type), " ", io.name, " : ", semantic, ";")
				}
				dedent(&e.w)
				write_line(&e.w, "};")
				write_line(&e.w, "")
			}
		} else {
			e.emitted_structs[input_name] = true
			write_line(&e.w, "struct ", input_name, " {")
			indent(&e.w)
			// For fragment shaders, inject SV_Position if not already present
			if fn.stage == .Fragment {
				has_sv_position := false
				for io in fn.inputs {
					if io.builtin == "frag_coord" {
						has_sv_position = true
						break
					}
				}
				if !has_sv_position {
					write_line(&e.w, "float4 _sv_position : SV_Position;")
				}
			}
			for io in fn.inputs {
				semantic := hlsl_input_semantic(io, fn.stage)
				write_line(&e.w, resolved_type_to_hlsl(io.type), " ", io.name, " : ", semantic, ";")
			}
			dedent(&e.w)
			write_line(&e.w, "};")
			write_line(&e.w, "")
		}
	}

	if len(fn.outputs) > 0 {
		output_name := fn.return_type != nil ? hlsl_struct_name(fn.return_type) : "PS_OUTPUT"
		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 {
			semantic := hlsl_output_semantic(io, fn.stage)
			write_line(&e.w, resolved_type_to_hlsl(io.type), " ", io.name, " : ", semantic, ";")
		}
		dedent(&e.w)
		write_line(&e.w, "};")
		write_line(&e.w, "")
	}
}

// -- Functions --

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

	if fn.is_entry {
		emit_hlsl_entry_point(e, fn)
	} else {
		emit_hlsl_helper_function(e, fn)
	}

	e.current_fn = nil
}

@(private = "file")
emit_hlsl_entry_point :: proc(e: ^Hlsl_Emitter, fn: ^IR_Function) {
	// Compute shader — emit [numthreads] and builtin params
	if fn.stage == .Compute {
		ws := fn.workgroup_size
		write_line(&e.w, "[numthreads(", ws[0], ", ", ws[1] > 0 ? ws[1] : 1, ", ", ws[2] > 0 ? ws[2] : 1, ")]")
		write(&e.w, "void ", fn.name, "(")
		first := true
		for io in fn.inputs {
			if io.builtin == "" do continue
			if !first do write(&e.w, ", ")
			first = false
			write(&e.w, resolved_type_to_hlsl(io.type), " ", io.name, " : ", hlsl_builtin_semantic(io.builtin, true))
		}
		write(&e.w, ") {\n")
		indent(&e.w)
		emit_hlsl_stmts(e, fn.body[:])
		dedent(&e.w)
		write_line(&e.w, "}")
		return
	}

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

	for _ in 0 ..< e.w.indent do write(&e.w, "\t")
	write(&e.w, output_name, " ", fn.name, "(", input_name, " ")
	if len(fn.params) > 0 {
		write(&e.w, fn.params[0].name)
	} else {
		write(&e.w, "input")
	}
	write(&e.w, ") {\n")
	indent(&e.w)

	if len(fn.outputs) > 0 {
		write_line(&e.w, output_name, " __luma_output;")
	}
	emit_hlsl_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_hlsl_helper_function :: proc(e: ^Hlsl_Emitter, fn: ^IR_Function) {
	for _ in 0 ..< e.w.indent do write(&e.w, "\t")
	write(&e.w, resolved_type_to_hlsl(fn.return_type), " ", fn.name, "(")
	for p, i in fn.params {
		if i > 0 do write(&e.w, ", ")
		write(&e.w, resolved_type_to_hlsl(p.type), " ", p.name)
	}
	write(&e.w, ") {\n")
	indent(&e.w)

	emit_hlsl_stmts(e, fn.body[:])

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

// -- Statements --

@(private = "file")
emit_hlsl_stmts :: proc(e: ^Hlsl_Emitter, stmts: []IR_Stmt) {
	for stmt in stmts {
		emit_hlsl_stmt(e, stmt)
	}
}

@(private = "file")
emit_hlsl_stmt :: proc(e: ^Hlsl_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_hlsl(s.type), " ", s.name, " = ")
		emit_hlsl_expr(e, s.value)
		write(&e.w, ";\n")

	case ^IR_Assign:
		for _ in 0 ..< e.w.indent do write(&e.w, "\t")
		emit_hlsl_expr(e, s.target)
		write(&e.w, " = ")
		emit_hlsl_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_hlsl_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_hlsl_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_hlsl_expr(e, s.condition)
		write(&e.w, ") {\n")
		indent(&e.w)
		emit_hlsl_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_hlsl_expr(e, ei.condition)
			write(&e.w, ") {\n")
			indent(&e.w)
			emit_hlsl_stmts(e, ei.body[:])
			dedent(&e.w)
		}
		if len(s.else_body) > 0 {
			write_line(&e.w, "} else {")
			indent(&e.w)
			emit_hlsl_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_hlsl_expr(e, s.start)
		write(&e.w, "; ", s.var_name, " <= ")
		emit_hlsl_expr(e, s.stop)
		write(&e.w, "; ", s.var_name)
		if s.step != nil {
			write(&e.w, " += ")
			emit_hlsl_expr(e, s.step)
		} else {
			write(&e.w, "++")
		}
		write(&e.w, ") {\n")
		indent(&e.w)
		emit_hlsl_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_hlsl_expr(e, s.condition)
		write(&e.w, ") {\n")
		indent(&e.w)
		emit_hlsl_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_hlsl_expr(e, s.expr)
		write(&e.w, ";\n")

	case ^IR_Barrier:
		write_line(&e.w, "GroupMemoryBarrierWithGroupSync();")

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

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

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

// -- Expressions --

@(private = "file")
emit_hlsl_expr :: proc(e: ^Hlsl_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:
		// HLSL requires mul() for matrix operations
		if d.op == .Mul && needs_hlsl_mul(d.left, d.right) {
			write(&e.w, "mul(")
			emit_hlsl_expr(e, d.left)
			write(&e.w, ", ")
			emit_hlsl_expr(e, d.right)
			write(&e.w, ")")
		} else {
			write(&e.w, "(")
			emit_hlsl_expr(e, d.left)
			write(&e.w, " ", ir_op_to_hlsl(d.op), " ")
			emit_hlsl_expr(e, d.right)
			write(&e.w, ")")
		}

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

	case ^IR_Call:
		hlsl_name := d.is_builtin ? builtin_to_hlsl(d.name) : d.name
		// Handle texture sampling — HLSL uses tex.Method(sampler, ...) syntax
		if d.is_builtin && (d.name == "sample" || d.name == "sample_level" || d.name == "sample_shadow") && len(d.args) >= 2 {
			method := d.name == "sample" ? "Sample" : (d.name == "sample_level" ? "SampleLevel" : "SampleCmpLevelZero")
			if lb, ok := d.args[0].derived.(^IR_Load_Binding); ok {
				tex_name, samp_name := find_split_bindings(e.module, lb.name)
				write(&e.w, tex_name, ".", method, "(", samp_name, ", ")
			} else {
				emit_hlsl_expr(e, d.args[0])
				write(&e.w, ".", method, "(default_samp, ")
			}
			emit_hlsl_expr(e, d.args[1])
			// Extra args after uv (lod for sample_level, ref for sample_shadow)
			for i := 2; i < len(d.args); i += 1 {
				write(&e.w, ", ")
				emit_hlsl_expr(e, d.args[i])
			}
			write(&e.w, ")")
			return
		}
		write(&e.w, hlsl_name, "(")
		for arg, i in d.args {
			if i > 0 do write(&e.w, ", ")
			emit_hlsl_expr(e, arg)
		}
		write(&e.w, ")")

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

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

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

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

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

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

	case ^IR_Type_Cast:
		write(&e.w, "(", resolved_type_to_hlsl(expr.type), ")")
		emit_hlsl_expr(e, d.value)

	case ^IR_Load_Binding:
		// For uniform blocks, we access fields directly with the binding prefix
		// Check if this binding is a uniform block
		for &b in e.module.bindings {
			if b.name == d.name && b.kind == .Uniform {
				// Will be accessed as binding_field, handled by field access
				write(&e.w, d.name)
				return
			}
		}
		write(&e.w, d.name)

	case ^IR_Input_Field:
		// HLSL entry points pass struct directly
		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, "(")
		emit_hlsl_expr(e, d.condition)
		write(&e.w, " ? ")
		emit_hlsl_expr(e, d.true_val)
		write(&e.w, " : ")
		emit_hlsl_expr(e, d.false_val)
		write(&e.w, ")")
	}
}

// -- Helpers --

@(private = "file")
ir_op_to_hlsl :: 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")
needs_hlsl_mul :: proc(left, right: ^IR_Expr) -> bool {
	if left == nil || right == nil { return false }
	return is_matrix(left.type) || is_matrix(right.type)
}

@(private = "file")
builtin_to_hlsl :: proc(name: string) -> string {
	switch name {
	case "sample":          return "Sample" // handled specially
	case "sample_level":    return "SampleLevel"
	case "mix":             return "lerp"
	case "fract":           return "frac"
	case "mod":             return "fmod"
	case "inversesqrt":     return "rsqrt"
	case "dfdx":            return "ddx"
	case "dfdy":            return "ddy"
	case "atan2":           return "atan2"
	case "sample_grad":     return "SampleGrad"
	case "sample_compare":  return "SampleCmpLevelZero"
	case "texel_fetch":     return "Load"
	case "texture_size":    return "GetDimensions"
	}
	return name
}

@(private = "file")
hlsl_input_semantic :: proc(io: IR_IO_Var, stage: Shader_Stage) -> string {
	if io.builtin != "" {
		return hlsl_builtin_semantic(io.builtin, true)
	}
	return fmt.aprintf("TEXCOORD%d", io.location)
}

@(private = "file")
hlsl_output_semantic :: proc(io: IR_IO_Var, stage: Shader_Stage) -> string {
	if io.builtin != "" {
		return hlsl_builtin_semantic(io.builtin, false)
	}
	if stage == .Fragment {
		return fmt.aprintf("SV_Target%d", io.location)
	}
	return fmt.aprintf("TEXCOORD%d", io.location)
}

@(private = "file")
hlsl_builtin_semantic :: proc(name: string, is_input: bool) -> string {
	switch name {
	case "position":              return "SV_Position"
	case "vertex_id":             return "SV_VertexID"
	case "instance_id":           return "SV_InstanceID"
	case "frag_coord":            return "SV_Position"
	case "front_facing":          return "SV_IsFrontFace"
	case "local_invocation_id":   return "SV_GroupThreadID"
	case "local_invocation_index": return "SV_GroupIndex"
	case "global_invocation_id":  return "SV_DispatchThreadID"
	case "workgroup_id":          return "SV_GroupID"
	}
	return name
}

resolved_type_to_hlsl :: 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.rows, v.cols)
	case Type_Struct_Resolved:
		return v.name
	case Type_Array_Resolved:
		elem_str := resolved_type_to_hlsl(v.elem)
		if v.size == 0 {
			return elem_str // unsized arrays handled differently in HLSL
		}
		return fmt.aprintf("%s[%d]", elem_str, v.size) // note: HLSL array syntax is type name[size]
	case Type_Sampler:
		switch v.kind {
		case .Sampler2D:       return "Texture2D"
		case .Sampler3D:       return "Texture3D"
		case .SamplerCube:     return "TextureCube"
		case .Sampler2DArray:  return "Texture2DArray"
		case .Sampler2DShadow: return "Texture2D" // with comparison sampler
		}
	case Type_Void:
		return "void"
	}
	return "void"
}

@(private = "file")
hlsl_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")
hlsl_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_hlsl(v.elem)
		if v.size == 0 {
			return elem_str, "[]"
		}
		return elem_str, fmt.aprintf("[%d]", v.size)
	}
	return resolved_type_to_hlsl(t), ""
}