maximum_active_blocks wrappers pass int into CudaHostAdapter* and no longer compile
Description
Since 3.4, GemmUniversalBase::maximum_active_blocks takes a CudaHostAdapter*:
// include/cutlass/gemm/device/gemm_universal_base.h:350
static int maximum_active_blocks(CudaHostAdapter *cuda_adapter = nullptr)but the 2.x-style device wrappers still declare the pre-3.4 signature (int smem_capacity = -1) and forward the int into that pointer parameter, which does not convert:
static int maximum_active_blocks(int smem_capacity = -1) {
return UnderlyingOperator::maximum_active_blocks(smem_capacity); // int -> CudaHostAdapter*: error
}Calling any of these wrappers fails to compile:
error: argument of type "int" is incompatible with parameter of type "cutlass::CudaHostAdapter *"Affected call sites (all forwarding smem_capacity):
gemm/device/gemm_universal.h:393-394gemm/device/gemm_universal_adapter.h:738-739(2.x alias path)gemm/device/rank_k.h:462gemm/device/rank_2k.h:499gemm/device/symm.h:554gemm/device/gemm_universal_with_broadcast.h:337gemm/device/gemm_universal_streamk_with_broadcast.h:337gemm/device/gemm_universal_with_absmax.h:355gemm/device/gemm_with_k_reduction.h:366gemm/device/gemm_layernorm_mainloop_fusion.h:336-337
The 3.x adapter (gemm_universal_adapter.h:270) already updated its own declaration correctly (maximum_active_blocks(int /* smem_capacity */ = -1) ignoring the value), so only these forwarders were missed. Nothing in-tree calls them anymore (the profiler moved to the new API), which is why CI does not see it.
Reproduction: compile a TU that calls e.g.
#include "cutlass/gemm/device/gemm.h"
using Gemm = cutlass::gemm::device::Gemm<float, cutlass::layout::ColumnMajor,
float, cutlass::layout::ColumnMajor,
float, cutlass::layout::RowMajor>;
int n = Gemm::maximum_active_blocks();with nvcc or g++ against the current headers; it errors at the forwarded argument.
Suggested fix
Drop the unused parameter in the ten wrappers to match the base signature (maximum_active_blocks() with no arguments), mirroring what gemm_universal_adapter.h:270 already does.
Source: NVIDIA/cutlass