magpie-engineer

name	magpie-engineer
description	Comprehensive guide for writing Magpie (.mp) programs, debugging compiler errors, and using the Magpie CLI. Designed for agents who have never seen Magpie before.

Magpie Language Programming Guide

Use this skill whenever you write, debug, or review Magpie .mp programs.

---

0) Non-negotiable Rules

1. Binary-first: treat compiler diagnostics as ground truth.
   • Run magpie explain <CODE> for any error code you don't recognize.
   • Use --output json for machine-readable diagnostics.
2. Diagnose by error code first, then apply the smallest fix.
3. Smallest reproducer first -- one tiny .mp file, then grow.
4. Change one dimension at a time (syntax, then types, then ownership, then backend).
5. Prove fixes with commands and exit codes.

---

1) What is Magpie

Magpie is a compiled, SSA-based language with explicit ownership, ARC memory management, and first-class GPU support. Programs are written in a low-level IR-like textual format (CSNF).

Key characteristics:

• SSA form -- every value is defined once, used by name (%x, %y)
• Explicit ownership -- you manually borrow.shared, borrow.mut, share, clone.shared
• ARC-managed heap -- the compiler inserts reference counting; you manage ownership transitions
• Block-structured -- functions contain labeled basic blocks (bb0:, bb1:) with explicit terminators
• No implicit conversions -- types must match exactly; const suffixes must match declared types
• 5 GPU backends -- SPIR-V, Metal (MSL), PTX (CUDA), HIP (AMD), WGSL (WebGPU)

How Magpie differs from familiar languages:

• There are no expressions or statements -- only SSA instructions and terminators
• Control flow uses br, cbr, switch -- not if/else/for/while
• Variables are SSA locals (%name) -- assigned exactly once
• Functions are @name -- always prefixed with @
• Types are TName -- user-defined types start with T
• Comments use ; (not // or #); doc comments use ;;

---

2) Module Structure

Every .mp file must have a strict header in exactly this order:

module <module.path>
exports { <exported symbols> }
imports { <import groups> }
digest "<hex string>"

The order is mandatory. Swapping any header line causes a parse error (MPP0002).

Minimal template (copy-paste this to start any program)

module demo.main
exports { @main }
imports { }
digest "0000000000000000"

fn @main() -> i32 {
bb0:
  ret const.i32 0
}

Exports

List functions (@name) and types (TName) the module makes visible:

exports { @main, @helper, TPoint }

Imports

Group imports by source module:

imports { std.io::{@println}, util.math::{@sum, TVector} }

Use imported functions with their full qualified path in calls:

call_void std.io.@println { args=[%msg] }

Digest

A hex string for content hashing. Use "0000000000000000" as a placeholder during development.

---

3) Functions

Function kinds

fn @name(%param: Type) -> ReturnType { ... }
async fn @name(%param: Type) -> ReturnType { ... }
unsafe fn @name(%param: Type) -> ReturnType { ... }
gpu fn @name(%param: Type) -> unit target(msl) { ... }

Parameters and return types

fn @add(%a: i64, %b: i64) -> i64 {
bb0:
  %r: i64 = i.add { lhs=%a, rhs=%b }
  ret %r
}

Parameters are SSA names (%name: Type). Return type follows ->.

Meta blocks (optional)

fn @compute(%x: i64) -> i64 meta {
  uses { @helper, @util }
  effects { io, alloc }
  cost { time=5, space=2 }
} {
bb0:
  %r: i64 = call @helper { x=%x }
  ret %r
}

Basic blocks and terminators

Every function body contains one or more blocks. Each block has a label (bbN:) and ends with exactly one terminator.

bb0:              ; block label (bb followed by a number)
  <instructions>  ; zero or more SSA assignments or void ops
  <terminator>    ; exactly one: ret, br, cbr, switch, or unreachable

Terminators

Terminator	Syntax	Description
ret	ret %value or ret const.i32 0	Return a value
ret (void)	ret	Return unit (for -> unit functions)
br	br bb1	Unconditional branch
cbr	cbr %cond bb_true bb_false	Conditional branch (%cond must be bool)
switch	switch %val { case 0 -> bb1 case 1 -> bb2 } else bb3	Multi-way branch
unreachable	unreachable	Mark unreachable code path

Branching example

fn @abs(%x: i64) -> i64 {
bb0:
  %neg: bool = icmp.slt { lhs=%x, rhs=const.i64 0 }
  cbr %neg bb1 bb2

bb1:
  %zero: i64 = const.i64 0
  %r: i64 = i.sub { lhs=%zero, rhs=%x }
  ret %r

bb2:
  ret %x
}

Phi nodes (merging values from multiple predecessors)

fn @max(%a: i64, %b: i64) -> i64 {
bb0:
  %cond: bool = icmp.sgt { lhs=%a, rhs=%b }
  cbr %cond bb1 bb2

bb1:
  br bb3

bb2:
  br bb3

bb3:
  %result: i64 = phi i64 { [bb1:%a], [bb2:%b] }
  ret %result
}

phi selects a value based on which predecessor block was taken. Borrow values cannot appear in phi nodes (MPO0102).

---

4) Type System

4.1 Primitive types

Category	Types
Signed integers	i1 i8 i16 i32 i64 i128
Unsigned integers	u1 u8 u16 u32 u64 u128
Floats	f16 f32 f64 bf16
Boolean	bool
Unit	unit

4.2 Builtin heap types

Type	Description
Str	UTF-8 string (built-in hash, eq, ord -- no impl needed)
Array<T>	Dynamic array
Map<K, V>	Hash map (K must have hash + eq impls)
TOption<T>	Optional value -- variants: Some { v: T }, None { }
TResult<Ok, Err>	Result value -- variants: Ok { v: Ok }, Err { e: Err }
TStrBuilder	Mutable string builder
TMutex<T>	Mutual exclusion lock
TRwLock<T>	Reader-writer lock
TCell<T>	Interior mutability cell
TFuture<T>	Async future
TChannelSend<T>	Channel send endpoint
TChannelRecv<T>	Channel receive endpoint
TCallable<TSigName>	First-class callable (closure/function pointer)

4.3 Ownership modifiers

Place these before a type to change ownership mode:

Modifier	Meaning	Example
(none)	Unique owner (default)	%p: TPoint
shared	Reference-counted shared	%s: shared TPoint
borrow	Shared read-only borrow	%b: borrow TPoint
mutborrow	Exclusive mutable borrow	%m: mutborrow TPoint
weak	Weak reference (non-owning)	%w: weak TPoint

Important restrictions:

• TOption and TResult reject shared and weak prefixes (MPT0002, MPT0003)
• Borrows cannot cross block boundaries, appear in phi, or be returned from functions

4.4 User-defined types

Heap struct (reference-counted, can contain any types)

heap struct TPoint {
  field x: i64
  field y: i64
}

Value struct (stack-allocated, cannot contain heap handles)

value struct TColor {
  field r: u8
  field g: u8
  field b: u8
}

If a value struct field contains a heap type (like Str), you get MPT1005. Use heap struct instead.

Heap enum

