Skip to content

Core IR (cheetah.internal.core_ir)

The core IR is a strongly-typed representation of values and control flow.

Types

  • PrimType(prim): primitive scalars (bool, i{8,16,32,64}, u{8,16,32,64}, bf16, f{16,32,64}, and vector pairs)
  • PtrType(mut, elem): pointer with mutability (const/mut/volatile), elem may be None for void*
  • TensorType(mut, elem): tensor view typed by primitive element

Type handles are uniqued (UniqueCache) and compared by identity.

Values and expressions

Note that Exprs - Literal scalar values: IntLit, FloatLit, BoolLit - Pointer ops: PtrOffsetExpr, PtrReadExpr, PtrWriteExpr, PtrCastExpr - Arithmetic/logical ops: BinOpExpr, UnOpExpr with strict type rules (check_binop, check_unop) - Casts: PrimCastExpr - Calls: FuncVoidPtrExpr (address-of), FuncCallExpr - Externs: ExternVarExpr for CUDA builtins (e.g., threadIdx.x) - Raw: RawExpr and RawStat with $name placeholder substitution - Inline PTX: RawAsmExpr with out/in operands and constraints

Blocks and control

  • Blocks have a role (ROOT, IF_THEN, IF_ELSE, FOR), instruction list, and optional return value
  • Structured building via Context.begin_if/finish_if, begin_for/finish_for ensures blocks are well-formed

Functions

  • Declared with argument types and optional return type (kernels must return void)
  • Defined by pushing a Body into the context, emitting instructions, and closing blocks
  • Internal names can be set; export names are assigned later during codegen

Validation

  • Centralized checks ensure type correctness, scope membership, and control structure matching
  • Epoch tracking (control/memory) feeds the inlining analysis in lowering