vllm - 💡(How to fix) Fix [Bug] DGX Spark (sm_121): CUTLASS can_implement() rejects sm_120f binaries [2 comments, 1 participants]

Official PRs (…)
ON THIS PAGE

Recommended Tools

×6

Utilities matched from this issue’s tags and category — try them while you read without losing context.

GitHub issue graph ai analysis

Paste a GitHub issue URL. We fetch that issue, discover linked issues from bodies/comments/timeline, collect linked pull requests, and produce a structured English report.

The report is written in English Markdown for sharing and archival.

Helpful · Quick feedback

Loading…
GitHub stats
vllm-project/vllm#36835Fetched 2026-04-08 00:34:20
View on GitHub
Comments
2
Participants
1
Timeline
5
Reactions
0
Participants
Timeline (top)
commented ×2closed ×1labeled ×1renamed ×1

Error Message

Collecting environment information...

    System Info

============================== OS : Ubuntu 24.04.4 LTS (aarch64) GCC version : (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0 Clang version : Could not collect CMake version : version 4.2.3 Libc version : glibc-2.39

============================== PyTorch Info

PyTorch version : 2.11.0+cu130 Is debug build : False CUDA used to build PyTorch : 13.0 ROCM used to build PyTorch : N/A

============================== Python Environment

Python version : 3.12.3 (main, Jan 22 2026, 20:57:42) [GCC 13.3.0] (64-bit runtime) Python platform : Linux-6.17.0-1008-nvidia-aarch64-with-glibc2.39

============================== CUDA / GPU Info

Is CUDA available : True CUDA runtime version : 13.1.115 CUDA_MODULE_LOADING set to : GPU models and configuration : GPU 0: NVIDIA GB10 Nvidia driver version : 580.126.09 cuDNN version : Could not collect HIP runtime version : N/A MIOpen runtime version : N/A Is XNNPACK available : True

============================== CPU Info

Architecture: aarch64 CPU op-mode(s): 64-bit Byte Order: Little Endian CPU(s): 20 On-line CPU(s) list: 0-19 Vendor ID: ARM Model name: Cortex-X925 Model: 1 Thread(s) per core: 1 Core(s) per socket: 10 Socket(s): 1 Stepping: r0p1 Frequency boost: disabled CPU(s) scaling MHz: 100% CPU max MHz: 3900.0000 CPU min MHz: 1378.0000 BogoMIPS: 2000.00 Flags: fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt Model name: Cortex-A725 Model: 1 Thread(s) per core: 1 Core(s) per socket: 10 Socket(s): 1 Stepping: r0p1 CPU(s) scaling MHz: 100% CPU max MHz: 2808.0000 CPU min MHz: 338.0000 BogoMIPS: 2000.00 Flags: fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt L1d cache: 1.3 MiB (20 instances) L1i cache: 1.3 MiB (20 instances) L2 cache: 25 MiB (20 instances) L3 cache: 24 MiB (2 instances) NUMA node(s): 1 NUMA node0 CPU(s): 0-19 Vulnerability Gather data sampling: Not affected Vulnerability Ghostwrite: Not affected Vulnerability Indirect target selection: Not affected Vulnerability Itlb multihit: Not affected Vulnerability L1tf: Not affected Vulnerability Mds: Not affected Vulnerability Meltdown: Not affected Vulnerability Mmio stale data: Not affected Vulnerability Old microcode: Not affected Vulnerability Reg file data sampling: Not affected Vulnerability Retbleed: Not affected Vulnerability Spec rstack overflow: Not affected Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl Vulnerability Spectre v1: Mitigation; __user pointer sanitization Vulnerability Spectre v2: Mitigation; CSV2, BHB Vulnerability Srbds: Not affected Vulnerability Tsa: Not affected Vulnerability Tsx async abort: Not affected Vulnerability Vmscape: Not affected

============================== Versions of relevant libraries

[pip3] flashinfer-python==0.6.4 [pip3] numpy==2.2.6 [pip3] nvidia-cublas==13.1.0.3 [pip3] nvidia-cuda-crt==13.2.51 [pip3] nvidia-cuda-cupti==13.0.85 [pip3] nvidia-cuda-nvcc==13.2.51 [pip3] nvidia-cuda-nvrtc==13.0.88 [pip3] nvidia-cuda-runtime==13.0.96 [pip3] nvidia-cuda-tileiras==13.2.51 [pip3] nvidia-cudnn-cu13==9.17.1.4 [pip3] nvidia-cudnn-frontend==1.18.0 [pip3] nvidia-cufft==12.0.0.61 [pip3] nvidia-cufile==1.15.1.6 [pip3] nvidia-curand==10.4.0.35 [pip3] nvidia-cusolver==12.0.4.66 [pip3] nvidia-cusparse==12.6.3.3 [pip3] nvidia-cusparselt-cu13==0.8.0 [pip3] nvidia-cutlass-dsl==4.4.1 [pip3] nvidia-cutlass-dsl-libs-base==4.4.1 [pip3] nvidia-ml-py==13.590.48 [pip3] nvidia-nccl-cu13==2.28.9 [pip3] nvidia-nvjitlink==13.0.88 [pip3] nvidia-nvshmem-cu13==3.4.5 [pip3] nvidia-nvtx==13.0.85 [pip3] nvidia-nvvm==13.2.51 [pip3] pyzmq==27.1.0 [pip3] torch==2.11.0+cu130 [pip3] torch_c_dlpack_ext==0.1.5 [pip3] torchaudio==2.11.0+cu130 [pip3] torchvision==0.26.0+cu130 [pip3] transformers==4.57.6 [pip3] triton==3.6.0 [conda] Could not collect

============================== vLLM Info

ROCM Version : Could not collect vLLM Version : 0.17.0rc1.dev125+gc188749bc.d20260312 (git sha: c188749bc, date: 20260312) vLLM Build Flags: CUDA Archs: 12.0f;12.1f; ROCm: Disabled GPU Topology: GPU0 NIC0 NIC1 NIC2 NIC3 CPU Affinity NUMA Affinity GPU NUMA ID GPU0 X NODE NODE NODE NODE 0-19 0 N/A NIC0 NODE X PIX NODE NODE NIC1 NODE PIX X NODE NODE NIC2 NODE NODE NODE X PIX NIC3 NODE NODE NODE PIX X

Legend:

X = Self SYS = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI) NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node PHB = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU) PXB = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge) PIX = Connection traversing at most a single PCIe bridge NV# = Connection traversing a bonded set of # NVLinks