heap enum TShape {
  variant Circle { field radius: f64 }
  variant Rect { field w: f64, field h: f64 }
}

Value enum (deferred in v0.1 -- use heap enum instead)

value enum will produce MPT1020. Use heap enum for now.

4.5 Signatures and callables

sig TBinaryOp(i64, i64) -> i64

Then use TCallable<TBinaryOp> as the type for first-class function values.

4.6 Raw pointers

rawptr<T>   ; e.g., rawptr<i64>, rawptr<u8>

Raw pointer ops require unsafe context.

4.7 Type hierarchy diagram

                          Types
           ┌───────────────┼───────────────┐
       Primitive        Builtin          User-defined
     i8..i128        Str, Array<T>      heap struct
     u8..u128        Map<K,V>           value struct
     f16 f32 f64     TOption<T>         heap enum
     bf16 bool       TResult<Ok,Err>    (value enum: v0.1 deferred)
     unit            TStrBuilder
                     TMutex<T> ...      rawptr<T>
                     TCallable<TSig>

4.8 Trait requirements for collections

Type	hash	eq	ord	Notes
All primitives	yes	yes	yes	Built-in
bool	yes	yes	yes	Built-in
Str	yes	yes	yes	Built-in -- no explicit impl needed
User heap struct	no	no	no	Must provide explicit impl declarations
User heap enum	no	no	no	Must provide explicit impl declarations

Map<Str, V> works with no extra code. Map<TMyStruct, V> requires:

fn @hash_my(%v: borrow TMyStruct) -> u64 { bb0: ret const.u64 0 }
impl hash for TMyStruct = @hash_my

fn @eq_my(%a: borrow TMyStruct, %b: borrow TMyStruct) -> bool { bb0: ret const.bool true }
impl eq for TMyStruct = @eq_my

Trait impl signature requirements:

• hash: (%v: borrow T) -> u64
• eq: (%a: borrow T, %b: borrow T) -> bool
• ord: (%a: borrow T, %b: borrow T) -> i32

4.9 Extern declarations (FFI)

extern "c" module libc {
  fn @malloc(%size: u64) -> rawptr<u8> attrs { returns="owned" }
  fn @free(%ptr: rawptr<u8>) -> unit
}

Extern functions returning rawptr must have attrs { returns="owned" } or attrs { returns="borrowed" } (MPF0001).

4.10 Globals

global @MAX_SIZE: i64 = const.i64 1024

---

5) Ownership & Borrowing Model

5.1 Ownership state machine

             ┌──────────────────────┐
             │       Unique         │ <── initial state after `new`
             │    (sole owner)      │
             └───┬────────┬─────┬──┘
                 │        │     │
        share{}  │        │     │ borrow.mut{}
    (consumes    │        │     │
     original)   │        │     v
                 v        │  ┌──────────────────┐
          ┌──────────┐    │  │    MutBorrow      │
          │  Shared  │    │  │ (exclusive r/w)   │
          │(ref-cnt) │    │  │ block-scoped only │
          └────┬─────┘    │  └──────────────────┘
               │          │
    clone      │          │ borrow.shared{}
    .shared{}  │          │
               v          v
          ┌──────────┐  ┌──────────────────┐
          │  Shared  │  │     Borrow       │
          │  (copy)  │  │  (shared r/o)    │
          └──────────┘  │  block-scoped    │
                        └──────────────────┘

5.2 Key rules

Rule	Consequence
share { v=%x }	Consumes %x. %x is dead after this.
borrow.mut { v=%x }	Creates exclusive mutable access, scoped to current block only
borrow.shared { v=%x }	Creates shared read access, scoped to current block only
clone.shared { v=%x }	Creates a new shared ref. Original survives.
Borrows cannot cross block boundaries	MPO0101
Borrows cannot appear in phi nodes	MPO0102
Borrows cannot be returned from functions	MPHIR03
Cannot use a value after it has been moved/consumed	MPO0007
Cannot move a value while a borrow is active	MPO0011

5.3 Borrow lifecycle per block

  Block entry
      │
      │  %x: TPoint = ...              (unique value)
      │
      │  %xb: borrow TPoint =          (borrow starts here ─┐
      │    borrow.shared { v=%x }                             │ valid scope
      │                                                       │
      │  getfield { obj=%xb, field=y }  (OK: same block)     │
      │                                                       │
      │  br bb1                         (block exit ──────────┘ borrow dies)
      │
      v
  bb1:
      │  %xb is DEAD here. Cannot reference it.
      │  To borrow again: %xb2 = borrow.shared { v=%x }

5.4 Correct multi-block borrow pattern

Wrong -- borrow crosses block boundary (MPO0101):

bb0:
  %pb: borrow TPoint = borrow.shared { v=%p }
  br bb1
bb1:
  %x: i64 = getfield { obj=%pb, field=x }   ; ERROR: %pb from bb0

Correct -- re-borrow in each block:

bb0:
  %pb: borrow TPoint = borrow.shared { v=%p }
  %x: i64 = getfield { obj=%pb, field=x }   ; OK: same block
  br bb1
bb1:
  %pb2: borrow TPoint = borrow.shared { v=%p }
  %y: i64 = getfield { obj=%pb2, field=y }  ; OK: new borrow in bb1
  ret %y

5.5 Mutation then read pattern (split across blocks)

To mutate a struct then read it, you must split into separate blocks:

bb0:
  %p: TPoint = new TPoint { x=const.i64 1, y=const.i64 2 }
  %pm: mutborrow TPoint = borrow.mut { v=%p }
  setfield { obj=%pm, field=y, val=const.i64 99 }
  br bb1                    ; end mutborrow scope

bb1:
  %pb: borrow TPoint = borrow.shared { v=%p }
  %y: i64 = getfield { obj=%pb, field=y }   ; reads 99
  ret %y

5.6 Mutation vs read receiver requirements

Operation	Required receiver
getfield	borrow T or mutborrow T
setfield	mutborrow T
arr.push, arr.set, arr.sort, arr.pop	mutborrow Array<T>
arr.len, arr.get, arr.slice, arr.contains, arr.map, arr.filter, arr.reduce	borrow Array<T> (or mutborrow)
map.set, map.delete, map.delete_void	mutborrow Map<K,V>
map.len, map.get, map.get_ref, map.contains_key, map.keys, map.values	borrow Map<K,V> (or mutborrow)
str.builder.append_*	mutborrow TStrBuilder (or unique)

5.7 Move/consume semantics

These operations consume (move) their input -- the value is dead afterward:

• share { v=%x } -- consumes %x
• new TFoo { field=%x } -- consumes each field value
• enum.new<V> { v=%x } -- consumes payload values
• callable.capture @fn { cap=%x } -- consumes captured values
• str.concat { a=%x, b=%y } -- consumes both strings
• str.builder.build { b=%sb } -- consumes the builder
• setfield { obj=%m, field=f, val=%x } -- consumes %x (the new value)
• arr.push { arr=%m, val=%x } -- consumes %x
• map.set { map=%m, key=%k, val=%v } -- consumes %k and %v
• ret %x -- consumes %x for move-only types

---

6) Complete Opcode Reference

