Sarek_hip.Hip_coopmatThe HIP cooperative-matrix probe (specs/coopmat-hip-wmma.md).
HIP has no query that lists matrix-instruction configurations the way vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR does, so the probe is a rule over two facts the runtime does report: the gfx family, read from hipDeviceComputeCapability (gfx1100 reports 11.0, the Raphael iGPU gfx1036 reports 10.3), and the wavefront size. The configurations offered are the ones Sarek's HIP lowering emits (Sarek_ir_cuda_coopmat) and that were executed on gfx1100: the gfx11 WMMA builtins are wave32-only.
Every other device is reported as probed with NO configuration — Some with an empty list, not None — so a refusal names the rule instead of "not probed". That includes devices that do have matrix units this lowering does not emit: gfx12 (WMMA with different builtins and layouts) and CDNA (MFMA). They are refused until they are lowered and measured.
module C = Sarek_coopmatval shape : C.shapeval config :
C.component_type ->
C.component_type ->
C.component_type ->
C.configval gfx11_configs : C.config list__builtin_amdgcn_wmma_f32_16x16x16_f16_w32, its bf16 twin __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 (specs/matrix-pattern-rewrite.md) and __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32 (both operand signednesses are immediates of the latter; a 32-bit wrapping sum has the same bits whether it is read as s32 or u32).
val support :
compute_capability:(int * int) ->
warp_size:int ->
C.device_supportsupport ~compute_capability ~warp_size is the rule, separated from the device record so a test can drive it with any device's facts.
val probe : Hip_api.Device.t -> C.device_support