Struct KernelDescriptor

Source

pub struct KernelDescriptor {
    pub id: String,
    pub bindings: BindingLayout,
    pub dispatch: Dispatch,
    pub body: KernelBody,
}

Expand description

The full kernel descriptor.

Fields§

§id: String

Stable kernel identifier (for caching). Computed from the content hash by lower::lower. Empty string until lowering assigns it.

§bindings: BindingLayout§dispatch: Dispatch§body: KernelBody

Implementations§

Source §

impl KernelDescriptor

Source

pub fn summary(&self) -> String

One-line human-readable summary. Useful for diagnostic output. Format: "<id>: N ops, M bindings, K child bodies, L literals, dispatch [x, y, z]".

Source

pub fn summary_compact(&self) -> String

Terser alternative to Self::summary. Format: "<id>(N ops, M bindings)". Useful for compact terminal output where the full summary is too noisy.

Source

pub fn total_ops(&self) -> usize

Total op count across the parent body AND every nested child body, recursively. The parent-only body.ops.len() is the flat count; this is the deep count.

Source

pub fn body_count(&self) -> usize

Total number of bodies (the parent counts as 1, plus each nested child recursively). Useful for “how nested is this kernel?” telemetry - a kernel with one big flat body has body_count() == 1; one with deep control flow has more.

Source

pub fn max_body_depth(&self) -> usize

Maximum nesting depth of child bodies. A flat kernel returns 0. A kernel with one If returns 1. An If-inside-an-If returns 2. Useful for routing decisions (deeply-nested kernels may need a different optimization strategy).

Source

pub fn body_at(&self, path: &[usize]) -> Option<&KernelBody>

Look up a body by its path (a Vec of child-body indices). Empty path returns the parent body. Each element of path indexes into the child_bodies of the body it descends into. Returns None if any index is out of range.

Matches the body_path shape used by verify::VerifyError, so tooling can take a verify error and resolve it to the actual body the error refers to.

Source

pub fn is_empty(&self) -> bool

True iff the descriptor has no ops at all (no parent ops AND no ops in any child body). The dispatch geometry and bindings can still be populated - this only asks about op content.

Source

pub fn is_pure(&self) -> bool

True iff the descriptor is pure - no side-effecting ops anywhere. Inverse of has_side_effects. Pure kernels can be safely cached by descriptor identity since they produce no observable output (the only “result” is whatever value-flow the consumer inspects, which is fully determined by the descriptor).

Source

pub fn ops_iter(&self) -> KernelOpsIter<'_> ⓘ

Iterator over every KernelOp in the descriptor (parent body

every nested child body, depth-first pre-order). Useful for tooling that wants to walk all ops without writing the recursion themselves.

Source

pub fn find_op_by_id(&self, id: u32) -> Option<&KernelOp>

Find the first op anywhere in the descriptor whose result matches id. Per-body id space means an id may be reused across child bodies - this returns the FIRST match in DFS pre-order. For a given body’s view, callers should iterate body.ops directly.

Source

pub fn dispatch_total_threads(&self) -> u32

Total threads per workgroup (the product of dispatch.workgroup_size). Saturates on overflow rather than wrapping. Useful for per-dispatch resource calculations (shared memory budget, register pressure, etc.).

Source

pub fn with_id(&self, id: impl Into<String>) -> Self

Return a clone of this descriptor with a new id field. Body, bindings, dispatch all unchanged. Useful for tooling that wants to fork a descriptor for ablation testing or versioning.

Source

pub fn has_side_effects(&self) -> bool

True iff the descriptor has at least one side-effecting op (memory write, atomic, sync/async op, barrier, control exit, indirect dispatch, call, or opaque extension). A pure descriptor with no side effects produces no observable output, so a caller is free to drop it entirely.

“Side effect” here means observable-or-cross-thread: AsyncLoad writes shared memory other threads read, AsyncWait/Barrier are sync points, and IndirectDispatch reconfigures the grid, all are unsafe to drop even though none produces a global-buffer write.