Enum KernelOpKind 

pub enum KernelOpKind {
    Literal,
    Copy,
    LocalInvocationId,
    GlobalInvocationId,
    WorkgroupId,
    SubgroupLocalId,
    SubgroupSize,
    LoopIndex {
        loop_var: Name,
    },
    LoopCarrierInit {
        name: Name,
    },
    LoopCarrier {
        name: Name,
    },
    LoopCarrierEnd {
        name: Name,
    },
    LoadGlobal,
    LoadShared,
    LoadConstant,
    BufferLength,
    StoreGlobal,
    StoreShared,
    BinOpKind(BinOp),
    UnOpKind(UnOp),
    Fma,
    MatrixMma {
        shape: MatrixMmaShape,
        a_layout: MatrixMmaLayout,
        b_layout: MatrixMmaLayout,
        a_type: MatrixMmaElement,
        b_type: MatrixMmaElement,
        accum_type: MatrixMmaElement,
    },
    Select,
    Cast {
        target: DataType,
    },
    Atomic {
        op: AtomicOp,
        ordering: MemoryOrdering,
    },
    SubgroupBallot,
    SubgroupShuffle,
    SubgroupBroadcast,
    SubgroupReduce {
        op: SubgroupReduceOp,
    },
    StructuredIfThen,
    StructuredIfThenElse,
    StructuredForLoop {
        loop_var: Name,
    },
    StructuredBlock,
    Return,
    Barrier {
        ordering: MemoryOrdering,
    },
    Region {
        generator: Name,
    },
    AsyncLoad {
        tag: Name,
    },
    AsyncStore {
        tag: Name,
    },
    AsyncWait {
        tag: Name,
    },
    Trap {
        tag: Name,
    },
    Resume {
        tag: Name,
    },
    IndirectDispatch {
        count_offset: u64,
    },
    Call {
        op_id: Name,
    },
    OpaqueExpr(Box<OpaqueExprData>),
    OpaqueNode(Box<OpaqueNodeData>),
}

Lowered op kinds. Closed enum but covers the entire vyre IR surface. Adding a new vyre IR variant requires a matching variant here AND emit rules in every vyre-emit-* crate - that’s the cost of substrate parity.

Operand semantics are documented per variant. Reading a kind without reading its operand contract gives wrong code.

Variants

Literal

Operand 0 = index into KernelBody.literals. Result is the literal value typed per the literal pool entry.

Copy

Snapshot a result value. Operand 0 = source result id. Result is a fresh SSA value with the source value at this program point. This is required when a source-level Let captures a mutable loop carrier: aliasing the carrier result id would read the later carrier value after a subsequent Assign.

LocalInvocationId

LocalInvocationId.x/y/z. Operand 0 = axis (0/1/2) as a small inline literal (NOT a literal-pool reference - emit picks the builtin directly).

GlobalInvocationId

GlobalInvocationId.x/y/z.

WorkgroupId

WorkgroupId.x/y/z.

SubgroupLocalId

Subgroup local invocation id (a.k.a. lane id).

SubgroupSize

Subgroup size.

LoopIndex

Current induction value for the nearest structured loop that declared this variable. Produced as the first op in that loop’s child body so uses of Expr::Var(loop_var) remain SSA-shaped instead of resolving to the loop’s lower bound.

Fields

loop_var: Name

LoopCarrierInit

Initialize the loop-carrier slot for name from the pre-loop SSA value. Emitted ONCE in the PARENT body before the StructuredForLoop op. Operands: [seed_value_id]. No result. Emitters allocate a function-scope LocalVariable keyed by name (if not already allocated) and Store(local, seed_value) in the parent block.

Fields

name: Name

LoopCarrier

Pure read of the carrier slot for name. Operands: []. Result: the SSA id that in-loop reads of the source-level variable resolve to. Emit semantics: Load from the function-local allocated by the matching LoopCarrierInit. Used in three places per loop: (a) once at the top of each iteration so per-iteration reads resolve to the latest stored value; (b) in the parent body AFTER the loop so post-loop readers observe the loop’s final value. Without this op, Node::Assign inside a loop body would have no observable effect on subsequent iterations - name resolution would always pick the pre-loop SSA, which is baked at lowering time.

Fields

name: Name

LoopCarrierEnd

Loop-carried-variable write at iteration end. Operands: [final_value_id]. No result. Pairs with LoopCarrier { name } to commit the iteration’s final value of name back to the carrier local so the next iteration (or the post-loop reader) observes it.

Fields

name: Name

LoadGlobal

load(buf, index). Operands: [binding_slot, index_op_id]. Result is the loaded value, dtype = binding’s element type.

LoadShared

load(buf, index) for a workgroup-shared binding.

LoadConstant

load(buf, index) for a constant/uniform binding.

BufferLength

Buffer length (number of elements). Operand 0 = binding_slot inline. Result is u32.

StoreGlobal

store(buf, index, value). Operands: [binding_slot, index_op_id, value_op_id]. Result: None.

StoreShared

store(buf, index, value) for a workgroup-shared binding.

BinOpKind(BinOp)

Binary op. Operands: [left_op_id, right_op_id]. Result has the dtype dictated by the operand dtypes (per vyre-spec rules).

UnOpKind(UnOp)

Unary op. Operands: operand_op_id. Result dtype per spec.

Fma

Fused multiply-add: a * b + c. Operands: [a_id, b_id, c_id].

MatrixMma

Matrix multiply-accumulate fragment op.