Every instruction is either a value-producing op (appears in an SSA assignment) or a void op (standalone statement).

SSA assignment form: %name: Type = <value-op> Void op form: <void-op> (no assignment)

6.1 Constants

%x: i32  = const.i32 42
%y: i64  = const.i64 -1
%z: f64  = const.f64 3.14
%b: bool = const.bool true
%s: Str  = const.Str "hello"
%u: u32  = const.u32 7

The type suffix on const must exactly match the declared SSA type. const.i32 for i32, const.i64 for i64, etc. Mismatch causes MPT2014/MPT2015.

6.2 Integer arithmetic / bitwise

All take { lhs=V, rhs=V } and produce the same integer type.

Op	Description
i.add	Addition
i.sub	Subtraction
i.mul	Multiplication
i.sdiv	Signed division
i.udiv	Unsigned division
i.srem	Signed remainder
i.urem	Unsigned remainder
i.add.wrap	Wrapping add
i.sub.wrap	Wrapping sub
i.mul.wrap	Wrapping mul
i.add.checked	Checked add (traps on overflow)
i.sub.checked	Checked sub
i.mul.checked	Checked mul
i.and	Bitwise AND
i.or	Bitwise OR
i.xor	Bitwise XOR
i.shl	Shift left
i.lshr	Logical shift right
i.ashr	Arithmetic shift right

Example:

%r: i64 = i.add { lhs=%a, rhs=%b }
%c: i64 = i.mul.checked { lhs=%r, rhs=const.i64 2 }

6.3 Float arithmetic

All take { lhs=V, rhs=V } and produce the same float type.

Op	Description
f.add	Addition
f.sub	Subtraction
f.mul	Multiplication
f.div	Division
f.rem	Remainder
f.add.fast	Fast-math add
f.sub.fast	Fast-math sub
f.mul.fast	Fast-math mul
f.div.fast	Fast-math div

6.4 Comparisons

All take { lhs=V, rhs=V } and produce bool.

Integer comparisons (icmp.*):

Op	Meaning
icmp.eq	Equal
icmp.ne	Not equal
icmp.slt	Signed less than
icmp.sgt	Signed greater than
icmp.sle	Signed less or equal
icmp.sge	Signed greater or equal
icmp.ult	Unsigned less than
icmp.ugt	Unsigned greater than
icmp.ule	Unsigned less or equal
icmp.uge	Unsigned greater or equal

Float comparisons (fcmp.*):

Op	Meaning
fcmp.oeq	Ordered equal
fcmp.one	Ordered not equal
fcmp.olt	Ordered less than
fcmp.ogt	Ordered greater than
fcmp.ole	Ordered less or equal
fcmp.oge	Ordered greater or equal

6.5 Calls

Op	Syntax	Produces value?
call	call @fn { key=V, ... }	Yes
call_void	call_void @fn { key=V, ... }	No (void)
call with type args	call @fn<i64> { key=V, ... }	Yes
try	try @fn { key=V, ... }	Yes

Qualified calls to imported functions:

%r: i64 = call util.math.@sum { a=%x, b=%y }
call_void std.io.@println { args=[%msg] }

Note: call argument keys are used for readability but lowering preserves argument order. Keep key order stable.

6.6 Async / suspend

Op	Syntax	Produces value?
suspend.call	suspend.call @fn { key=V, ... }	Yes
suspend.await	suspend.await { fut=V }	Yes

async fn @fetch(%id: u64) -> Str meta { uses { @get_data } } {
bb0:
  %data: Str = suspend.call @get_data { id=%id }
  ret %data
}

6.7 Indirect calls (callables)

Op	Syntax	Produces value?
call.indirect	call.indirect %callable { key=V, ... }	Yes
call_void.indirect	call_void.indirect %callable { key=V, ... }	No (void)

6.8 Struct / enum ops

Op	Syntax	Produces value?
new	new TFoo { field1=V, field2=V }	Yes
getfield	getfield { obj=V, field=name }	Yes
setfield	setfield { obj=V, field=name, val=V }	No (void)
phi	phi Type { [bb0:V], [bb1:V] }	Yes
enum.new	enum.new<Variant> { key=V, ... }	Yes
enum.tag	enum.tag { v=V }	Yes (i32)
enum.is	enum.is<Variant> { v=V }	Yes (bool)
enum.payload	enum.payload<Variant> { v=V }	Yes

getfield requires a borrow/mutborrow receiver (MPHIR01). setfield requires a mutborrow receiver (MPHIR02). setfield uses val= key (not value=). Both getfield and setfield accept keys in any order.

6.9 Ownership conversion ops

Op	Syntax	Description
share	share { v=V }	Unique -> Shared (consumes input)
clone.shared	clone.shared { v=V }	Copy a shared ref (original survives)
clone.weak	clone.weak { v=V }	Copy a weak ref
weak.downgrade	weak.downgrade { v=V }	Shared -> Weak
weak.upgrade	weak.upgrade { v=V }	Weak -> TOption
borrow.shared	borrow.shared { v=V }	Create shared read-only borrow
borrow.mut	borrow.mut { v=V }	Create exclusive mutable borrow
cast	cast<FromPrim, ToPrim> { v=V }	Primitive-to-primitive cast

cast only works between primitive types (MPT2010/MPT2011):

%narrow: i32 = cast<i64, i32> { v=%wide }

6.10 Raw pointer ops (unsafe context required)

Op	Syntax	Produces value?
ptr.null	ptr.null<Type>	Yes
ptr.addr	ptr.addr<Type> { p=V }	Yes (u64)
ptr.from_addr	ptr.from_addr<Type> { addr=V }	Yes (rawptr)
ptr.add	ptr.add<Type> { p=V, count=V }	Yes (rawptr)
ptr.load	ptr.load<Type> { p=V }	Yes
ptr.store	ptr.store<Type> { p=V, v=V }	No (void)

Using any ptr.* outside unsafe fn or unsafe { } causes MPS0024.

6.11 Callable capture

%cb: TCallable<TSigName> = callable.capture @fn_name { cap1=%val1, cap2=%val2 }

Use { } for no captures. The captures are consumed (moved into the callable).

6.12 Array ops

Value-producing (read):

Op	Syntax	Receiver
arr.new<T>	arr.new<T> { cap=V }	n/a (creates new)
arr.len	arr.len { arr=V }	borrow/mutborrow
arr.get	arr.get { arr=V, idx=V }	borrow/mutborrow
arr.pop	arr.pop { arr=V }	mutborrow
arr.slice	arr.slice { arr=V, start=V, end=V }	borrow/mutborrow
arr.contains	arr.contains { arr=V, val=V }	borrow/mutborrow (elem needs eq)
arr.map	arr.map { arr=V, fn=V }	borrow/mutborrow
arr.filter	arr.filter { arr=V, fn=V }	borrow/mutborrow
arr.reduce	arr.reduce { arr=V, init=V, fn=V }	borrow/mutborrow

Void (mutation):