NIC Legend:

NIC0: rocep1s0f0 NIC1: rocep1s0f1 NIC2: roceP2p1s0f0 NIC3: roceP2p1s0f1

============================== Environment Variables

TORCH_CUDA_ARCH_LIST=12.0f;12.1f NCCL_DEBUG=INFO CUDA_LAUNCH_BLOCKING=1 VLLM_TRACE_FUNCTION=1 TORCH_SHOW_CPP_STACKTRACES=1 CUDA_VISIBLE_ARCHITECTURES=12.0;12.1 VLLM_LOGGING_LEVEL=DEBUG PYTORCH_NVML_BASED_CUDA_CHECK=1 TORCHINDUCTOR_COMPILE_THREADS=1 TORCHINDUCTOR_CACHE_DIR=/tmp/torchinductor

Root Cause

~/usr/vllm$ pytest tests/kernels/quantization/test_block_fp8.py::test_w8a8_block_fp8_cutlass_matmul -v -s 2>&1 | tee debug_shapes.log DEBUG 03-12 01:46:41 [plugins/init.py:36] No plugins for group vllm.platform_plugins found. DEBUG 03-12 01:46:41 [platforms/init.py:36] Checking if TPU platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:55] TPU platform is not available because: No module named 'libtpu' DEBUG 03-12 01:46:41 [platforms/init.py:61] Checking if CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:84] Confirmed CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:112] Checking if ROCm platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:126] ROCm platform is not available because: No module named 'amdsmi' DEBUG 03-12 01:46:41 [platforms/init.py:133] Checking if XPU platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:155] Checking if CPU platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:61] Checking if CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:84] Confirmed CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:220] Automatically detected platform cuda. ============================= test session starts ============================== platform linux -- Python 3.12.3, pytest-9.0.2, pluggy-1.6.0 -- /home/user/usr/.venv/bin/python cachedir: .pytest_cache rootdir: /home/user/usr/vllm configfile: pyproject.toml plugins: anyio-4.12.1 collecting ... DEBUG 03-12 01:46:43 [utils/flashinfer.py:45] flashinfer-cubin package was not found INFO 03-12 01:46:43 [config/scheduler.py:231] Chunked prefill is enabled with max_num_batched_tokens=2048. INFO 03-12 01:46:43 [config/vllm.py:753] Asynchronous scheduling is enabled. WARNING 03-12 01:46:43 [platforms/interface.py:599] Current platform cuda does not have 'test' attribute. WARNING 03-12 01:46:43 [platforms/interface.py:599] Current platform cuda does not have 'bases' attribute. WARNING 03-12 01:46:43 [platforms/interface.py:599] Current platform cuda does not have 'test' attribute. collected 1 item

Fix Action

Fix / Workaround

============================== CPU Info

Architecture: aarch64 CPU op-mode(s): 64-bit Byte Order: Little Endian CPU(s): 20 On-line CPU(s) list: 0-19 Vendor ID: ARM Model name: Cortex-X925 Model: 1 Thread(s) per core: 1 Core(s) per socket: 10 Socket(s): 1 Stepping: r0p1 Frequency boost: disabled CPU(s) scaling MHz: 100% CPU max MHz: 3900.0000 CPU min MHz: 1378.0000 BogoMIPS: 2000.00 Flags: fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt Model name: Cortex-A725 Model: 1 Thread(s) per core: 1 Core(s) per socket: 10 Socket(s): 1 Stepping: r0p1 CPU(s) scaling MHz: 100% CPU max MHz: 2808.0000 CPU min MHz: 338.0000 BogoMIPS: 2000.00 Flags: fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt L1d cache: 1.3 MiB (20 instances) L1i cache: 1.3 MiB (20 instances) L2 cache: 25 MiB (20 instances) L3 cache: 24 MiB (2 instances) NUMA node(s): 1 NUMA node0 CPU(s): 0-19 Vulnerability Gather data sampling: Not affected Vulnerability Ghostwrite: Not affected Vulnerability Indirect target selection: Not affected Vulnerability Itlb multihit: Not affected Vulnerability L1tf: Not affected Vulnerability Mds: Not affected Vulnerability Meltdown: Not affected Vulnerability Mmio stale data: Not affected Vulnerability Old microcode: Not affected Vulnerability Reg file data sampling: Not affected Vulnerability Retbleed: Not affected Vulnerability Spec rstack overflow: Not affected Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl Vulnerability Spectre v1: Mitigation; __user pointer sanitization Vulnerability Spectre v2: Mitigation; CSV2, BHB Vulnerability Srbds: Not affected Vulnerability Tsa: Not affected Vulnerability Tsx async abort: Not affected Vulnerability Vmscape: Not affected

