Sarek_cuda.Cuda_coopmatThe CUDA/C cooperative-matrix probe (specs/coopmat-hip-wmma.md, specs/matrix-pattern-rewrite.md).
The CUDA driver has no query that lists matrix-instruction configurations, so the probe is a rule over two facts it does report: the compute capability and the warp size. The configurations offered are the ones Sarek_ir_cuda_coopmat's nvcuda::wmma arm emits: 16x16x16 on 32-wide warps, f16 operands with an f32 accumulator from sm_70, and 8-bit operands of the SAME signedness with a 32-bit integer accumulator from sm_72 (the mixed-sign arms and bf16 loaded from float32 are #error there, so they are not offered).
Only the CUDA/C backend reports this: CUDA/PTX has no coopmat lowering and keeps reporting None. Executed in this repository only through ZLUDA on gfx1100 (which reports compute capability 8.8), never on NVIDIA hardware.
module C = Sarek_coopmatval shape : C.shapeval config :
C.component_type ->
C.component_type ->
C.component_type ->
C.configval f16_configs : C.config listval int8_configs : C.config listZLUDA (a CUDA driver API over ROCm) rejects every wmma.* instruction and 8-bit mma with unsigned operands: the module fails to load (cuModuleGetFunction -> CUDA_ERROR_NOT_FOUND), measured on gfx1100 (specs/zluda-mma-nvrtc13.md, F-1; docs/measurements/zluda-mma-probe-2026-10-11). Its devices carry "ZLUDA" in Its devices carry "ZLUDA" in their name and report compute capability 8.8, which no NVIDIA part does.
val support :
device_name:string ->
compute_capability:(int * int) ->
warp_size:int ->
C.device_supportsupport ~device_name ~compute_capability ~warp_size is the rule, separated from the device record so a test can drive it with any device's facts.