Op	Syntax	Receiver
arr.push	arr.push { arr=V, val=V }	mutborrow
arr.set	arr.set { arr=V, idx=V, val=V }	mutborrow
arr.sort	arr.sort { arr=V }	mutborrow (elem needs ord)
arr.foreach	arr.foreach { arr=V, fn=V }	borrow/mutborrow

6.13 Map ops

Value-producing (read):

Op	Syntax	Receiver
map.new<K, V>	map.new<K, V> { }	n/a (K needs hash + eq)
map.len	map.len { map=V }	borrow/mutborrow
map.get	map.get { map=V, key=V }	borrow/mutborrow (V must be Dupable, else MPO0103)
map.get_ref	map.get_ref { map=V, key=V }	borrow/mutborrow
map.delete	map.delete { map=V, key=V }	mutborrow
map.contains_key	map.contains_key { map=V, key=V }	borrow/mutborrow
map.keys	map.keys { map=V }	borrow/mutborrow
map.values	map.values { map=V }	borrow/mutborrow

Void (mutation):

Op	Syntax	Receiver
map.set	map.set { map=V, key=V, val=V }	mutborrow
map.delete_void	map.delete_void { map=V, key=V }	mutborrow

6.14 String / builder / JSON ops

Value-producing:

Op	Syntax	Notes
str.concat	str.concat { a=V, b=V }	Consumes both inputs
str.len	str.len { s=V }	s must be borrow/mutborrow
str.eq	str.eq { a=V, b=V }	Both must be borrow/mutborrow
str.slice	str.slice { s=V, start=V, end=V }	s must be borrow/mutborrow
str.bytes	str.bytes { s=V }	s must be borrow/mutborrow
str.builder.new	str.builder.new { }	Creates new builder
str.builder.build	str.builder.build { b=V }	Consumes builder, produces Str
str.parse_i64	str.parse_i64 { s=V }	Returns TResult<i64, Str>
str.parse_u64	str.parse_u64 { s=V }	Returns TResult<u64, Str>
str.parse_f64	str.parse_f64 { s=V }	Returns TResult<f64, Str>
str.parse_bool	str.parse_bool { s=V }	Returns TResult<bool, Str>
json.encode<T>	json.encode<T> { v=V }	Returns TResult<Str, Str>
json.decode<T>	json.decode<T> { s=V }	Returns TResult<rawptr<u8>, Str>

Void (builder append):

Op	Syntax
str.builder.append_str	str.builder.append_str { b=V, s=V }
str.builder.append_i64	str.builder.append_i64 { b=V, v=V }
str.builder.append_i32	str.builder.append_i32 { b=V, v=V }
str.builder.append_f64	str.builder.append_f64 { b=V, v=V }
str.builder.append_bool	str.builder.append_bool { b=V, v=V }

6.15 GPU ops

Value-producing:

Op	Syntax	Notes
gpu.thread_id	gpu.thread_id { dim=V }	dim: const.u32 0/1/2 (x/y/z)
gpu.workgroup_id	gpu.workgroup_id { dim=V }	dim: 0/1/2
gpu.workgroup_size	gpu.workgroup_size { dim=V }	dim: 0/1/2
gpu.global_id	gpu.global_id { dim=V }	dim: 0/1/2
gpu.buffer_load<T>	gpu.buffer_load<T> { buf=V, idx=V }	Typed load
gpu.buffer_len<T>	gpu.buffer_len<T> { buf=V }	Buffer length
gpu.shared<count, T>	gpu.shared<256, f32>	Shared memory
gpu.launch	gpu.launch { device=V, kernel=@fn, grid=V, block=V, args=V }	Strict key order
gpu.launch_async	gpu.launch_async { device=V, kernel=@fn, grid=V, block=V, args=V }	Strict key order

Void:

Op	Syntax
gpu.barrier	gpu.barrier
gpu.buffer_store<T>	gpu.buffer_store<T> { buf=V, idx=V, v=V }

6.16 Other void ops

Op	Syntax
panic	panic { msg=V }

6.17 Arg value grammar

Arguments in call/GPU forms can be:

1. Value ref: %x or const.i64 42
2. List: [%a, %b, @fn_ref]
3. Function ref: @fn or module.@fn

6.18 Key-ordering summary

Op family	Key order required?	Notes
getfield	No -- any order	Parser accepts keys in any order
setfield	No -- any order	Uses val= key (not value=)
gpu.launch	Yes -- strict	device, kernel, grid, block, args
gpu.launch_async	Yes -- strict	device, kernel, grid, block, args
All other ops	Stable recommended	Keys preserved by order

---

7) GPU Programming

7.1 Overview

Magpie supports 5 GPU backends:

Backend	Target	Output	Control flow
SPIR-V	spv	Binary SPIR-V module	Native CFG
Metal	msl	.metal source	Structurized
PTX	ptx	LLVM IR -> PTX via llc	Native
HIP	hip	LLVM IR -> HSACO via llc+lld	Native
WGSL	wgsl	.wgsl source	Structurized

7.2 GPU function syntax

gpu fn @kernel_name(%param: rawptr<f32>) -> unit target(msl) {
bb0:
  ret
}

With workgroup size and capability requirements:

gpu fn @vector_add(%a: rawptr<f32>, %b: rawptr<f32>, %out: rawptr<f32>) -> unit target(msl) workgroup(64, 1, 1) {
bb0:
  %gid: u32 = gpu.global_id { dim=const.u32 0 }
  %x: f32 = gpu.buffer_load<f32> { buf=%a, idx=%gid }
  %y: f32 = gpu.buffer_load<f32> { buf=%b, idx=%gid }
  %z: f32 = f.add { lhs=%x, rhs=%y }
  gpu.buffer_store<f32> { buf=%out, idx=%gid, v=%z }
  ret
}

Unsafe GPU function with capability requirements:

unsafe gpu fn @advanced_kernel(%buf: rawptr<i32>) -> unit target(ptx) workgroup(256, 1, 1) requires(device_malloc) {
bb0:
  ret
}

7.3 GPU kernel restrictions

GPU kernels cannot use:

Restriction	Diagnostic
new (heap allocation)	MPG_CORE_1100
Str type	MPG_CORE_1104
Array<T>	MPG_CORE_1105
Map<K,V>	MPG_CORE_1106
TCallable (dynamic dispatch)	MPG_CORE_1102, MPG_CORE_1107
Recursion	MPG_CORE_1103

Use rawptr<T> for buffer parameters and primitive types for computation.

7.4 Launching kernels from host code

%result: unit = gpu.launch {
  device=%dev,
  kernel=@vector_add,
  grid=[const.u32 256, const.u32 1, const.u32 1],
  block=[const.u32 64, const.u32 1, const.u32 1],
  args=[%buf_a, %buf_b, %buf_out]
}

Grid and block dimensions are 3D: [x, y, z]. Key order is strict: device, kernel, grid, block, args.

7.5 Building GPU programs

magpie --entry src/main.mp --emit exe,msl build      # Metal
magpie --entry src/main.mp --emit exe,spv build      # Vulkan/SPIR-V
magpie --entry src/main.mp --emit exe,ptx build      # NVIDIA
magpie --entry src/main.mp --emit exe,hip build      # AMD
magpie --entry src/main.mp --emit exe,wgsl build     # WebGPU