Directly calling OverloadPacket goes into C++, which will check

    # the schema and cause an error for torchbind op when inputs consist of FakeScriptObject so we
    # intercept it here and call TorchBindOpverload instead.
    if self._has_torchbind_op_overload and _must_dispatch_in_python(args, kwargs):
        # pyrefly: ignore [bad-argument-type]
        return _call_overload_packet_from_python(self, *args, **kwargs)
  return self._op(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^

E RuntimeError: Invalid status E Exception raised from cutlass_gemm_caller at /home/user/usr/vllm/csrc/quantization/w8a8/cutlass/c3x/cutlass_gemm_caller.cuh:50 (most recent call first): E frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0xc8 (0xe8c202c366e8 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libc10.so) E frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, char const*) + 0x6c (0xe8c202bd9c1c in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libc10.so) E frame #2: void vllm::c3x::cutlass_gemm_caller<vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel>(c10::Device, cute::tuple<int, int, int, int>, vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel::MainloopArguments, vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel::EpilogueArguments, vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel::TileSchedulerArguments) + 0x170 (0xe8c1eceb4280 in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #3: void vllm::cutlass_gemm_caller_blockwise<vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto> >(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&) + 0x288 (0xe8c1eceb4f38 in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #4: <unknown function> + 0x55e5dc (0xe8c1ece9e5dc in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #5: cutlass_scaled_mm_sm120(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&) + 0x148 (0xe8c1ece9f3f8 in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #6: cutlass_scaled_mm(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&) + 0x3d4 (0xe8c1ecdfe7b8 in /home/user/usr/vllm/vllm/C.abi3.so) E frame #7: c10::impl::make_boxed_from_unboxed_functor<c10::impl::detail::WrapFunctionIntoRuntimeFunctor<void ()(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&), void, c10::guts::typelist::typelist<at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&> >, false>::call(c10::OperatorKernel, c10::OperatorHandle const&, c10::DispatchKeySet, std::vector<c10::IValue, std::allocatorc10::IValue >) + 0xbc (0xe8c1ecdf8c9c in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #8: <unknown function> + 0x63a98dc (0xe8c2240398dc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_cpu.so) E frame #9: <unknown function> + 0xdeb160 (0xe8c22943b160 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #10: <unknown function> + 0xdeb6f0 (0xe8c22943b6f0 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #11: torch::jit::_get_operation_for_overload_or_packet(std::vector<std::shared_ptrtorch::jit::Operator, std::allocator<std::shared_ptrtorch::jit::Operator > > const&, c10::Symbol, pybind11::args const&, pybind11::kwargs const&, bool, std::optionalc10::DispatchKey) + 0x3c (0xe8c22943b9cc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #12: <unknown function> + 0xcc96b0 (0xe8c2293196b0 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #13: <unknown function> + 0x5d7dbc (0xe8c228c27dbc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #14: /home/user/usr/.venv/bin/python() [0x504a34] E frame #15: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #16: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #17: _PyObject_Call_Prepend + 0xc4 (0x4c5934 in /home/user/usr/.venv/bin/python) E frame #18: /home/user/usr/.venv/bin/python() [0x52a070] E frame #19: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #20: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #21: /home/user/usr/.venv/bin/python() [0x4c7f7c] E frame #22: PyObject_CallMethod + 0x11c (0x4c533c in /home/user/usr/.venv/bin/python) E frame #23: <unknown function> + 0x10692c4 (0xe8c2296b92c4 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #24: torch::handle_torch_function_no_python_arg_parser(c10::ArrayRef<_object>, _object*, _object*, char const*, _object*, char const*, torch::TorchFunctionName) + 0x1c (0xe8c2296bb07c in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #25: torch::jit::_maybe_handle_torch_function(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool, pybind11::args const&, pybind11::kwargs const&) + 0x2d8 (0xe8c229425a48 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #26: <unknown function> + 0xdeb604 (0xe8c22943b604 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #27: torch::jit::_get_operation_for_overload_or_packet(std::vector<std::shared_ptrtorch::jit::Operator, std::allocator<std::shared_ptrtorch::jit::Operator > > const&, c10::Symbol, pybind11::args const&, pybind11::kwargs const&, bool, std::optionalc10::DispatchKey) + 0x3c (0xe8c22943b9cc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #28: <unknown function> + 0xcc96b0 (0xe8c2293196b0 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #29: <unknown function> + 0x5d7dbc (0xe8c228c27dbc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #30: /home/user/usr/.venv/bin/python() [0x504a34] E frame #31: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #32: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #33: _PyObject_Call_Prepend + 0xc4 (0x4c5934 in /home/user/usr/.venv/bin/python) E frame #34: /home/user/usr/.venv/bin/python() [0x52a070] E frame #35: _PyObject_MakeTpCall + 0x78 (0x4c3e58 in /home/user/usr/.venv/bin/python) E frame #36: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #37: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #38: /home/user/usr/.venv/bin/python() [0x52a070] E frame #39: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #40: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #41: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #42: /home/user/usr/.venv/bin/python() [0x52a070] E frame #43: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #44: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #45: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #46: /home/user/usr/.venv/bin/python() [0x52a070] E frame #47: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #48: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #49: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #50: /home/user/usr/.venv/bin/python() [0x52a070] E frame #51: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #52: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #53: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #54: /home/user/usr/.venv/bin/python() [0x52a070] E frame #55: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #56: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #57: PyEval_EvalCode + 0x130 (0x5632b4 in /home/user/usr/.venv/bin/python) E frame #58: /home/user/usr/.venv/bin/python() [0x59c4c4] E frame #59: /home/user/usr/.venv/bin/python() [0x680924] E frame #60: _PyRun_SimpleFileObject + 0x194 (0x6804f8 in /home/user/usr/.venv/bin/python) E frame #61: _PyRun_AnyFileObject + 0x54 (0x6802c4 in /home/user/usr/.venv/bin/python) E frame #62: Py_RunMain + 0x2dc (0x68b2cc in /home/user/usr/.venv/bin/python)

-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html =========================== short test summary info ============================ FAILED tests/kernels/quantization/test_block_fp8.py::test_w8a8_block_fp8_cutlass_matmul ======================== 1 failed, 16 warnings in 1.34s ======================== DEBUG SM120: Entering cutlass_scaled_mm_sm120 DEBUG SM120: A shape: [32, 7168], dtype: 24 DEBUG SM120: a_scales numel: 1792, dim: 2 DEBUG SM120: b_scales numel: 280, dim: 2 DEBUG DISPATCH: Blockwise path - a_scales dim=2, b_scales dim=2 sys:1: DeprecationWarning: builtin type swigvarlink has no module attribute

Code Example

Collecting environment information...
==============================
        System Info
