sharkk/gpu public

branch main
showing the latest snapshot on main
ir.odin 13.8 KB · Plain text
package shader

// IR data structures for the Luma shader compiler.
// This is a structured IR that preserves control flow nesting (if/for/while)
// rather than using CFG basic blocks. This makes text backend emission trivial.
// Phase 3 can lower to CFG for SPIR-V.

IR_Module :: struct {
	functions:      [dynamic]IR_Function,
	structs:        [dynamic]IR_Struct,
	bindings:       [dynamic]IR_Binding,
	spec_constants: [dynamic]IR_Spec_Constant,
	shared_vars:    [dynamic]IR_Shared_Var,
}

IR_Address_Space :: enum {
	Private,
	Uniform,
	Storage,
	Workgroup,
	Input,
	Output,
	Push_Constant,
}

IR_Shared_Var :: struct {
	name: string,
	type: ^Resolved_Type,
}

IR_Var_Id :: distinct int

IR_Var_Decl :: struct {
	id:   IR_Var_Id,
	name: string,
	type: ^Resolved_Type,
}

Barrier_Kind :: enum {
	Workgroup,
}

IR_Spec_Constant :: struct {
	name:          string,
	spec_id:       int,
	type:          ^Resolved_Type,
	default_value: IR_Const_Value,
}

IR_Const_Value :: union {
	i64,
	f64,
	bool,
}

IR_Function :: struct {
	name:           string,
	params:         []IR_Param,
	return_type:    ^Resolved_Type,
	body:           [dynamic]IR_Stmt,
	is_entry:       bool,
	stage:          Shader_Stage,
	inputs:         [dynamic]IR_IO_Var, // flattened entry point inputs
	outputs:        [dynamic]IR_IO_Var, // flattened entry point outputs
	workgroup_size: [3]int,
	var_decls:      [dynamic]IR_Var_Decl, // all local variables, indexed by IR_Var_Id
}

IR_Param :: struct {
	name: string,
	type: ^Resolved_Type,
	id:   IR_Var_Id,
}

// Flattened entry point I/O variable
IR_IO_Var :: struct {
	name:     string,
	type:     ^Resolved_Type,
	location: int,    // -1 if builtin
	builtin:  string, // "" if location-based
}

IR_Struct :: struct {
	name:   string,
	fields: []IR_Struct_Field,
}

IR_Struct_Field :: struct {
	name:       string,
	type:       ^Resolved_Type,
	attributes: []Ast_Attribute,
}

IR_Binding_Kind :: enum {
	Uniform,
	Buffer,
	Texture,
	Sampler,
	Push_Constant,
}

IR_Binding :: struct {
	name:          string,
	type:          ^Resolved_Type,
	kind:          IR_Binding_Kind,
	group:         int,
	binding_num:   int,
	struct_ref:    ^Type_Struct_Resolved, // for uniform/buffer blocks
	combined_name: string,                // original name for split sampler pairs, "" otherwise
	address_space: IR_Address_Space,
}

// Statements — structured control flow preserved

IR_Discard :: struct {
	span: Source_Span,
}

IR_Break :: struct {
	span: Source_Span,
}

IR_Continue :: struct {
	span: Source_Span,
}

IR_Barrier :: struct {
	kind: Barrier_Kind,
	span: Source_Span,
}

IR_Stmt :: union {
	^IR_Let,
	^IR_Assign,
	^IR_Return,
	^IR_If,
	^IR_For,
	^IR_While,
	^IR_Store_Output,
	^IR_Expr_Stmt,
	^IR_Barrier,
	^IR_Discard,
	^IR_Break,
	^IR_Continue,
}

IR_Let :: struct {
	name:    string,
	id:      IR_Var_Id,
	value:   ^IR_Expr,
	type:    ^Resolved_Type,
	span:    Source_Span,
	mutable: bool,
}

IR_Assign :: struct {
	target: ^IR_Expr,
	value:  ^IR_Expr,
	span:   Source_Span,
}

IR_Return :: struct {
	value: ^IR_Expr, // nil for void return
	span:  Source_Span,
}