7.6 Manifest GPU config

In Magpie.toml:

[gpu]
enabled = true
backend = "metal"
device_index = 0

---

8) Complete Working Examples

8.1 Minimal hello (i32 return)

module demo.main
exports { @main }
imports { }
digest "0000000000000000"

fn @main() -> i32 {
bb0:
  ret const.i32 0
}

8.2 Arithmetic

module demo.math
exports { @compute }
imports { }
digest "0000000000000000"

fn @compute(%a: i64, %b: i64) -> i64 {
bb0:
  %sum: i64 = i.add { lhs=%a, rhs=%b }
  %doubled: i64 = i.mul { lhs=%sum, rhs=const.i64 2 }
  %checked: i64 = i.add.checked { lhs=%doubled, rhs=const.i64 1 }
  ret %checked
}

8.3 Struct create + field access (borrow pattern)

module demo.structs
exports { @get_y }
imports { }
digest "0000000000000000"

heap struct TPoint {
  field x: i64
  field y: i64
}

fn @get_y() -> i64 {
bb0:
  %p: TPoint = new TPoint { x=const.i64 10, y=const.i64 20 }
  %pb: borrow TPoint = borrow.shared { v=%p }
  %y: i64 = getfield { obj=%pb, field=y }
  ret %y
}

8.4 Struct mutation (mutborrow, split blocks)

module demo.mutation
exports { @mutate_point }
imports { }
digest "0000000000000000"

heap struct TPoint {
  field x: i64
  field y: i64
}

fn @mutate_point() -> i64 {
bb0:
  %p: TPoint = new TPoint { x=const.i64 1, y=const.i64 2 }
  %pm: mutborrow TPoint = borrow.mut { v=%p }
  setfield { obj=%pm, field=y, val=const.i64 99 }
  br bb1

bb1:
  %pb: borrow TPoint = borrow.shared { v=%p }
  %y: i64 = getfield { obj=%pb, field=y }
  ret %y
}

8.5 TOption pattern

module demo.option
exports { @maybe }
imports { }
digest "0000000000000000"

fn @maybe(%flag: bool) -> i64 {
bb0:
  cbr %flag bb1 bb2

bb1:
  %some: TOption<i64> = enum.new<Some> { v=const.i64 42 }
  br bb3

bb2:
  %none: TOption<i64> = enum.new<None> { }
  br bb3

bb3:
  %opt: TOption<i64> = phi TOption<i64> { [bb1:%some], [bb2:%none] }
  %is_some: bool = enum.is<Some> { v=%opt }
  cbr %is_some bb4 bb5

bb4:
  %val: i64 = enum.payload<Some> { v=%opt }
  ret %val

bb5:
  ret const.i64 0
}

8.6 TResult pattern

module demo.result
exports { @safe_divide }
imports { }
digest "0000000000000000"

fn @safe_divide(%a: i32, %b: i32) -> TResult<i32, Str> {
bb0:
  %is_zero: bool = icmp.eq { lhs=%b, rhs=const.i32 0 }
  cbr %is_zero bb1 bb2

bb1:
  %msg: Str = const.Str "division by zero"
  %err: TResult<i32, Str> = enum.new<Err> { e=%msg }
  ret %err

bb2:
  %result: i32 = i.sdiv { lhs=%a, rhs=%b }
  %ok: TResult<i32, Str> = enum.new<Ok> { v=%result }
  ret %ok
}

8.7 Heap enum with tag dispatch

module demo.enums
exports { @classify }
imports { }
digest "0000000000000000"

heap enum TShape {
  variant Circle { field radius: f64 }
  variant Rect { field w: f64, field h: f64 }
}

fn @classify(%s: borrow TShape) -> i32 {
bb0:
  %tag: i32 = enum.tag { v=%s }
  %is_circle: bool = icmp.eq { lhs=%tag, rhs=const.i32 0 }
  cbr %is_circle bb1 bb2

bb1:
  ret const.i32 1

bb2:
  ret const.i32 2
}

8.8 Array push/len (mutborrow then borrow across blocks)

module demo.arrays
exports { @arr_demo }
imports { }
digest "0000000000000000"

fn @arr_demo() -> i64 {
bb0:
  %arr: Array<i64> = arr.new<i64> { cap=const.i64 8 }
  %arrm: mutborrow Array<i64> = borrow.mut { v=%arr }
  arr.push { arr=%arrm, val=const.i64 10 }
  arr.push { arr=%arrm, val=const.i64 20 }
  arr.push { arr=%arrm, val=const.i64 30 }
  br bb1

bb1:
  %arrb: borrow Array<i64> = borrow.shared { v=%arr }
  %len: i64 = arr.len { arr=%arrb }
  ret %len
}

8.9 Map with Str keys (no impl needed)

module demo.maps
exports { @map_demo }
imports { }
digest "0000000000000000"

fn @map_demo() -> i64 {
bb0:
  %m: Map<Str, i64> = map.new<Str, i64> { }
  %mb: mutborrow Map<Str, i64> = borrow.mut { v=%m }
  map.set { map=%mb, key=const.Str "alice", val=const.i64 100 }
  map.set { map=%mb, key=const.Str "bob", val=const.i64 200 }
  br bb1

bb1:
  %mr: borrow Map<Str, i64> = borrow.shared { v=%m }
  %len: i64 = map.len { map=%mr }
  ret %len
}

8.10 String ops and builder

module demo.strings
exports { @build_greeting }
imports { }
digest "0000000000000000"

fn @build_greeting() -> Str {
bb0:
  %sb: TStrBuilder = str.builder.new { }
  %hello: Str = const.Str "Hello, "
  %hb: borrow Str = borrow.shared { v=%hello }
  str.builder.append_str { b=%sb, s=%hb }
  str.builder.append_i32 { b=%sb, v=const.i32 42 }
  %result: Str = str.builder.build { b=%sb }
  ret %result
}

8.11 String parsing (TResult)

module demo.parse
exports { @parse_demo }
imports { }
digest "0000000000000000"

fn @parse_demo() -> i64 {
bb0:
  %s: Str = const.Str "123"
  %sb: borrow Str = borrow.shared { v=%s }
  %parsed: TResult<i64, Str> = str.parse_i64 { s=%sb }
  %is_ok: bool = enum.is<Ok> { v=%parsed }
  cbr %is_ok bb1 bb2

bb1:
  %val: i64 = enum.payload<Ok> { v=%parsed }
  ret %val

bb2:
  ret const.i64 -1
}

8.12 TCallable capture + indirect call

module demo.callable
exports { @callable_demo }
imports { }
digest "0000000000000000"

sig TUnaryOp(i64) -> i64

fn @double(%x: i64) -> i64 {
bb0:
  %r: i64 = i.mul { lhs=%x, rhs=const.i64 2 }
  ret %r
}

fn @callable_demo() -> i64 {
bb0:
  %cb: TCallable<TUnaryOp> = callable.capture @double { }
  %result: i64 = call.indirect %cb { x=const.i64 21 }
  ret %result
}

8.13 Ownership: share, clone.shared

