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), "" }