IR_If :: struct {
	condition:      ^IR_Expr,
	then_body:      [dynamic]IR_Stmt,
	elseif_clauses: []IR_Elseif,
	else_body:      [dynamic]IR_Stmt,
	span:           Source_Span,
}

IR_Elseif :: struct {
	condition: ^IR_Expr,
	body:      [dynamic]IR_Stmt,
}

IR_For :: struct {
	var_name: string,
	var_id:   IR_Var_Id,
	start:    ^IR_Expr,
	stop:     ^IR_Expr,
	step:     ^IR_Expr, // nil if default step of 1
	body:     [dynamic]IR_Stmt,
	span:     Source_Span,
}

IR_While :: struct {
	condition: ^IR_Expr,
	body:      [dynamic]IR_Stmt,
	span:      Source_Span,
}

IR_Store_Output :: struct {
	io_index: int,     // index into IR_Function.outputs
	value:    ^IR_Expr,
	span:     Source_Span,
}

IR_Expr_Stmt :: struct {
	expr: ^IR_Expr,
	span: Source_Span,
}

// Expressions

IR_Expr :: struct {
	kind:          IR_Expr_Kind,
	type:          ^Resolved_Type,
	derived:       IR_Expr_Derived,
}

IR_Expr_Kind :: enum {
	Literal,
	Var_Ref,
	Binary,
	Unary,
	Call,
	Field_Access,
	Swizzle,
	Index,
	Construct,
	Type_Cast,
	Load_Binding,
	Input_Field,
	Builtin_Var,
	Composite_Extract, // value-semantic field access (struct.field on SSA values)
	Vector_Shuffle,    // swizzle/reorder vector components by index
	Shared_Ref,        // reference to a shared (workgroup) variable
	Select,            // ternary conditional: select(cond, true, false)
}

IR_Expr_Derived :: union {
	^IR_Literal,
	^IR_Var_Ref,
	^IR_Binary,
	^IR_Unary,
	^IR_Call,
	^IR_Field_Access,
	^IR_Swizzle,
	^IR_Index,
	^IR_Construct,
	^IR_Type_Cast,
	^IR_Load_Binding,
	^IR_Input_Field,
	^IR_Builtin_Var,
	^IR_Composite_Extract,
	^IR_Vector_Shuffle,
	^IR_Shared_Ref,
	^IR_Select,
}

IR_Select :: struct {
	condition: ^IR_Expr,
	true_val:  ^IR_Expr,
	false_val: ^IR_Expr,
}

IR_Shared_Ref :: struct {
	name: string,
}

IR_Literal :: struct {
	value: IR_Literal_Value,
}

IR_Literal_Value :: union {
	i64,
	f64,
	bool,
}

IR_Var_Ref :: struct {
	name: string,
	id:   IR_Var_Id,
}

IR_Binary :: struct {
	op:    IR_Op,
	left:  ^IR_Expr,
	right: ^IR_Expr,
}

IR_Unary :: struct {
	op:      IR_Op,
	operand: ^IR_Expr,
}

IR_Call :: struct {
	name:       string,
	args:       []^IR_Expr,
	is_builtin: bool,
}

IR_Field_Access :: struct {
	object:     ^IR_Expr,
	field_name: string,
}

IR_Swizzle :: struct {
	object:     ^IR_Expr,
	components: string, // e.g. "xyz", "rgb"
}

IR_Index :: struct {
	object: ^IR_Expr,
	index:  ^IR_Expr,
}

IR_Construct :: struct {
	type_name: string,
	args:      []^IR_Expr,
}

IR_Type_Cast :: struct {
	value: ^IR_Expr,
}

IR_Load_Binding :: struct {
	name: string,
}

IR_Input_Field :: struct {
	param_name: string,
	field_name: string,
}

IR_Builtin_Var :: struct {
	name:     string,
	stage:    Shader_Stage,
	is_input: bool,
}

// Value-semantic field extraction from a composite (struct/vector) SSA value.
// Maps to SPIR-V OpCompositeExtract. Text backends emit as obj.field.
IR_Composite_Extract :: struct {
	object:     ^IR_Expr,
	index:      int,    // field index in the composite type
	field_name: string, // field name for text backend emission
}