module demo.sharing
exports { @share_demo }
imports { }
digest "0000000000000000"

heap struct TPerson {
  field name: Str
  field age: i32
}

fn @share_demo() -> i32 {
bb0:
  %name: Str = const.Str "Alice"
  %p: TPerson = new TPerson { name=%name, age=const.i32 30 }
  ; share consumes %p -- %p is dead after this
  %sp: shared TPerson = share { v=%p }
  ; clone.shared copies the shared ref -- %sp survives
  %sp2: shared TPerson = clone.shared { v=%sp }
  %b: borrow TPerson = borrow.shared { v=%sp }
  %age: i32 = getfield { obj=%b, field=age }
  ret %age
}

8.14 Async function with suspend.call

module demo.async_example
exports { @fetch_data }
imports { }
digest "0000000000000000"

fn @get_connection() -> Str {
bb0:
  %s: Str = const.Str "conn:1"
  ret %s
}

fn @process(%conn: Str, %id: u64) -> Str {
bb0:
  ret %conn
}

async fn @fetch_data(%id: u64) -> Str meta { uses { @get_connection, @process } } {
bb0:
  %conn: Str = suspend.call @get_connection { }
  %result: Str = call @process { conn=%conn, id=%id }
  ret %result
}

8.15 GPU kernel with buffer ops

module demo.gpu
exports { @vector_add }
imports { }
digest "0000000000000000"

gpu fn @vector_add(%a: rawptr<f32>, %b: rawptr<f32>, %out: rawptr<f32>) -> unit target(msl) workgroup(64, 1, 1) {
bb0:
  %gid: u32 = gpu.global_id { dim=const.u32 0 }
  %x: f32 = gpu.buffer_load<f32> { buf=%a, idx=%gid }
  %y: f32 = gpu.buffer_load<f32> { buf=%b, idx=%gid }
  %sum: f32 = f.add { lhs=%x, rhs=%y }
  gpu.buffer_store<f32> { buf=%out, idx=%gid, v=%sum }
  ret
}

8.16 Unsafe raw pointer block

module demo.unsafe_demo
exports { @raw_demo }
imports { }
digest "0000000000000000"

unsafe fn @raw_demo() -> i64 {
bb0:
  %p: rawptr<i64> = ptr.null<i64>
  %addr: u64 = ptr.addr<i64> { p=%p }
  ret const.i64 0
}

Or inline unsafe block in a normal function:

fn @inline_unsafe() -> i64 {
bb0:
  unsafe {
    %p: rawptr<i64> = ptr.null<i64>
  }
  ret const.i64 0
}

8.17 Trait impl for Map key (user struct)

module demo.traits
exports { @trait_demo }
imports { }
digest "0000000000000000"

heap struct TKey {
  field id: i64
}

fn @hash_key(%k: borrow TKey) -> u64 {
bb0:
  %kb: borrow TKey = borrow.shared { v=%k }
  %id: i64 = getfield { obj=%kb, field=id }
  %h: u64 = cast<i64, u64> { v=%id }
  ret %h
}
impl hash for TKey = @hash_key

fn @eq_key(%a: borrow TKey, %b: borrow TKey) -> bool {
bb0:
  %ab: borrow TKey = borrow.shared { v=%a }
  %bb: borrow TKey = borrow.shared { v=%b }
  %aid: i64 = getfield { obj=%ab, field=id }
  %bid: i64 = getfield { obj=%bb, field=id }
  %eq: bool = icmp.eq { lhs=%aid, rhs=%bid }
  ret %eq
}
impl eq for TKey = @eq_key

fn @trait_demo() -> i64 {
bb0:
  %m: Map<TKey, i64> = map.new<TKey, i64> { }
  %mb: mutborrow Map<TKey, i64> = borrow.mut { v=%m }
  %k: TKey = new TKey { id=const.i64 1 }
  map.set { map=%mb, key=%k, val=const.i64 42 }
  br bb1

bb1:
  %mr: borrow Map<TKey, i64> = borrow.shared { v=%m }
  %len: i64 = map.len { map=%mr }
  ret %len
}

---

9) Error Code Reference

9.1 Diagnostic triage flowchart

Got an error code?
       │
       v
  ┌───────────┐
  │ Code      │
  │ prefix?   │
  └─────┬─────┘
        │
   ┌────┼────┬────┬────┬────┐
   │    │    │    │    │    │
  MPP  MPS  MPT  MPHIR MPO  MPG/MPL
   │    │    │    │    │    │
   v    v    v    v    v    v
 Parse Resolve Type HIR  Own  GPU/
 error SSA/  mis-  bor- ship  Link/
 head  dom   match row  mode  Lint
 order unsafe trait viol move
 comma  ctx  impl  ation issue

9.2 Complete error code table

Parse (MPP*)

Code	Description
MPP0001	Source I/O / lexer read issue
MPP0002	Syntax error (missing comma, wrong header order, etc.)
MPP0003	Artifact emission failure

Resolve / SSA (MPS*)

Code	Description
MPS0000	Generic module resolution failure
MPS0001	Duplicate definition (module path or SSA local)
MPS0002	Unresolved reference / use-before-def
MPS0003	Dominance violation
MPS0004	Import/local namespace conflict
MPS0005	Type import/local type conflict
MPS0006	Ambiguous import name
MPS0008	Invalid CFG target / phi type legality
MPS0009	Duplicate block label
MPS0010	Structural/type invariant failure
MPS0011	Duplicate SSA local in lowering
MPS0012	Call arity mismatch
MPS0013	Expected single arg value
MPS0014	Invalid fn ref in scalar position
MPS0015	Invalid fn ref inside list
MPS0016	Invalid plain-value argument uses fn ref
MPS0017	Invalid branch predicate type (must be bool)
MPS0020	Duplicate function/global symbol
MPS0021	Duplicate type symbol
MPS0022	Duplicate @ namespace symbol
MPS0023	Duplicate sig symbol
MPS0024	ptr.* outside unsafe context
MPS0025	Unsafe fn call outside unsafe context

FFI (MPF*)

Code	Description
MPF0001	Extern rawptr return missing attrs { returns="owned" }

Type (MPT*)

Code	Description
MPT0001	Unknown primitive type
MPT0002	shared/weak invalid on TOption
MPT0003	shared/weak invalid on TResult
MPT1005	Value struct contains heap handle
MPT1020	Value enum deferred in v0.1
MPT1021	Aggregate type deferred in v0.1
MPT1023	Missing required trait impl (hash/eq/ord)
MPT1030	suspend.call on non-function target in v0.1
MPT1200	Orphan impl
MPT2001	Call arity mismatch
MPT2002	Call arg unknown type
MPT2003	Call arg type mismatch
MPT2004	Call target not found
MPT2005	Invalid generic type argument
MPT2006	getfield object unknown type
MPT2007	getfield requires borrow/mutborrow struct
MPT2008	getfield target not a struct
MPT2009	Missing struct field
MPT2010	cast operand unknown
MPT2011	cast only primitive-to-primitive
MPT2012	Numeric lhs unknown type
MPT2013	Numeric rhs unknown type
MPT2014	Numeric operands type mismatch
MPT2015	Wrong primitive family for numeric op
MPT2016	Duplicate field in constructor
MPT2017	Unknown field in constructor
MPT2018	Field value unknown type
MPT2019	Field type mismatch in constructor
MPT2020	Missing required field in constructor
MPT2021	new target must be struct
MPT2022	Unknown struct target
MPT2023	Invalid variant for TOption
MPT2024	Invalid variant for TResult
MPT2025	enum.new result type not enum
MPT2026	User enum variant not found
MPT2027	enum.new target must be enum
MPT2028	Trait impl parameter count mismatch
MPT2029	Trait impl return type mismatch
MPT2030	Trait impl first param must be borrow T
MPT2031	Trait impl params must both match borrow T
MPT2032	Impl target function missing
MPT2033	Parse/JSON result type shape mismatch
MPT2034	Parse/JSON input must be Str/borrow Str
MPT2035	json.encode<T> value type mismatch

