[Profiler] Skip cuBLAS verification when C and D element types differ

[LiRunGuo]

Motivation

cutlass_profiler aborts the entire profiling session (SIGABRT, exit code 134) as soon as it profiles an SM90 GEMM that writes narrow FP8 output from an FP32 C:

$ ./build/tools/profiler/cutlass_profiler \
    --kernels=cutlass3x_sm90_tensorop_gemm_e4m3_e5m2_f32_f32_e4m3_128x128x128_1x2x1_0_tnn_align16_warpspecialized_epi_nosmem \
    --m=4096 --n=4096 --k=4096
...
          Status: Success
    Verification: ON
     Disposition: Passed
terminate called after throwing an instance of 'std::runtime_error'
  what():  Failed to allocate workspace
Aborted (core dumped)

Because the profiler is usually run over the whole kernel set, this kills the run at the first such kernel and all already-collected results are lost (the final report is never printed).

Root cause. In GemmOperationProfiler::verify_with_cublas_, the Reference workspace is allocated with desc.D.element and the D layout, but that single buffer is handed to cuBLAS as both the C and the D operand, while cublasLtGemmExDispatcher builds Cdesc and Ddesc with data_type_C:

gemm_workspace_.arguments.C = gemm_workspace_.Reference->data();
gemm_workspace_.arguments.D = gemm_workspace_.Reference->data();
...
cublasLtMatrixLayoutCreate(&Cdesc, data_type_C, m, n, configuration.ldc);
cublasLtMatrixLayoutCreate(&Ddesc, data_type_C, m, n, configuration.ldd);

When the kernel's C and D element types differ (FP32 C + FP8 D), cuBLASLt walks the 1-byte-per-element Reference buffer as if it held 4-byte elements, faults with an illegal memory access, and leaves a sticky cudaErrorIllegalAddress (700) in the context. The next cudaMalloc therefore fails, so the abort surfaces as a completely misleading "Failed to allocate workspace" from Handle::set_workspace_size, far away from the real fault. Confirmed with a cudaMalloc interposer:

[hook] cudaMalloc size=0.134 GB rc=0
[hook] cudaMalloc size=0.004 GB rc=700   <-- sticky illegal-address from the cuBLAS verification
terminate called after throwing an instance of 'std::runtime_error'
  what():  Failed to allocate workspace

The existing guard in cublas_satisfies() only covered one FP8 operand ordering (A=E5M2, B=E4M3, C=F32, D!=F32, tools/profiler/src/cublas_helpers.cu:272). The symmetric ordering (A=E4M3, B=E5M2) and even the homogeneous E4M3 x E4M3 -> FP32 C -> E4M3 D kernels are not covered and take the same crashing path.

Modifications

• tools/profiler/src/cublas_helpers.cu: cublas_satisfies() now returns kErrorNotSupported whenever desc.C.element != desc.D.element. This is the exact precondition of the single-buffer/single-CType verification path, so those problems are reported as not verified instead of being sent to a cuBLAS call that cannot represent them.
• tools/library/src/handle.cu: the workspace allocation failure message now reports the requested size and the CUDA error name/string, so a poisoned context can no longer be mistaken for a workspace sizing problem.

The guard affects 212 of the 604 SM90 GEMM kernels of the default instantiation level, all of which have a narrow FP8 D and a void/s8/f16/bf16/f32 C. It cannot remove verification coverage anywhere else: no SM80/SM100 kernel in the library has C.element != D.element, and every C!=D kernel that did complete before the fix reports exactly the same dispositions afterwards.

Verification

H200 (sm90a), CUDA 12.8, driver 580, CUTLASS 4.8.0 (147295a3), -DCUTLASS_NVCC_ARCHS=90a.

Single kernel, --m=4096 --n=4096 --k=4096 (the kernel above):

	before	after
exit code	134 (Aborted (core dumped))	0
disposition	none (session aborted)	Passed (cuBLAS: not run, device reference: passed)
performance	–	0.288 ms / 477 TFLOP/s

Full SM90 dense GEMM sweep, --kernels="cutlass3x_sm90_*" --profiling-iterations=1, 604 kernels:

shape	before	after
4096x4096x4096	aborted after 434/604 problems (exit 134, no report)	604/604, exit 0, 602 passed / 2 not verified
1024x1024x1024	–	604/604, exit 0
16384x128x1024	–	604/604, exit 0
128x128x64	–	604/604, exit 0
8x8192x8192	–	536/536 runnable kernels, exit 0
256x2048x3520	604/604, exit 0	604/604, exit 0 (identical per-kernel dispositions/cuBLAS results)

Per-kernel comparison for the 434 problems that completed before the abort: 0 differences in Disposition, cuBLAS, and reference_device results; 170 previously unreachable kernels are now profiled.

Regression tests on the same host (all green, none of these paths are touched by the patch): cutlass_test_unit_gemm_device_tensorop_sm90 92/92, ..._alignx_sm90 77/77, ..._epilogue_fusion_sm90 65/65, ..._gmma_rs_warpspecialized_sm90 16/16, ..._sm90_blockwise 3/3, ..._sm90_stream_k 35/35 in 25 min. (..._sm90_group_gemm and ..._sm90_ptr_array contain only SM100 test cases, so they register 0 tests when built for 90a.)

Notes

A complete fix would let the profiler verify these kernels by allocating separate, correctly typed C and D buffers (Ddesc would then need the D element type and a type-aware comparison against the CUTLASS result). That is a larger change to the verification harness, so this PR takes the conservative route and reports the affected problems as not verified, mirroring how C=void and C=s8 kernels are already handled.