[Bug]: [CUDA] cute::make_int_sequence<cute::size(...)> fails in dependent template contexts

System Info

• Discovered on Windows Server 2025 with MSVC as the NVCC host compiler
• Visual Studio Build Tools with MSVC 19.51.36248
• MSVC toolset 14.51.36231
• CUDA/NVCC 13.2, V13.2.51
• C++20
• --gpu-architecture=sm_89
• --expt-relaxed-constexpr
• --expt-extended-lambda
• /Zc:preprocessor
• CuTe from the CUTLASS dependency vendored with FlashInfer
• CUTLASS revision e6e2cc29
• The reproduced diagnostic is emitted by NVCC during dependent-template instantiation
• The defect does not appear to be specific to Windows or MSVC

Who can help?

No response

Information

• The official example scripts
• My own modified scripts

Tasks

• An officially supported task in the examples folder (such as GLUE/SQuAD, ...)
• My own task or dataset (give details below)

Reproduction

Background

Hi!

I encountered this issue on Windows while building SystemPanic/vllm-windows for a DeepSeek-V4-Flash-0731 deployment. The affected FP8 blockscale GEMM kernel originates from TensorRT-LLM and is vendored through FlashInfer, so I am reporting the issue here in the hope that the portable source fix can be incorporated upstream. Although the failure was reproduced with NVCC 13.2 using MSVC as the host compiler, the diagnostic occurs during NVCC dependent-template instantiation and may not be Windows-specific.

The affected TensorRT-LLM source file is:

cpp/tensorrt_llm/kernels/cutlass_kernels/fp8_blockscale_gemm/ada_blockwise_gemm/sm89_fp8_gemm_1d1d.cuh

The corresponding vendored FlashInfer path is:

csrc/nv_internal/tensorrt_llm/kernels/cutlass_kernels/fp8_blockscale_gemm/ada_blockwise_gemm/sm89_fp8_gemm_1d1d.cuh

The issue was reproduced with CUDA/NVCC 13.2 using MSVC as the host compiler. The diagnostic is emitted during NVCC dependent-template instantiation.

The affected SM89 kernel passes the CuTe integral-constant wrapper returned by cute::size(...) directly as the non-type template argument to cute::make_int_sequence<N>:

cute::for_each(
    cute::make_int_sequence<cute::size(epi)>{},
    auto i { epi(i) = ElementOutput(accum(i)); });

cute::for_each(
    cute::make_int_sequence<cute::size(scale)>{},
    auto i { scale(i) = tXrSFA(i) * tXrSFB(0); });

The failure occurs when the object passed to cute::size(...) is a dependent function parameter or local object inside the surrounding template context.

Problem

cute::size(object) returns a class-type integral-constant wrapper such as cute::C<N>. The wrapper provides a static value member and a constant conversion operator.

cute::make_int_sequence<N> ultimately forms an integer sequence whose non-type template argument must be a valid converted constant expression of the expected integral type.

In a dependent-template context, NVCC does not accept cute::size(object) directly as that template argument because the expression syntactically depends on the function parameter or local object.

NVCC reports:

error: expression must have a constant value
note #2689-D: the value of parameter "obj" cannot be used as a constant

The same expression compiles in a non-dependent context, isolating the reproduced failure to dependent-template constant-expression evaluation.

Reproduction

Minimal dependent-template reproduction:

#include <cute/tensor.hpp>
#include <cute/numeric/integer_sequence.hpp>

using namespace cute;

template <class Tensor>
__host__ __device__ void use_seq(Tensor const& obj) {
  cute::for_each(
      cute::make_int_sequence<cute::size(obj)>{},
      auto i { (void)obj(i); });
}

__global__ void launcher() {
  auto frag = cute::make_fragment_like<float>(
      cute::make_shape(cute::Int<8>{}, cute::Int<4>{}));
  use_seq(frag);
}

Portable compile command:

nvcc -c -std=c++20 original_dep.cu -o original_dep.obj \
  --gpu-architecture=sm_89 \
  --expt-relaxed-constexpr \
  --expt-extended-lambda \
  -I<PATH_TO_CUTLASS>/include \
  -I<PATH_TO_CUDA>/include

The include paths should point to the CUTLASS/CuTe and CUDA installations being tested.

The reproduction used:

CUDA/NVCC 13.2
CUTLASS revision e6e2cc29
C++20
SM89 compilation target
MSVC as the NVCC host compiler