HIR verify (MPHIR*)

Code	Description
MPHIR01	getfield object must be borrow/mutborrow
MPHIR02	setfield object must be mutborrow
MPHIR03	Borrow escapes via return

Ownership (MPO*)

Code	Description
MPO0003	Borrow escapes scope
MPO0004	Wrong ownership mode for mut/read op
MPO0007	Use after move
MPO0011	Move while borrowed
MPO0101	Borrow crosses block boundary
MPO0102	Borrow in phi
MPO0103	map.get requires Dupable V
MPO0201	Spawn/send capture rule violation

MPIR / Link / Budget (MPM*, MPLINK*, MPL*)

Code	Description
MPM0001	MPIR lowering produced no modules
MPLINK01	Primary link path failed
MPLINK02	Fallback link unavailable
MPL0001	Unknown emit kind
MPL0002	Requested artifact missing
MPL0801	LLM budget too small
MPL0802	Tokenizer fallback

Lint (MPL2*)

Code	Description
MPL2001	Oversized function body
MPL2002	Unused/dead symbol
MPL2003	Unnecessary borrow
MPL2005	Empty block
MPL2007	Unreachable code
MPL2020	Monomorphization pressure too high
MPL2021	Mixed generics mode conflict

GPU (MPG*)

Code	Description
MPG_CORE_1100	new heap allocation forbidden in kernel
MPG_CORE_1101	gpu.TBuffer (Arc) forbidden in kernel
MPG_CORE_1102	Dynamic dispatch / TCallable forbidden
MPG_CORE_1103	Recursion forbidden in kernel
MPG_CORE_1104	Str type forbidden in kernel
MPG_CORE_1105	Array<T> forbidden in kernel
MPG_CORE_1106	Map<K,V> forbidden in kernel
MPG_CORE_1107	TCallable parameter forbidden in kernel
MPG_CORE_1200	unsafe gpu fn missing requires(...)
MPG_CORE_1201	Invalid capability in requires(...)
MPG_CORE_1301	PTX/HIP toolchain not found
MPG_TYP_*	GPU type errors (e.g. bf16 unsupported on backend)
MPG_KRN_*	Kernel validation errors
MPG_BUF_*	GPU buffer errors
MPG_SYN_*	GPU sync errors (barrier misuse)
MPG_CAP_*	GPU capability errors
MPG_LNK_*	GPU link errors (llc/lld not found)
MPG_PRF_*	GPU profiling errors

9.3 Fix playbooks (most common errors)

MPHIR01 -- getfield object must be borrow/mutborrow

Bad:

%p: TPoint = new TPoint { x=const.i64 1, y=const.i64 2 }
%x: i64 = getfield { obj=%p, field=x }

Fix:

%p: TPoint = new TPoint { x=const.i64 1, y=const.i64 2 }
%pb: borrow TPoint = borrow.shared { v=%p }
%x: i64 = getfield { obj=%pb, field=x }

MPHIR02 -- setfield object must be mutborrow

Bad:

%pb: borrow TPoint = borrow.shared { v=%p }
setfield { obj=%pb, field=x, val=const.i64 3 }

Fix:

%pm: mutborrow TPoint = borrow.mut { v=%p }
setfield { obj=%pm, field=x, val=const.i64 3 }

MPHIR03 -- borrow escapes via return

Bad:

fn @leak(%p: TPoint) -> borrow TPoint {
bb0:
  %pb: borrow TPoint = borrow.shared { v=%p }
  ret %pb
}

Fix (return the field value instead):

fn @ok(%p: TPoint) -> i64 {
bb0:
  %pb: borrow TPoint = borrow.shared { v=%p }
  %x: i64 = getfield { obj=%pb, field=x }
  ret %x
}

MPO0007 -- use after move

Bad:

%sp: shared TPoint = share { v=%p }
%pb: borrow TPoint = borrow.shared { v=%p }  ; ERROR: %p consumed by share

Fix:

%sp: shared TPoint = share { v=%p }
%cp: shared TPoint = clone.shared { v=%sp }
; use %sp or %cp, not %p

MPO0011 -- move while borrowed

Bad:

%pb: borrow TPoint = borrow.shared { v=%p }
%sp: shared TPoint = share { v=%p }   ; ERROR: %p still borrowed

Fix (finish borrow, branch, then move):

bb0:
  %pb: borrow TPoint = borrow.shared { v=%p }
  %x: i64 = getfield { obj=%pb, field=x }
  br bb1
bb1:
  %sp: shared TPoint = share { v=%p }  ; OK: borrow ended at bb0 exit

MPO0101 -- borrow crosses block boundary

Bad:

bb0:
  %pb: borrow TPoint = borrow.shared { v=%p }
  br bb1
bb1:
  %x: i64 = getfield { obj=%pb, field=x }   ; ERROR: %pb from bb0

Fix:

bb0:
  br bb1
bb1:
  %pb: borrow TPoint = borrow.shared { v=%p }
  %x: i64 = getfield { obj=%pb, field=x }   ; OK: %pb defined in bb1

MPO0102 -- borrow in phi

Bad:

%pb: borrow TPoint = phi borrow TPoint { [bb1:%p1b], [bb2:%p2b] }

Fix (phi the owned value, then borrow locally):

%p: TPoint = phi TPoint { [bb1:%p1], [bb2:%p2] }
%pb: borrow TPoint = borrow.shared { v=%p }

MPT2014 -- numeric operands type mismatch

Bad:

%r: i64 = i.add { lhs=const.i64 1, rhs=const.i32 2 }

Fix:

%r: i64 = i.add { lhs=const.i64 1, rhs=const.i64 2 }

MPT0002 -- shared/weak invalid on TOption

Bad:

%x: shared TOption<i64> = enum.new<None> { }

Fix:

%x: TOption<i64> = enum.new<None> { }

MPT1023 -- missing required trait impl

Bad (Map<TKey, V> without hash/eq):

%m: Map<TKey, i64> = map.new<TKey, i64> { }

Fix:

fn @hash_key(%k: borrow TKey) -> u64 { bb0: ret const.u64 0 }
impl hash for TKey = @hash_key

fn @eq_key(%a: borrow TKey, %b: borrow TKey) -> bool { bb0: ret const.bool true }
impl eq for TKey = @eq_key

