Sarek_cuda.Cuda_coopmatThe CUDA cooperative-matrix probe (specs/coopmat-hip-wmma.md, specs/matrix-pattern-rewrite.md), shared by the CUDA/C and CUDA/PTX backends.
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 both lowerings emit (Sarek_ir_cuda_coopmat's nvcuda::wmma arm and Sarek_ir_ptx_coopmat's wmma.* instructions): 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 refused by both, so they are not offered). lowering only labels the evidence.
Executed in this repository only through ZLUDA on gfx1100 (which reports compute capability 8.8 and rejects wmma.*), 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 lowering_name : lowering -> stringval 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_for :
lowering:lowering ->
device_name:string ->
compute_capability:(int * int) ->
warp_size:int ->
C.device_supportsupport_for ~lowering ~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.
val support :
device_name:string ->
compute_capability:(int * int) ->
warp_size:int ->
C.device_supportThe CUDA/C rule.