// Swizzle/reorder vector components by integer indices.
// Maps to SPIR-V OpVectorShuffle.
IR_Vector_Shuffle :: struct {
	object:     ^IR_Expr,
	components: []int, // e.g. [0, 1, 2] for .xyz
}

IR_Op :: enum {
	// Arithmetic
	Add, Sub, Mul, Div, Mod,
	// Comparison
	Eq, Neq, Lt, Gt, Lte, Gte,
	// Logical
	And, Or,
	// Unary
	Neg, Not,
}

// -- Debug printer --

ir_to_string :: proc(module: ^IR_Module, allocator := context.allocator) -> string {
	w := writer_init(allocator)

	write_line(&w, "=== IR Module ===")
	write_line(&w, "")

	// Structs
	for s in module.structs {
		write_line(&w, "struct ", s.name)
		indent(&w)
		for f in s.fields {
			write_line(&w, f.name, ": ", type_to_string(f.type))
		}
		dedent(&w)
		write_line(&w, "end")
		write_line(&w, "")
	}

	// Spec constants
	for sc in module.spec_constants {
		write_line(&w, "@spec(", sc.spec_id, ") const ", sc.name, ": ", type_to_string(sc.type), " = ", ir_const_value_to_string(sc.default_value))
	}
	if len(module.spec_constants) > 0 {
		write_line(&w, "")
	}

	// Shared variables
	for sv in module.shared_vars {
		write_line(&w, "shared ", sv.name, ": ", type_to_string(sv.type))
	}
	if len(module.shared_vars) > 0 {
		write_line(&w, "")
	}

	// Bindings
	for b in module.bindings {
		kind_str: string
		switch b.kind {
		case .Uniform:       kind_str = "uniform"
		case .Buffer:        kind_str = "buffer"
		case .Texture:       kind_str = "texture"
		case .Sampler:       kind_str = "sampler"
		case .Push_Constant: kind_str = "push_constant"
		}
		if b.kind == .Push_Constant {
			write_line(&w, "@push_constant uniform ", b.name, ": ", type_to_string(b.type))
		} else {
			write_line(&w, "@group(", b.group, ") @binding(", b.binding_num, ") ", kind_str, " ", b.name, ": ", type_to_string(b.type))
		}
	}
	if len(module.bindings) > 0 {
		write_line(&w, "")
	}

	// Functions
	for &fn in module.functions {
		if fn.is_entry {
			write(&w, "@entry(")
			#partial switch fn.stage {
			case .Vertex:   write(&w, "vertex")
			case .Fragment: write(&w, "fragment")
			case .Compute:  write(&w, "compute")
			case:           write(&w, "unknown")
			}
			write(&w, ") ")
		}
		write(&w, "function ", fn.name, "(")
		for p, i in fn.params {
			if i > 0 do write(&w, ", ")
			write(&w, p.name, ": ", type_to_string(p.type))
		}
		write(&w, ") -> ", type_to_string(fn.return_type), "\n")

		if fn.is_entry {
			indent(&w)
			write_line(&w, "-- inputs:")
			for io, i in fn.inputs {
				if io.builtin != "" {
					write_line(&w, "  [", i, "] @builtin(", io.builtin, ") ", io.name, ": ", type_to_string(io.type))
				} else {
					write_line(&w, "  [", i, "] @location(", io.location, ") ", io.name, ": ", type_to_string(io.type))
				}
			}
			write_line(&w, "-- outputs:")
			for io, i in fn.outputs {
				if io.builtin != "" {
					write_line(&w, "  [", i, "] @builtin(", io.builtin, ") ", io.name, ": ", type_to_string(io.type))
				} else {
					write_line(&w, "  [", i, "] @location(", io.location, ") ", io.name, ": ", type_to_string(io.type))
				}
			}
			dedent(&w)
		}

		indent(&w)
		ir_print_stmts(&w, fn.body[:])
		dedent(&w)
		write_line(&w, "end")
		write_line(&w, "")
	}

	return writer_to_string(w)
}