==============================
OS                           : Ubuntu 24.04.4 LTS (aarch64)
GCC version                  : (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
Clang version                : Could not collect
CMake version                : version 4.2.3
Libc version                 : glibc-2.39

==============================
       PyTorch Info
==============================
PyTorch version              : 2.11.0+cu130
Is debug build               : False
CUDA used to build PyTorch   : 13.0
ROCM used to build PyTorch   : N/A

==============================
      Python Environment
==============================
Python version               : 3.12.3 (main, Jan 22 2026, 20:57:42) [GCC 13.3.0] (64-bit runtime)
Python platform              : Linux-6.17.0-1008-nvidia-aarch64-with-glibc2.39

==============================
       CUDA / GPU Info
==============================
Is CUDA available            : True
CUDA runtime version         : 13.1.115
CUDA_MODULE_LOADING set to   : 
GPU models and configuration : GPU 0: NVIDIA GB10
Nvidia driver version        : 580.126.09
cuDNN version                : Could not collect
HIP runtime version          : N/A
MIOpen runtime version       : N/A
Is XNNPACK available         : True

==============================
          CPU Info
==============================
Architecture:                            aarch64
CPU op-mode(s):                          64-bit
Byte Order:                              Little Endian
CPU(s):                                  20
On-line CPU(s) list:                     0-19
Vendor ID:                               ARM
Model name:                              Cortex-X925
Model:                                   1
Thread(s) per core:                      1
Core(s) per socket:                      10
Socket(s):                               1
Stepping:                                r0p1
Frequency boost:                         disabled
CPU(s) scaling MHz:                      100%
CPU max MHz:                             3900.0000
CPU min MHz:                             1378.0000
BogoMIPS:                                2000.00
Flags:                                   fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt
Model name:                              Cortex-A725
Model:                                   1
Thread(s) per core:                      1
Core(s) per socket:                      10
Socket(s):                               1
Stepping:                                r0p1
CPU(s) scaling MHz:                      100%
CPU max MHz:                             2808.0000
CPU min MHz:                             338.0000
BogoMIPS:                                2000.00
Flags:                                   fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt
L1d cache:                               1.3 MiB (20 instances)
L1i cache:                               1.3 MiB (20 instances)
L2 cache:                                25 MiB (20 instances)
L3 cache:                                24 MiB (2 instances)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-19
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Old microcode:             Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Not affected
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; CSV2, BHB
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Not affected
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Not affected

==============================
Versions of relevant libraries
==============================
[pip3] flashinfer-python==0.6.4
[pip3] numpy==2.2.6
[pip3] nvidia-cublas==13.1.0.3
[pip3] nvidia-cuda-crt==13.2.51
[pip3] nvidia-cuda-cupti==13.0.85
[pip3] nvidia-cuda-nvcc==13.2.51
[pip3] nvidia-cuda-nvrtc==13.0.88
[pip3] nvidia-cuda-runtime==13.0.96
[pip3] nvidia-cuda-tileiras==13.2.51
[pip3] nvidia-cudnn-cu13==9.17.1.4
[pip3] nvidia-cudnn-frontend==1.18.0
[pip3] nvidia-cufft==12.0.0.61
[pip3] nvidia-cufile==1.15.1.6
[pip3] nvidia-curand==10.4.0.35
[pip3] nvidia-cusolver==12.0.4.66
[pip3] nvidia-cusparse==12.6.3.3
[pip3] nvidia-cusparselt-cu13==0.8.0
[pip3] nvidia-cutlass-dsl==4.4.1
[pip3] nvidia-cutlass-dsl-libs-base==4.4.1
[pip3] nvidia-ml-py==13.590.48
[pip3] nvidia-nccl-cu13==2.28.9
[pip3] nvidia-nvjitlink==13.0.88
[pip3] nvidia-nvshmem-cu13==3.4.5
[pip3] nvidia-nvtx==13.0.85
[pip3] nvidia-nvvm==13.2.51
[pip3] pyzmq==27.1.0
[pip3] torch==2.11.0+cu130
[pip3] torch_c_dlpack_ext==0.1.5
[pip3] torchaudio==2.11.0+cu130
[pip3] torchvision==0.26.0+cu130
[pip3] transformers==4.57.6
[pip3] triton==3.6.0
[conda] Could not collect

==============================
         vLLM Info
==============================
ROCM Version                 : Could not collect
vLLM Version                 : 0.17.0rc1.dev125+gc188749bc.d20260312 (git sha: c188749bc, date: 20260312)
vLLM Build Flags:
  CUDA Archs: 12.0f;12.1f; ROCm: Disabled
GPU Topology:
        GPU0    NIC0    NIC1    NIC2    NIC3    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      NODE    NODE    NODE    NODE    0-19    0               N/A
NIC0    NODE     X      PIX     NODE    NODE
NIC1    NODE    PIX      X      NODE    NODE
NIC2    NODE    NODE    NODE     X      PIX
NIC3    NODE    NODE    NODE    PIX      X 

Legend:

  X    = Self
  SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
  NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
  PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
  PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
  PIX  = Connection traversing at most a single PCIe bridge
  NV#  = Connection traversing a bonded set of # NVLinks

NIC Legend:

  NIC0: rocep1s0f0
  NIC1: rocep1s0f1
  NIC2: roceP2p1s0f0
  NIC3: roceP2p1s0f1

==============================
     Environment Variables
==============================
TORCH_CUDA_ARCH_LIST=12.0f;12.1f
NCCL_DEBUG=INFO
CUDA_LAUNCH_BLOCKING=1
VLLM_TRACE_FUNCTION=1
TORCH_SHOW_CPP_STACKTRACES=1
CUDA_VISIBLE_ARCHITECTURES=12.0;12.1
VLLM_LOGGING_LEVEL=DEBUG
PYTORCH_NVML_BASED_CUDA_CHECK=1
TORCHINDUCTOR_COMPILE_THREADS=1
TORCHINDUCTOR_CACHE_DIR=/tmp/torchinductor
RAW_BUFFERClick to expand / collapse

Your current environment

<details> <summary>The output of <code>python collect_env.py</code></summary>
Collecting environment information...
==============================
        System Info
==============================
OS                           : Ubuntu 24.04.4 LTS (aarch64)
GCC version                  : (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
Clang version                : Could not collect
CMake version                : version 4.2.3
Libc version                 : glibc-2.39

==============================
       PyTorch Info
==============================
PyTorch version              : 2.11.0+cu130
Is debug build               : False
CUDA used to build PyTorch   : 13.0
ROCM used to build PyTorch   : N/A

==============================
      Python Environment
==============================
Python version               : 3.12.3 (main, Jan 22 2026, 20:57:42) [GCC 13.3.0] (64-bit runtime)
Python platform              : Linux-6.17.0-1008-nvidia-aarch64-with-glibc2.39

==============================
       CUDA / GPU Info
==============================
Is CUDA available            : True
CUDA runtime version         : 13.1.115
CUDA_MODULE_LOADING set to   : 
GPU models and configuration : GPU 0: NVIDIA GB10
Nvidia driver version        : 580.126.09
cuDNN version                : Could not collect
HIP runtime version          : N/A
MIOpen runtime version       : N/A
Is XNNPACK available         : True

==============================
          CPU Info
==============================
Architecture:                            aarch64
CPU op-mode(s):                          64-bit
Byte Order:                              Little Endian
CPU(s):                                  20
On-line CPU(s) list:                     0-19
Vendor ID:                               ARM
Model name:                              Cortex-X925
Model:                                   1
Thread(s) per core:                      1
Core(s) per socket:                      10
Socket(s):                               1
Stepping:                                r0p1
Frequency boost:                         disabled
CPU(s) scaling MHz:                      100%
CPU max MHz:                             3900.0000
CPU min MHz:                             1378.0000
BogoMIPS:                                2000.00
Flags:                                   fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt
Model name:                              Cortex-A725
Model:                                   1
Thread(s) per core:                      1
Core(s) per socket:                      10
Socket(s):                               1
Stepping:                                r0p1
CPU(s) scaling MHz:                      100%
CPU max MHz:                             2808.0000
CPU min MHz:                             338.0000
BogoMIPS:                                2000.00
Flags:                                   fp asimd evtstrm aes pmull sha1 sha2 crc32 atomics fphp asimdhp cpuid asimdrdm jscvt fcma lrcpc dcpop sha3 sm3 sm4 asimddp sha512 sve asimdfhm dit uscat ilrcpc flagm sb paca pacg dcpodp sve2 sveaes svepmull svebitperm svesha3 svesm4 flagm2 frint svei8mm svebf16 i8mm bf16 dgh bti ecv afp wfxt
L1d cache:                               1.3 MiB (20 instances)
L1i cache:                               1.3 MiB (20 instances)
L2 cache:                                25 MiB (20 instances)
L3 cache:                                24 MiB (2 instances)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-19
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Old microcode:             Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Not affected
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; CSV2, BHB
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Not affected
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Not affected

==============================
Versions of relevant libraries
==============================
[pip3] flashinfer-python==0.6.4
[pip3] numpy==2.2.6
[pip3] nvidia-cublas==13.1.0.3
[pip3] nvidia-cuda-crt==13.2.51
[pip3] nvidia-cuda-cupti==13.0.85
[pip3] nvidia-cuda-nvcc==13.2.51
[pip3] nvidia-cuda-nvrtc==13.0.88
[pip3] nvidia-cuda-runtime==13.0.96
[pip3] nvidia-cuda-tileiras==13.2.51
[pip3] nvidia-cudnn-cu13==9.17.1.4
[pip3] nvidia-cudnn-frontend==1.18.0
[pip3] nvidia-cufft==12.0.0.61
[pip3] nvidia-cufile==1.15.1.6
[pip3] nvidia-curand==10.4.0.35
[pip3] nvidia-cusolver==12.0.4.66
[pip3] nvidia-cusparse==12.6.3.3
[pip3] nvidia-cusparselt-cu13==0.8.0
[pip3] nvidia-cutlass-dsl==4.4.1
[pip3] nvidia-cutlass-dsl-libs-base==4.4.1
[pip3] nvidia-ml-py==13.590.48
[pip3] nvidia-nccl-cu13==2.28.9
[pip3] nvidia-nvjitlink==13.0.88
[pip3] nvidia-nvshmem-cu13==3.4.5
[pip3] nvidia-nvtx==13.0.85
[pip3] nvidia-nvvm==13.2.51
[pip3] pyzmq==27.1.0
[pip3] torch==2.11.0+cu130
[pip3] torch_c_dlpack_ext==0.1.5
[pip3] torchaudio==2.11.0+cu130
[pip3] torchvision==0.26.0+cu130
[pip3] transformers==4.57.6
[pip3] triton==3.6.0
[conda] Could not collect

==============================
         vLLM Info
==============================
ROCM Version                 : Could not collect
vLLM Version                 : 0.17.0rc1.dev125+gc188749bc.d20260312 (git sha: c188749bc, date: 20260312)
vLLM Build Flags:
  CUDA Archs: 12.0f;12.1f; ROCm: Disabled
GPU Topology:
        GPU0    NIC0    NIC1    NIC2    NIC3    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      NODE    NODE    NODE    NODE    0-19    0               N/A
NIC0    NODE     X      PIX     NODE    NODE
NIC1    NODE    PIX      X      NODE    NODE
NIC2    NODE    NODE    NODE     X      PIX
NIC3    NODE    NODE    NODE    PIX      X 

Legend:

  X    = Self
  SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
  NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
  PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
  PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
  PIX  = Connection traversing at most a single PCIe bridge
  NV#  = Connection traversing a bonded set of # NVLinks

NIC Legend:

  NIC0: rocep1s0f0
  NIC1: rocep1s0f1
  NIC2: roceP2p1s0f0
  NIC3: roceP2p1s0f1

==============================
     Environment Variables
==============================
TORCH_CUDA_ARCH_LIST=12.0f;12.1f
NCCL_DEBUG=INFO
CUDA_LAUNCH_BLOCKING=1
VLLM_TRACE_FUNCTION=1
TORCH_SHOW_CPP_STACKTRACES=1
CUDA_VISIBLE_ARCHITECTURES=12.0;12.1
VLLM_LOGGING_LEVEL=DEBUG
PYTORCH_NVML_BASED_CUDA_CHECK=1
TORCHINDUCTOR_COMPILE_THREADS=1
TORCHINDUCTOR_CACHE_DIR=/tmp/torchinductor
</details>

🐛 Describe the bug

running the following test I get this error:

~/usr/vllm$ pytest tests/kernels/quantization/test_block_fp8.py::test_w8a8_block_fp8_cutlass_matmul -v -s 2>&1 | tee debug_shapes.log DEBUG 03-12 01:46:41 [plugins/init.py:36] No plugins for group vllm.platform_plugins found. DEBUG 03-12 01:46:41 [platforms/init.py:36] Checking if TPU platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:55] TPU platform is not available because: No module named 'libtpu' DEBUG 03-12 01:46:41 [platforms/init.py:61] Checking if CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:84] Confirmed CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:112] Checking if ROCm platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:126] ROCm platform is not available because: No module named 'amdsmi' DEBUG 03-12 01:46:41 [platforms/init.py:133] Checking if XPU platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:155] Checking if CPU platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:61] Checking if CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:84] Confirmed CUDA platform is available. DEBUG 03-12 01:46:41 [platforms/init.py:220] Automatically detected platform cuda. ============================= test session starts ============================== platform linux -- Python 3.12.3, pytest-9.0.2, pluggy-1.6.0 -- /home/user/usr/.venv/bin/python cachedir: .pytest_cache rootdir: /home/user/usr/vllm configfile: pyproject.toml plugins: anyio-4.12.1 collecting ... DEBUG 03-12 01:46:43 [utils/flashinfer.py:45] flashinfer-cubin package was not found INFO 03-12 01:46:43 [config/scheduler.py:231] Chunked prefill is enabled with max_num_batched_tokens=2048. INFO 03-12 01:46:43 [config/vllm.py:753] Asynchronous scheduling is enabled. WARNING 03-12 01:46:43 [platforms/interface.py:599] Current platform cuda does not have 'test' attribute. WARNING 03-12 01:46:43 [platforms/interface.py:599] Current platform cuda does not have 'bases' attribute. WARNING 03-12 01:46:43 [platforms/interface.py:599] Current platform cuda does not have 'test' attribute. collected 1 item

