Module Sarek_metal.Metal_coopmat

The Metal cooperative-matrix capability rule (specs/coopmat-metal-simdgroup.md).

Metal has no query that lists matrix configurations, and MSL offers exactly one lowered here: simdgroup_half8x8 x half8x8 + float8x8 -> float8x8 (docs/design/f16-relaxed-accuracy.md §2). The probe is therefore a compile check on the device rather than a family table: a kernel using that configuration is compiled on the device (compile evidence, not execution), and its pipeline reports the SIMD-group width. The pure rule below turns those two facts into a Sarek_coopmat.device_support, so it can be driven without a device. A device whose compiler rejects the probe is reported as probed with NO configuration — Some with an empty list — so a refusal says why instead of "not probed"; that pre-Apple7 GPUs and Intel/AMD Macs are in this class is inferred, not observed. The launch check and Sarek_gemm.Coopmat assume whole 32-wide SIMD-groups, so a device reporting another width is likewise offered no configuration.

module C = Sarek_coopmat
val shape : C.shape
val f16_config : C.config

simdgroup_half8x8 x half8x8 -> float8x8, 8x8x8.

val probe_kernel_name : string
val probe_source : string

Exercises the lowered configuration with the same statements the generator emits: declare, load, multiply-accumulate, store.

val required_simd_width : int

The SIMD-group width the launch check and the GEMM builder assume.

val support : simd_width:int -> probe:(unit, string) Stdlib.result -> C.device_support

support ~simd_width ~probe is the device's support given the SIMD-group width its pipeline reported and the outcome of compiling probe_source.