@(private = "file")
ir_print_stmts :: proc(w: ^Writer, stmts: []IR_Stmt) {
	for stmt in stmts {
		ir_print_stmt(w, stmt)
	}
}

@(private = "file")
ir_print_stmt :: proc(w: ^Writer, stmt: IR_Stmt) {
	switch s in stmt {
	case ^IR_Let:
		for _ in 0 ..< w.indent do write(w, "\t")
		write(w, "let ", s.name, ": ", type_to_string(s.type), " = ")
		ir_print_expr(w, s.value)
		write(w, "\n")

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

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

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

	case ^IR_For:
		for _ in 0 ..< w.indent do write(w, "\t")
		write(w, "for ", s.var_name, " = ")
		ir_print_expr(w, s.start)
		write(w, ", ")
		ir_print_expr(w, s.stop)
		if s.step != nil {
			write(w, ", ")
			ir_print_expr(w, s.step)
		}
		write(w, " do\n")
		indent(w)
		ir_print_stmts(w, s.body[:])
		dedent(w)
		write_line(w, "end")

	case ^IR_While:
		for _ in 0 ..< w.indent do write(w, "\t")
		write(w, "while ")
		ir_print_expr(w, s.condition)
		write(w, " do\n")
		indent(w)
		ir_print_stmts(w, s.body[:])
		dedent(w)
		write_line(w, "end")

	case ^IR_Store_Output:
		for _ in 0 ..< w.indent do write(w, "\t")
		write(w, "store_output[", s.io_index, "] = ")
		ir_print_expr(w, s.value)
		write(w, "\n")

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

	case ^IR_Barrier:
		for _ in 0 ..< w.indent do write(w, "\t")
		write_line(w, "barrier()")

	case ^IR_Discard:
		for _ in 0 ..< w.indent do write(w, "\t")
		write_line(w, "discard")

	case ^IR_Break:
		for _ in 0 ..< w.indent do write(w, "\t")
		write_line(w, "break")

	case ^IR_Continue:
		for _ in 0 ..< w.indent do write(w, "\t")
		write_line(w, "continue")
	}
}

@(private = "file")
ir_print_expr :: proc(w: ^Writer, expr: ^IR_Expr) {
	if expr == nil {
		write(w, "<nil>")
		return
	}

	switch e in expr.derived {
	case ^IR_Literal:
		switch v in e.value {
		case i64:  write(w, v)
		case f64:  write(w, v)
		case bool: write(w, v ? "true" : "false")
		}

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

	case ^IR_Binary:
		write(w, "(")
		ir_print_expr(w, e.left)
		write(w, " ", ir_op_to_string(e.op), " ")
		ir_print_expr(w, e.right)
		write(w, ")")

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

	case ^IR_Call:
		if e.is_builtin do write(w, "@")
		write(w, e.name, "(")
		for arg, i in e.args {
			if i > 0 do write(w, ", ")
			ir_print_expr(w, arg)
		}
		write(w, ")")

	case ^IR_Field_Access:
		ir_print_expr(w, e.object)
		write(w, ".", e.field_name)

	case ^IR_Swizzle:
		ir_print_expr(w, e.object)
		write(w, ".", e.components)

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

	case ^IR_Construct:
		write(w, e.type_name, "(")
		for arg, i in e.args {
			if i > 0 do write(w, ", ")
			ir_print_expr(w, arg)
		}
		write(w, ")")

	case ^IR_Type_Cast:
		write(w, type_to_string(expr.type), "(")
		ir_print_expr(w, e.value)
		write(w, ")")

	case ^IR_Load_Binding:
		write(w, "binding:", e.name)

	case ^IR_Input_Field:
		write(w, "input:", e.param_name, ".", e.field_name)

	case ^IR_Builtin_Var:
		write(w, "builtin:", e.name)

	case ^IR_Composite_Extract:
		ir_print_expr(w, e.object)
		write(w, ".", e.field_name)

	case ^IR_Vector_Shuffle:
		ir_print_expr(w, e.object)
		write(w, ".", swizzle_indices_to_string(e.components))

	case ^IR_Shared_Ref:
		write(w, "shared:", e.name)

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

ir_op_to_string :: 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 "?"
}