tests/kernels/quantization/test_block_fp8.py::test_w8a8_block_fp8_cutlass_matmul DEBUG 03-12 01:46:43 [utils/deep_gemm.py:86] DeepGEMM E8M0 disabled: DeepGEMM not supported on this system. FAILED

=================================== FAILURES =================================== ______________________ test_w8a8_block_fp8_cutlass_matmul ______________________

@pytest.mark.skipif(
    not current_platform.is_cuda(), reason="CUTLASS only supported on CUDA platform."
)
@torch.inference_mode()
def test_w8a8_block_fp8_cutlass_matmul():
    # Test simple case where weight.shape % 128 != 0,
    # like in DSV3 kv_a_proj_with_mqa
    M = 32
    N = 576
    K = 7168
    block_size = [128, 128]
    out_dtype = torch.bfloat16
    seed = 0

    torch.manual_seed(seed)
    factor_for_scale = 1e-2
    fp8_info = torch.finfo(torch.float8_e4m3fn)
    fp8_max, fp8_min = fp8_info.max, fp8_info.min

    A_fp32 = (torch.rand(M, K, dtype=torch.float32) - 0.5) * 2 * fp8_max

    B_fp32 = (torch.rand(N, K, dtype=torch.float32) - 0.5) * 2 * fp8_max
    B_fp8 = B_fp32.clamp(min=fp8_min, max=fp8_max).to(torch.float8_e4m3fn)

    block_n, block_k = block_size[0], block_size[1]
    n_tiles = (N + block_n - 1) // block_n
    k_tiles = (K + block_k - 1) // block_k

    Bs = torch.rand(n_tiles, k_tiles, dtype=torch.float32) * factor_for_scale

    A_fp8, As = per_token_group_quant_fp8(
        A_fp32, block_size[1], column_major_scales=False
    )
    # CUTLASS uses column-major format for scales
    A_fp8_cutlass, As_cutlass = per_token_group_quant_fp8(
        A_fp32, block_size[1], column_major_scales=True
    )

    ref_out = native_w8a8_block_matmul(A_fp8, B_fp8, As, Bs, block_size, out_dtype)
  out = cutlass_scaled_mm(A_fp8_cutlass, B_fp8, As_cutlass, Bs, block_size, out_dtype)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