%m: Map<TKey, i64> = map.new<TKey, i64> { }

MPT1005 -- value struct contains heap handle

Bad:

value struct TBad { field s: Str }

Fix:

heap struct TGood { field s: Str }

MPS0024 -- ptr.* outside unsafe context

Bad:

%p: rawptr<i64> = ptr.null<i64>

Fix:

unsafe {
  %p: rawptr<i64> = ptr.null<i64>
}

MPS0003 -- dominance violation

Bad (%x only defined in bb1, used in bb3 which also has bb2 as predecessor):

bb0:
  cbr %c bb1 bb2
bb1:
  %x: i64 = const.i64 1
  br bb3
bb2:
  br bb3
bb3:
  %y: i64 = i.add { lhs=%x, rhs=const.i64 1 }   ; ERROR

Fix (use phi):

bb0:
  cbr %c bb1 bb2
bb1:
  %x1: i64 = const.i64 1
  br bb3
bb2:
  %x2: i64 = const.i64 2
  br bb3
bb3:
  %x: i64 = phi i64 { [bb1:%x1], [bb2:%x2] }
  %y: i64 = i.add { lhs=%x, rhs=const.i64 1 }

MPP0002 -- syntax error (missing comma)

Bad:

%x: i64 = i.add { lhs=const.i64 1 rhs=const.i64 2 }

Fix:

%x: i64 = i.add { lhs=const.i64 1, rhs=const.i64 2 }

---

10) CLI Reference

10.1 Global flags (place BEFORE the subcommand)

magpie [global-flags] <subcommand> [subcommand-args]

Flag	Description	Default
--entry <path>	Entry source file	From Magpie.toml
--emit <kinds>	Artifact types (comma-separated)	—
--output <text|json|jsonl>	Output format	text
--color <auto|always|never>	Color mode	auto
--log-level <level>	error|warn|info|debug|trace	warn
--profile <dev|release>	Build profile	dev
--target <triple>	Target triple	Host
-j, --jobs <n>	Parallel jobs	—
--max-errors <n>	Max errors per pass	20
--llm	LLM-optimized output (implies --output json)	—
--llm-token-budget <n>	Token budget for LLM output	—
--shared-generics	Use vtable-based shared generics	—
--no-auto-fmt	Disable auto-format in LLM mode	—

10.2 Subcommands

Command	Description	Example
build	Compile the project	magpie --entry src/main.mp --emit exe build
run	Build and execute	magpie --entry src/main.mp run -- arg1 arg2
parse	Parse and emit AST	magpie --entry src/main.mp parse
fmt	Format source (CSNF)	magpie fmt
lint	Run linter	magpie --entry src/main.mp lint
test	Run tests	magpie test --filter pattern
explain <CODE>	Explain a diagnostic code	magpie explain MPO0007
new <name>	Create a new project	magpie new my_project
doc	Generate documentation	magpie doc
repl	Start REPL	magpie repl
mpir verify	Verify MPIR correctness	magpie --entry src/main.mp mpir verify
graph symbols	Symbol graph	magpie --entry src/main.mp graph symbols
graph deps	Dependency graph	magpie --entry src/main.mp graph deps
graph ownership	Ownership graph	magpie --entry src/main.mp graph ownership
graph cfg	Control flow graph	magpie --entry src/main.mp graph cfg
ffi import	Import C headers	magpie ffi import --header foo.h --out ffi.mp

10.3 Emit kinds

Kind	Description
exe	Native executable
llvm-ir	LLVM IR text
llvm-bc	LLVM bitcode
object	Object file
asm	Assembly
shared-lib	Shared library
mpir	Magpie IR
mpd	MPD debug info
mpdbg	Debug info
ast	AST dump
spv	SPIR-V (Vulkan)
msl	Metal Shading Language
ptx	PTX (NVIDIA)
hip	HIP (AMD)
wgsl	WGSL (WebGPU)
symgraph	Symbol graph
depsgraph	Dependency graph
ownershipgraph	Ownership graph
cfggraph	Control flow graph

10.4 Common workflows

; Build and check for errors (machine-readable)
magpie --entry src/main.mp --output json build

; Build with multiple artifacts
magpie --entry src/main.mp --emit exe,llvm-ir,mpir build

; Run the program
magpie --entry src/main.mp run

; Get help for a specific error
magpie explain MPO0101

; Format the source
magpie --entry src/main.mp fmt

; Build GPU program for Metal
magpie --entry src/main.mp --emit exe,msl build

---

11) Manifest Format (Magpie.toml)

[package]
name = "my_project"
version = "0.1.0"
edition = "2024"

[build]
entry = "src/main.mp"
profile_default = "dev"

[dependencies]
some_lib = "1.2.3"
util = { path = "../util", version = "1.2.3" }

[gpu]
enabled = true
backend = "metal"
device_index = 0

[llm]
mode_default = false
token_budget = 32000

---

12) Common Pitfalls (Top 15)

1. Wrong const suffix -- const.i32 0 for i32, const.i64 0 for i64. They must match exactly.

2. Comment syntax -- use ; for line comments, ;; for doc comments. Not // or #.

3. Borrow crossing blocks -- borrows are block-scoped. Re-borrow in each new block.

4. getfield without borrow -- you must borrow.shared or borrow.mut before field access.

5. setfield uses val= -- not value=. Write setfield { obj=%m, field=f, val=%x }.

6. Missing comma in ops -- { lhs=x rhs=y } is wrong; use { lhs=x, rhs=y }.

7. Header order -- must be exactly: module, exports, imports, digest.

8. Using value after move -- share { v=%x } consumes %x. Don't use %x after.

9. Borrow in phi -- phi borrow T { ... } is illegal. Phi the owned value, then borrow.

10. Map<TUserStruct, V> without impls -- user types need explicit impl hash and impl eq. (Str doesn't.)

11. gpu.launch key order -- must be exactly device, kernel, grid, block, args.

12. Value struct with heap field -- value struct cannot contain Str, Array, etc. Use heap struct.

13. Returning a borrow -- borrows cannot be function return types. Return the value instead.

14. TOption variant names -- use Some/None for TOption, Ok/Err for TResult. Don't mix them.

15. GPU kernel using heap types -- GPU kernels cannot use Str, Array, Map, or TCallable. Use rawptr<T> and primitives.

---

13) Checklist Before Declaring Done

• Program has correct header order: module, exports, imports, digest
• Every block ends with exactly one terminator (ret, br, cbr, switch, unreachable)
• Every const suffix matches the declared SSA type
• Every getfield uses a borrow/mutborrow receiver
• Every setfield uses a mutborrow receiver with val= key
• No borrow crosses a block boundary
• No borrow appears in a phi node
• No borrow is returned from a function
• No value is used after being consumed/moved
• Collection ops use correct receiver mode (mutborrow for mutation, borrow for reads)
• Map key types have hash + eq impls (unless using Str)
• GPU kernels use only primitives and rawptr<T> -- no heap types
• magpie --entry <file> --output json build produces zero errors
• magpie explain <CODE> consulted for any unfamiliar error