The original dependent form failed consistently in two attempts:

return code: 1

original_dep.cu(10): error: expression must have a constant value
    cute::for_each(cute::make_int_sequence<cute::size(obj)>{},
                                           ^

original_dep.cu(10): note #2689-D: the value of parameter "obj" (declared at line 9) cannot be used as a constant
    cute::for_each(cute::make_int_sequence<cute::size(obj)>{},
                                                      ^

detected during instantiation of:
void use_seq(const Tensor&) [
  with Tensor=cute::Tensor<
      cute::ArrayEngine<float, 32ULL>,
      cute::Layout<
          cute::tuple<cute::_8, cute::_4>,
          cute::tuple<cute::_1, cute::_8>>>
]

1 error detected in the compilation of original_dep.cu.

Controlled results:

ORIGINAL non-dependent:
  cute::make_int_sequence<cute::size(object)>
  return code: 0

CORRECTED non-dependent:
  cute::make_int_sequence<decltype(cute::size(object))::value>
  return code: 0

CONTROL:
  cute::make_int_sequence<32>
  return code: 0

ORIGINAL dependent, attempt 1:
  cute::make_int_sequence<cute::size(object)>
  return code: 1

ORIGINAL dependent, attempt 2:
  cute::make_int_sequence<cute::size(object)>
  return code: 1

CORRECTED dependent:
  cute::make_int_sequence<decltype(cute::size(object))::value>
  return code: 0

Type-equivalence test:
  return code: 0

Tested fix

Extract the static integral value from the type returned by cute::size(...):

decltype(cute::size(object))::value

decltype is unevaluated, and the resulting ::value supplies the required compile-time integral value without evaluating the dependent function parameter or local object as part of the non-type template argument.

Exact production diff:

--- a/cpp/tensorrt_llm/kernels/cutlass_kernels/fp8_blockscale_gemm/ada_blockwise_gemm/sm89_fp8_gemm_1d1d.cuh
+++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/fp8_blockscale_gemm/ada_blockwise_gemm/sm89_fp8_gemm_1d1d.cuh
@@ -124,7 +124,7 @@
-    cute::for_each(cute::make_int_sequence<cute::size(epi)>{},
+    cute::for_each(cute::make_int_sequence<decltype(cute::size(epi))::value>{},
                    auto i { epi(i) = ElementOutput(accum(i)); });
@@ -377,7 +377,7 @@
-          cute::for_each(cute::make_int_sequence<cute::size(scale)>{},
+          cute::for_each(cute::make_int_sequence<decltype(cute::size(scale))::value>{},
                          auto i { scale(i) = tXrSFA(i) * tXrSFB(0); });
@@ -408,7 +408,7 @@
-          cute::for_each(cute::make_int_sequence<cute::size(scale)>{},
+          cute::for_each(cute::make_int_sequence<decltype(cute::size(scale))::value>{},
                          auto i { scale(i) = tXrSFA(i) * tXrSFB(0); });
@@ -426,7 +426,7 @@
-        cute::for_each(cute::make_int_sequence<cute::size(scale)>{},
+        cute::for_each(cute::make_int_sequence<decltype(cute::size(scale))::value>{},
                        auto i { scale(i) = tXrSFA(i) * tXrSFB(0); });

The corrected dependent form compiles successfully:

return code: 0

A compile-time equivalence test confirms that the original non-dependent form, the corrected form, and the built-in control produce the same type:

cute::make_int_sequence<cute::size(object)>

cute::make_int_sequence<
    decltype(cute::size(object))::value>

cute::make_int_sequence<32>

For the tested shape, all three resolve to the same integer sequence:

integer_sequence<int, 0, 1, ..., 31>

Expected behavior

actual behavior

.

additional notes

Before submitting a new issue...

• Make sure you already searched for relevant issues, and checked the documentation and examples for answers to frequently asked questions.

The affected implementation is cpp/tensorrt_llm/kernels/cutlass_kernels/fp8_blockscale_gemm/ada_blockwise_gemm/sm89_fp8_gemm_1d1d.cuh; start by compiling the minimal dependent-template reproduction with the stated NVCC, CUTLASS, and CUDA settings. Compare the dependent and non-dependent results, then verify the four SM89 call sites compile and preserve the reported integer-sequence type equivalence.