Module avail 

2026-09-30: Which kernels a device can run: a rule’s kernel is available when the class compiles its module and names the function, no MODEL.toml declares it absent, the class’s build defines do not compile its region out, and the device runs natively the instruction kind its source guard stands for (kernels/DEVICES.toml [[guard]]).

Owner: metrale-circuit (hardware). Invariants:

• A kernel’s requirement is read from the source it compiles from: the preprocessor region holding its __global__ definition. No kernel is marked by hand.
• The build and the device must agree: a class that compiles a guarded region out while one of its devices claims the instruction, or keeps it while a device lacks it, is an error (check_build), never resolved in either’s favour.

Structs

Availability

2026-09-30: The kernels of a rule set on one device.

Enums

Absence

2026-09-30: Why a kernel is not available on the device.

Functions

availability

2026-09-30: The rule kernels device can run from sources of class.

check_build

2026-09-30: Fail when the class’s build defines contradict the device’s instruction kinds.

guard_of

2026-09-30: The guard whose region holds the __global__ definition of func in text, when every definition sits in one.

kernel_status

2026-09-30: Whether device can run kernel k from sources of class, and the instruction kind its source guard needs (None: unguarded or not compiled).

names_word

2026-09-30: word appears in text as a whole identifier.

without_fp4_kernel

2026-10-01: The nodes of c that multiply an NVFP4 activation (nvfp4_mma) and that no compiled FP4 block-scale kernel of the class runs: no rule with a pattern element that fits the node has every kernel available and one that needs the FP4 block-scale MMA, and no family that implements the node has such a kernel available. kernel answers for one kernel: None when the device cannot run it, else whether its source guard needs the FP4 block-scale MMA. Row ranges and policy conditions are not consulted: they choose among the kernels a class compiles, not which kernels it compiles.