Operand contract for M16N8K16/F16/F16/F32: [a0,a1,a2,a3, b0,b1, c0,c1,c2,c3], where a* and b* are packed 16-bit fragment words and c* are f32 accumulators. result is the first of four consecutive result ids (result..result+4). This keeps the descriptor SSA-shaped without adding backend-specific register-fragment objects to the neutral IR.

Fields

shape: MatrixMmaShape

a_layout: MatrixMmaLayout

b_layout: MatrixMmaLayout

a_type: MatrixMmaElement

b_type: MatrixMmaElement

accum_type: MatrixMmaElement

Select

Conditional select: if cond { true_val } else { false_val }. Operands: [cond_id, true_val_id, false_val_id].

Cast

Type cast. Operands: value_id. The target dtype is on the op.

Fields

target: DataType

Atomic

Atomic op. Operands: [binding_slot, index_op_id, value_op_id] for most ops. CompareExchange variants prepend expected_op_id: [binding_slot, index_op_id, expected_op_id, value_op_id].

Fields

op: AtomicOp

ordering: MemoryOrdering

SubgroupBallot

Operand 0 = bool-typed cond_op_id. Result is u32 ballot mask.

SubgroupShuffle

Operands: [value_op_id, lane_op_id]. Result has the value’s dtype.

SubgroupBroadcast

Operands: [value_op_id, lane_op_id]. Broadcasts value from the lane named by lane (uniform) to every lane; result has the value’s dtype. Distinct from SubgroupShuffle (per-lane source), broadcast requires a uniform source lane and emits subgroupBroadcast.

SubgroupReduce

Operand 0 = value_op_id. Reduces across the subgroup with op; result has the value’s dtype.

Fields

op: SubgroupReduceOp

StructuredIfThen

if (cond) { body }. Operands: [cond_op_id, child_body_index]. child_body_index references KernelBody.child_bodies. Result: None.

StructuredIfThenElse

if (cond) { then } else { otherwise }. Operands: [cond_op_id, then_body_index, otherwise_body_index].

StructuredForLoop

for (var = lo; var < hi; ++var) { body }. Operands: [lo_op_id, hi_op_id, body_index]. The loop variable name is embedded on the op (preserved for debug, not for codegen).

Fields

loop_var: Name

StructuredBlock

Inline statement block - explicit grouping; semantically a no-op (body is flattened during emit). Operand 0 = body_index.

Return

Function/kernel return. Operands: empty. Result: None.

Barrier

Memory barrier with explicit ordering.

Fields

ordering: MemoryOrdering

Region

Tracing/grouping marker (vyre IR Node::Region). Operand 0 = body_index. Carries no execution semantics; emitters MAY pass through as a comment or annotation. SEPARATION_AUDIT S5 plans to move this to a sidecar; until then it’s an op so the descriptor preserves it round-trip.

Fields

generator: Name

AsyncLoad

cp.async-style global-to-shared copy. Operands: [src_binding, dst_binding, offset_op_id, size_op_id]. tag ties the load to a matching AsyncWait.

Fields

tag: Name

AsyncStore

Mirror of AsyncLoad for shared-to-global. Operands: [src_binding, dst_binding, offset_op_id, size_op_id].

Fields

tag: Name

AsyncWait

Wait on a previously-issued AsyncLoad/Store. Operands: empty.

Fields

tag: Name

Trap

Trap into a host-side effect handler. Operands: address_op_id.

Fields

tag: Name

Resume

Resume from a previously-trapped effect.

Fields

tag: Name

IndirectDispatch

Indirect-dispatch hint. The dispatch shape comes from count_buffer[count_offset]. Operand 0 = count_buffer binding_slot. Result: None.

Fields

count_offset: u64

Call

Call into a known op-id (e.g., a vyre-primitives builder surface). Operand list is the call’s args. The op_id picks the callee at emit time.

Fields

op_id: Name

OpaqueExpr(Box<OpaqueExprData>)

Opaque expression extension. The extension id resolves through vyre-core’s extension registry. Emitters that don’t recognize the extension MUST surface an error rather than silently emit nothing.

Boxed to keep the common-case KernelOpKind small: most ops are Literal/BinOp/Load/Store at ≤16 bytes; without boxing, every op in the ops Vec pays the 52-byte OpaqueExpr tax.

OpaqueNode(Box<OpaqueNodeData>)

Opaque statement-node extension.

Trait Implementations

impl Clone for KernelOpKind

fn clone(&self) -> KernelOpKind

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for KernelOpKind

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl<'de> Deserialize<'de> for KernelOpKind

fn deserialize<__D>(__deserializer: __D) -> Result<Self, __D::Error>

where __D: Deserializer<'de>,

Deserialize this value from the given Serde deserializer.

impl Eq for KernelOpKind

impl Hash for KernelOpKind

fn hash<H: Hasher>(&self, state: &mut H)

Feeds this value into the given Hasher.

fn hash_slice<H>(data: &[Self], state: &mut H)

where H: Hasher, Self: Sized,

Feeds a slice of this type into the given Hasher.

impl PartialEq for KernelOpKind

fn eq(&self, other: &KernelOpKind) -> bool

Equality operator ==.

fn ne(&self, other: &Rhs) -> bool

Inequality operator !=.

impl Serialize for KernelOpKind

fn serialize<__S>(&self, __serializer: __S) -> Result<__S::Ok, __S::Error>

where __S: Serializer,

Serialize this value into the given Serde serializer.

impl StructuralPartialEq for KernelOpKind