[bheu]

📌 Summary

Fix a GCC build failure in the CUDA EP for SM120 (Blackwell consumer) FP4 QMoE MoE GEMM TMA warp-specialized kernels. All 11 moe_gemm_tma_ws_sm120_*.generated.cu files fail to compile on Linux with:

moe_gemm_tma_ws_launcher.inl:82:8: error: expression list treated as compound
expression in functional cast [-fpermissive]
   82 |         return ReturnType(std::forward<Args>(args)...);

(repeated per std::forward argument, then cc1plus: all warnings being treated as errors → build of onnxruntime_providers_cuda_sm120_tma fails).

🔍 Root cause

moe_gemm_tma_ws_launcher.inl:535:

using EpilogueScalars = decltype(EpilogueArguments{}.thread);

EpilogueArguments{} is a prvalue, so the member access .thread yields an xvalue, and decltype therefore gives an rvalue reference. For the SM120 epilogue this resolves to

cutlass::epilogue::fusion::FusionCallbacks<
    cutlass::epilogue::Sm90PtrArrayTmaWarpSpecialized<3,2,4,true,false,2>,
    cutlass::epilogue::fusion::LinearCombination<half_t,float,half_t,float,round_to_nearest>,
    cute::tuple<...>, cute::tuple<...>>::Arguments&&

(the instantiated ReturnType = ...::Arguments&& is visible in the compiler
diagnostic above).

The 8-argument call site then expands construct_if_true<IsBlackwell, EpilogueScalars>(...)
into Arguments&&(alpha, beta, nullptr, nullptr, alpha_ptr_array, nullptr, dAlpha, dBeta) — a
multi-expression functional cast to a reference type, which is ill-formed in C++. GCC rejects
it under -fpermissive; with the project's -Werror the build fails.

The call was written as C++20 parenthesized aggregate initialization of the value type
FusionCallbacks::Arguments (whose 8 members are exactly alpha, beta, alpha_ptr, beta_ptr, alpha_ptr_array, beta_ptr_array, dAlpha, dBeta — see
cutlass/epilogue/fusion/sm90_callbacks_tma_warpspecialized.hpp), i.e. the reference qualifier is
unintended.

For SM90 the thread member is already a plain value type
(cutlass::epilogue::thread::LinearCombination<...>::Params), which is why the bug only surfaces
for the SM120 epilogue.

🛠️ Fix

One line — strip the reference so the intended value type is used:

using EpilogueScalars = std::remove_reference_t<decltype(EpilogueArguments{}.thread)>;

• SM120: FusionCallbacks<...>::Arguments&& → FusionCallbacks<...>::Arguments (value type,
  8-member aggregate — the parenthesized init is well-formed and maps member-by-member).
• SM90: already a value type, so remove_reference_t is an identity no-op (no behavior change).

No new includes needed; no functional change to the constructed values.

✅ Verification

Environment: Linux x86_64, GCC 13, CUDA 13.3, CUTLASS 4.7.0, CMAKE_CUDA_ARCHITECTURES=86-real;120-real.

Check	Before fix	After fix
11 × moe_gemm_tma_ws_sm120_*.generated.cu (sm_120a)	❌ expression list treated as compound expression in functional cast [-fpermissive] (all files)	✅ all compile, 0 errors
moe_gemm_tma_ws_sm90_*.generated.cu (sm_90a, from onnxruntime_providers_cuda_sm90_tma)	✅	✅ unchanged (identity transform)
Full ./build.sh shared lib + wheel	❌ fails at onnxruntime_providers_cuda_sm120_tma	✅ built and installed successfully

Note: the SM90 files are compiled by a separate object library with hardcoded 90a-real
architecture, so any build with ≥ one SM90+ architecture exercises both epilogue paths; both are
verified above.

📝 Note (not fixed here, follow-up candidate)

moe_gemm_tma_ws_mixed_input_launcher.inl:201 uses the same
decltype(EpilogueArguments{}.thread) pattern and then declares EpilogueScalars scalars; and
mutates its fields directly. It is currently safe because that launcher only instantiates SM90
epilogues (value type), but it would break the same way if an SM120 epilogue were ever routed
there. Kept out of this PR to keep it minimal.

[azure-pipelines]

Azure Pipelines:
There may be pipelines that require an authorized user to comment /azp run to run.

[Copilot]

🟢 Approval recommended

The one-line change is a well-understood, behavior-preserving reference-stripping transform that fixes the SM120 build without affecting the SM90 path, and both were verified to compile.

Pull request overview

This PR fixes a GCC -Werror build failure for the CUDA EP's SM120 (Blackwell consumer) FP4 QMoE MoE GEMM TMA warp-specialized kernels. In moe_gemm_tma_ws_launcher.inl, the alias EpilogueScalars was derived via decltype(EpilogueArguments{}.thread), which for the SM120 epilogue resolves to a reference type (...::Arguments&&). Feeding that reference type into the C++20 parenthesized aggregate-init helper construct_if_true (return ReturnType(args...)) produces a multi-argument functional cast to a reference — ill-formed C++ that GCC rejects, failing the build of onnxruntime_providers_cuda_sm120_tma. Wrapping the decltype in std::remove_reference_t restores the intended value type. The change is a no-op for the SM90 epilogue, whose thread member is already a value type.

Changes:

• Strip the reference qualifier from the EpilogueScalars alias so the value aggregate type is used for parenthesized aggregate initialization.

File summaries

File	Description
onnxruntime/contrib_ops/cuda/llm/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl	Wraps decltype(EpilogueArguments{}.thread) in std::remove_reference_t to yield the epilogue's value type, fixing the SM120 functional-cast build error while leaving SM90 unchanged.

I found no actionable issues. The fix is minimal and safe: std::remove_reference_t strips the offending reference for SM120 and acts as an identity transform for SM90's already-value type, so no constructed values change. std::remove_reference_t is available transitively (the file already relies on std::forward), and the author verified both epilogue paths build. The analogous pattern in moe_gemm_tma_ws_mixed_input_launcher.inl:201 is outside this diff and documented as intentionally deferred (currently safe since it only instantiates SM90 value-type epilogues).

Review details

• Files reviewed: 1/1 changed files
• Comments generated: 0
• Review effort level: Balanced

---

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.