NVIDIA / NVIDIA/cutlass

maximum_active_blocks wrappers pass int into CudaHostAdapter* and no longer compile

Open
#3,534 2 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

CUTLASS C++
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-394
  • gemm/device/gemm_universal_adapter.h:738-739 (2.x alias path)
  • gemm/device/rank_k.h:462
  • gemm/device/rank_2k.h:499
  • gemm/device/symm.h:554
  • gemm/device/gemm_universal_with_broadcast.h:337
  • gemm/device/gemm_universal_streamk_with_broadcast.h:337
  • gemm/device/gemm_universal_with_absmax.h:355
  • gemm/device/gemm_with_k_reduction.h:366
  • gemm/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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.