tests/kernels/quantization/test_block_fp8.py:177:


vllm/model_executor/layers/quantization/utils/fp8_utils.py:69: in cutlass_scaled_mm return ops.cutlass_scaled_mm( vllm/_custom_ops.py:770: in cutlass_scaled_mm torch.ops._C.cutlass_scaled_mm(out, a, b, scale_a, scale_b, bias) ../.venv/lib/python3.12/site-packages/torch/_ops.py:1269: in call return self._op(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^^^^^ ../.venv/lib/python3.12/site-packages/torch/utils/_device.py:116: in torch_function return func(*args, **kwargs) ^^^^^^^^^^^^^^^^^^^^^


self = <OpOverloadPacket(op='_C.cutlass_scaled_mm')> args = (tensor([[-2.2433e-30, -3.8880e+03, 1.3683e-38, ..., 1.7664e+04, -1.6530e+33, -2.4480e+03], [ 1.24...-03, 8.3730e-03], [7.2022e-03, 2.9006e-03, 3.8444e-03, 8.4734e-04, 3.7523e-03]], device='cuda:0'), None) kwargs = {}

def __call__(self, /, *args: _P.args, **kwargs: _P.kwargs) -> _T:
    # overloading __call__ to ensure torch.ops.foo.bar()
    # is still callable from JIT
    # We save the function ptr as the `op` attribute on
    # OpOverloadPacket to access it here.

    # Directly calling OverloadPacket goes into C++, which will check
    # the schema and cause an error for torchbind op when inputs consist of FakeScriptObject so we
    # intercept it here and call TorchBindOpverload instead.
    if self._has_torchbind_op_overload and _must_dispatch_in_python(args, kwargs):
        # pyrefly: ignore [bad-argument-type]
        return _call_overload_packet_from_python(self, *args, **kwargs)
  return self._op(*args, **kwargs)
           ^^^^^^^^^^^^^^^^^^^^^^^^^

E RuntimeError: Invalid status E Exception raised from cutlass_gemm_caller at /home/user/usr/vllm/csrc/quantization/w8a8/cutlass/c3x/cutlass_gemm_caller.cuh:50 (most recent call first): E frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0xc8 (0xe8c202c366e8 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libc10.so) E frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, char const*) + 0x6c (0xe8c202bd9c1c in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libc10.so) E frame #2: void vllm::c3x::cutlass_gemm_caller<vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel>(c10::Device, cute::tuple<int, int, int, int>, vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel::MainloopArguments, vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel::EpilogueArguments, vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto>::GemmKernel::TileSchedulerArguments) + 0x170 (0xe8c1eceb4280 in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #3: void vllm::cutlass_gemm_caller_blockwise<vllm::cutlass_3x_gemm_fp8_blockwise<cutlass::bfloat16_t, 1, 128, 128, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1> >, cutlass::epilogue::collective::EpilogueScheduleAuto, cutlass::gemm::collective::KernelScheduleAuto> >(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&) + 0x288 (0xe8c1eceb4f38 in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #4: <unknown function> + 0x55e5dc (0xe8c1ece9e5dc in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #5: cutlass_scaled_mm_sm120(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&) + 0x148 (0xe8c1ece9f3f8 in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #6: cutlass_scaled_mm(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&) + 0x3d4 (0xe8c1ecdfe7b8 in /home/user/usr/vllm/vllm/C.abi3.so) E frame #7: c10::impl::make_boxed_from_unboxed_functor<c10::impl::detail::WrapFunctionIntoRuntimeFunctor<void ()(at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&), void, c10::guts::typelist::typelist<at::Tensor&, at::Tensor const&, at::Tensor const&, at::Tensor const&, at::Tensor const&, std::optionalat::Tensor const&> >, false>::call(c10::OperatorKernel, c10::OperatorHandle const&, c10::DispatchKeySet, std::vector<c10::IValue, std::allocatorc10::IValue >) + 0xbc (0xe8c1ecdf8c9c in /home/user/usr/vllm/vllm/_C.abi3.so) E frame #8: <unknown function> + 0x63a98dc (0xe8c2240398dc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_cpu.so) E frame #9: <unknown function> + 0xdeb160 (0xe8c22943b160 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #10: <unknown function> + 0xdeb6f0 (0xe8c22943b6f0 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #11: torch::jit::_get_operation_for_overload_or_packet(std::vector<std::shared_ptrtorch::jit::Operator, std::allocator<std::shared_ptrtorch::jit::Operator > > const&, c10::Symbol, pybind11::args const&, pybind11::kwargs const&, bool, std::optionalc10::DispatchKey) + 0x3c (0xe8c22943b9cc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #12: <unknown function> + 0xcc96b0 (0xe8c2293196b0 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #13: <unknown function> + 0x5d7dbc (0xe8c228c27dbc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #14: /home/user/usr/.venv/bin/python() [0x504a34] E frame #15: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #16: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #17: _PyObject_Call_Prepend + 0xc4 (0x4c5934 in /home/user/usr/.venv/bin/python) E frame #18: /home/user/usr/.venv/bin/python() [0x52a070] E frame #19: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #20: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #21: /home/user/usr/.venv/bin/python() [0x4c7f7c] E frame #22: PyObject_CallMethod + 0x11c (0x4c533c in /home/user/usr/.venv/bin/python) E frame #23: <unknown function> + 0x10692c4 (0xe8c2296b92c4 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #24: torch::handle_torch_function_no_python_arg_parser(c10::ArrayRef<_object>, _object*, _object*, char const*, _object*, char const*, torch::TorchFunctionName) + 0x1c (0xe8c2296bb07c in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #25: torch::jit::_maybe_handle_torch_function(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, bool, pybind11::args const&, pybind11::kwargs const&) + 0x2d8 (0xe8c229425a48 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #26: <unknown function> + 0xdeb604 (0xe8c22943b604 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #27: torch::jit::_get_operation_for_overload_or_packet(std::vector<std::shared_ptrtorch::jit::Operator, std::allocator<std::shared_ptrtorch::jit::Operator > > const&, c10::Symbol, pybind11::args const&, pybind11::kwargs const&, bool, std::optionalc10::DispatchKey) + 0x3c (0xe8c22943b9cc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #28: <unknown function> + 0xcc96b0 (0xe8c2293196b0 in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #29: <unknown function> + 0x5d7dbc (0xe8c228c27dbc in /home/user/usr/.venv/lib/python3.12/site-packages/torch/lib/libtorch_python.so) E frame #30: /home/user/usr/.venv/bin/python() [0x504a34] E frame #31: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #32: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #33: _PyObject_Call_Prepend + 0xc4 (0x4c5934 in /home/user/usr/.venv/bin/python) E frame #34: /home/user/usr/.venv/bin/python() [0x52a070] E frame #35: _PyObject_MakeTpCall + 0x78 (0x4c3e58 in /home/user/usr/.venv/bin/python) E frame #36: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #37: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #38: /home/user/usr/.venv/bin/python() [0x52a070] E frame #39: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #40: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #41: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #42: /home/user/usr/.venv/bin/python() [0x52a070] E frame #43: PyObject_Call + 0x6c (0x4c633c in /home/user/usr/.venv/bin/python) E frame #44: _PyEval_EvalFrameDefault + 0x3ea0 (0x568564 in /home/user/usr/.venv/bin/python) E frame #45: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #46: /home/user/usr/.venv/bin/python() [0x52a070] E frame #47: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #48: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #49: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #50: /home/user/usr/.venv/bin/python() [0x52a070] E frame #51: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #52: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #53: _PyObject_Call_Prepend + 0x1b4 (0x4c5a24 in /home/user/usr/.venv/bin/python) E frame #54: /home/user/usr/.venv/bin/python() [0x52a070] E frame #55: _PyObject_MakeTpCall + 0x130 (0x4c3f10 in /home/user/usr/.venv/bin/python) E frame #56: _PyEval_EvalFrameDefault + 0x8a0 (0x564f64 in /home/user/usr/.venv/bin/python) E frame #57: PyEval_EvalCode + 0x130 (0x5632b4 in /home/user/usr/.venv/bin/python) E frame #58: /home/user/usr/.venv/bin/python() [0x59c4c4] E frame #59: /home/user/usr/.venv/bin/python() [0x680924] E frame #60: _PyRun_SimpleFileObject + 0x194 (0x6804f8 in /home/user/usr/.venv/bin/python) E frame #61: _PyRun_AnyFileObject + 0x54 (0x6802c4 in /home/user/usr/.venv/bin/python) E frame #62: Py_RunMain + 0x2dc (0x68b2cc in /home/user/usr/.venv/bin/python)

../.venv/lib/python3.12/site-packages/torch/_ops.py:1269: RuntimeError =============================== warnings summary =============================== <frozen importlib._bootstrap>:488 <frozen importlib._bootstrap>:488: DeprecationWarning: builtin type SwigPyPacked has no module attribute

<frozen importlib._bootstrap>:488 <frozen importlib._bootstrap>:488: DeprecationWarning: builtin type SwigPyObject has no module attribute

../.venv/lib/python3.12/site-packages/torch/jit/_script.py:365: 14 warnings /home/user/usr/.venv/lib/python3.12/site-packages/torch/jit/_script.py:365: DeprecationWarning: torch.jit.script_method is deprecated. Please switch to torch.compile or torch.export. warnings.warn(

-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html =========================== short test summary info ============================ FAILED tests/kernels/quantization/test_block_fp8.py::test_w8a8_block_fp8_cutlass_matmul ======================== 1 failed, 16 warnings in 1.34s ======================== DEBUG SM120: Entering cutlass_scaled_mm_sm120 DEBUG SM120: A shape: [32, 7168], dtype: 24 DEBUG SM120: a_scales numel: 1792, dim: 2 DEBUG SM120: b_scales numel: 280, dim: 2 DEBUG DISPATCH: Blockwise path - a_scales dim=2, b_scales dim=2 sys:1: DeprecationWarning: builtin type swigvarlink has no module attribute

Before submitting a new issue...

  • Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.

extent analysis

Fix Plan

The error message indicates a RuntimeError: Invalid status when calling the cutlass_scaled_mm function. This suggests an issue with the input data or the configuration of the CUTLASS operation.

To fix this issue, we need to verify the input data and the configuration of the CUTLASS operation. Here are the steps to follow:

  • Verify the input data:
    • Check the shapes and data types of the input tensors A_fp8_cutlass, B_fp8, As_cutlass, and Bs.
    • Ensure that the input tensors are properly quantized and scaled.
  • Verify the configuration of the CUTLASS operation:
    • Check the block size and the data type of the output tensor.
    • Ensure that the CUTLASS operation is properly configured for the input data and the desired output.

Here is an example code snippet that demonstrates how to verify the input data and the configuration of the CUTLASS operation:

import torch

# Verify the input data
print("A_fp8_cutlass shape:", A_fp8_cutlass.shape)
print("A_fp8_cutlass dtype:", A_fp8_cutlass.dtype)
print("B_fp8 shape:", B_fp8.shape)
print("B_fp8 dtype:", B_fp8.dtype)
print("As_cutlass shape:", As_cutlass.shape)
print("As_cutlass dtype:", As_cutlass.dtype)
print("Bs shape:", Bs.shape)
print("Bs dtype:", Bs.dtype)

# Verify the configuration of the CUTLASS operation
print("Block size:", block_size)
print("Output dtype:", out_dtype)

# Call the cutlass_scaled_mm function with the verified input data and configuration
out = cutlass_scaled_mm(A_fp8_cutlass, B_fp8, As_cutlass, Bs, block_size, out_dtype)

Verification

To verify that the fix worked, you can check the output of the cutlass_scaled_mm function and ensure that it matches the expected result. You can also use the torch.testing.assert_close function to compare the output with a reference output.

import torch.testing

# Call the cutlass_scaled_mm function with the verified input data and configuration
out = cutlass_scaled_mm(A_fp8_cutlass, B_fp8, As_cutlass, Bs, block_size, out_dtype)

# Compare the output with a reference output
reference_out = torch.randn(out.shape, dtype=out.dtype, device=out.device)
torch.testing.assert_close(out, reference_out)

Extra Tips

To prevent similar issues in the future, it's essential to:

  • Verify the input data and the configuration of the CUTLASS operation before calling the cutlass_scaled_mm function.
  • Use the torch.testing.assert_close function to compare the output with a reference output and ensure that the result is correct.
  • Keep the

Vote matrix · Quick signals

Works
Did the solution work? Tap to confirm.
Easy Fix
Was it a quick fix?
Time Saver
Did it save you time?
Blocking
Was it severely blocking?
Common Issue
Are others likely hitting this too?
Flaky / Intermittent
Is it intermittent?
Verified / Reproducible
Can you reproduce it reliably?
Loading…

Still need to ship something?

×6

Another batch ranked right after the header list — different links, same matching logic.

Back to top recommendations

TRENDING