maximum_active_blocks wrappers pass int into CudaHostAdapter* and no longer compile
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
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.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start with the affected wrappers in gemm/device/gemm_universal.h, gemm_universal_adapter.h, rank_k.h, rank_2k.h, symm.h, the broadcast and absmax headers, gemm_with_k_reduction.h, and gemm_layernorm_mainloop_fusion.h. Compare them with GemmUniversalBase::maximum_active_blocks and the corrected adapter declaration, then compile the provided Gemm reproduction to verify all wrappers use the current API.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- backend-api-design
- Issue type
- Bug
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Active
- Clarity
- Clearly specified
- Newbie friendliness
- 82/100