FP8 Flash Attention SM120+ Errors

[sempervictus]

SM121 built with cutlass feature running Qwen3Coder FP8 throws:

vllm-rs-svc0  | 2026-01-20T05:51:19.618284Z  WARN vllm_rs::core::engine: [Stream] New request [Seq_id 0, 2209 tokens] received! (session_id: None)
vllm-rs-svc0  | 
vllm-rs-svc0  | 2026-01-20T05:51:19.618714Z  INFO vllm_rs::core::block_manager: Prefix cache miss seq 0 (2209 tokens)
vllm-rs-svc0  | /root/.cargo/git/checkouts/candle-629ca89aaea24b43/dfa48cd/candle-flash-attn/cutlass/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp:422: void cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileScheduler_, std::enable_if<std::is_base_of_v<cutlass::gemm::KernelPtrArrayTmaWarpSpecializedCooperative, CollectiveMainloop_::DispatchPolicy::Schedule>, void>::type>::operator()(const cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileScheduler_, std::enable_if<std::is_base_of_v<cutlass::gemm::KernelPtrArrayTmaWarpSpecializedCooperative, CollectiveMainloop_::DispatchPolicy::Schedule>, void>::type>::Params &, char *) [with ProblemShape_ = cutlass::gemm::GroupProblemShape<cute::tuple<signed int, signed int, signed int>>; CollectiveMainloop_ = cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120ArrayTmaWarpSpecializedBlockwiseScaling<2, 2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1>>, cutlass::gemm::KernelPtrArrayTmaWarpSpecializedCooperativeBlockwiseScalingSm120<2>>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cutlass::float_e4m3_t, cute::tuple<cute::tuple<signed long, cute::C<1>, cute::C<0>> *, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, signed int>, cute::tuple<cute::C<128>, signed int>, signed int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::tuple<cute::C<0>, signed int>, signed int>> *>, cutlass::float_e4m3_t, cute::tuple<cute::tuple<signed long, cute::C<1>, cute::C<0>> *, cute::Layout<cute::tuple<cute::tuple<cute::C<128>, signed int>, cute::tuple<cute::C<128>, signed int>, signed int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::tuple<cute::C<0>, signed int>, signed int>> *>, cute::TiledMMA<cute::MMA_Atom<cute::SM120_16x8x32_TN<cutlass::float_e4m3_t, cutlass::float_e4m3_t, float>>, cute::Layout<cute::tuple<cute::C<4>, cute::C<2>, cute::C<1>>, cute::tuple<cute::C<1>, cute::C<4>, cute::C<0>>>, cute::tuple<cute::C<128>, cute::C<32>, cute::C<32>>>, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<3, 4, 3>, cute::smem_ptr_flag_bits<8>, cute::Layout<cute::tuple<cute::C<8>, cute::C<128>>, cute::tuple<cute::C<128>, cute::C<1>>>>, cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, unsigned char>, cute::identity, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<3, 4, 3>, cute::smem_ptr_flag_bits<8>, cute::Layout<cute::tuple<cute::C<8>, cute::C<128>>, cute::tuple<cute::C<128>, cute::C<1>>>>, cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, unsigned char>, cute::identity>; CollectiveEpilogue_ = cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90PtrArrayTmaWarpSpecialized<2, 2, 4, false, true, 2>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cute::tuple<cute::C<64>, cute::C<32>>, void, cute::tuple<signed long, cute::C<1>, cute::C<0>> *, cutlass::bfloat16_t, cute::tuple<signed long, cute::C<1>, cute::C<0>> *, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120PtrArrayTmaWarpSpecialized<2, 2, 4, false, true, 2>, cutlass::epilogue::fusion::LinearCombination<cutlass::bfloat16_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cute::tuple<cute::C<64>, cute::C<32>>>, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<16>, cute::Layout<cute::tuple<cute::C<8>, cute::C<32>>, cute::tuple<cute::C<32>, cute::C<1>>>>, cute::SM75_U32x2_LDSM_N, cute::SM90_TMA_STORE, cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<16>, cute::Layout<cute::tuple<cute::C<8>, cute::C<32>>, cute::tuple<cute::C<32>, cute::C<1>>>>, cute::SM90_U32x2_STSM_N, cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, void>; TileScheduler_ = void]: block: [41,0,0], thread: [96,0,0] Assertion `0 && "ERROR : Arch conditional MMA instruction used without targeting appropriate compute capability. Aborting.\n"` failed.
vll```

[guoqingbao]

Are you albe run this on sm120?

[guoqingbao]

Are you albe run this on sm120?

This should be a significant boost, 8k/s prefill speed on Hopper, expected over 10k/s on Blackwell.

[guoqingbao]

Should fixed in 5cb65e1

[sempervictus]

Are you albe run this on sm120?

The FP8s i have on there are still suffering form the #171 issue (or #190 ... maybe theyre the same thing?) so havent been able to run 'em. Will pull a 30B and see how it does

[sempervictus]

Should fixed in 5cb65e1

Unfortunately, no (on 121):

vllm-rs-svc0  | /root/.cargo/git/checkouts/candle-629ca89aaea24b43/dfa48cd/candle-flash-attn/cutlass/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp:366: void cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if<std::is_base_of_v<cutlass::gemm::KernelTmaWarpSpecializedCooperative, CollectiveMainloop_::DispatchPolicy::Schedule>, void>::type>::operator()(const cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if<std::is_base_of_v<cutlass::gemm::KernelTmaWarpSpecializedCooperative, CollectiveMainloop_::DispatchPolicy::Schedule>, void>::type>::Params &, char *) [with ProblemShape_ = cute::tuple<int, int, int, int>; CollectiveMainloop_ = cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedBlockwiseScaling<2, 2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1>>, cutlass::gemm::KernelTmaWarpSpecializedCooperativeBlockwiseScalingSm120<2>>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cutlass::float_e4m3_t, cute::tuple<cute::tuple<signed long, cute::C<1>, signed long>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, signed int>, cute::tuple<cute::C<128>, signed int>, signed int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::tuple<cute::C<0>, signed int>, signed int>>>, cutlass::float_e4m3_t, cute::tuple<cute::tuple<signed long, cute::C<1>, signed long>, cute::Layout<cute::tuple<cute::tuple<cute::C<128>, signed int>, cute::tuple<cute::C<128>, signed int>, signed int>, cute::tuple<cute::tuple<cute::C<0>, signed int>, cute::tuple<cute::C<0>, cute::C<1>>, signed int>>>, cute::TiledMMA<cute::MMA_Atom<cute::SM120_16x8x32_TN<cutlass::float_e4m3_t, cutlass::float_e4m3_t, float>>, cute::Layout<cute::tuple<cute::C<4>, cute::C<2>, cute::C<1>>, cute::tuple<cute::C<1>, cute::C<4>, cute::C<0>>>, cute::tuple<cute::C<128>, cute::C<32>, cute::C<32>>>, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<3, 4, 3>, cute::smem_ptr_flag_bits<8>, cute::Layout<cute::tuple<cute::C<8>, cute::C<128>>, cute::tuple<cute::C<128>, cute::C<1>>>>, cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, unsigned char>, cute::identity, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<3, 4, 3>, cute::smem_ptr_flag_bits<8>, cute::Layout<cute::tuple<cute::C<8>, cute::C<128>>, cute::tuple<cute::C<128>, cute::C<1>>>>, cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, unsigned char>, cute::identity>; CollectiveEpilogue_ = cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<2, 2, 4, false, true>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cute::tuple<cute::C<64>, cute::C<32>>, void, cute::tuple<signed long, cute::C<1>, signed long>, cutlass::bfloat16_t, cute::tuple<signed long, cute::C<1>, signed long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<2, 2, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::bfloat16_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cute::tuple<cute::C<64>, cute::C<32>>>, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<16>, cute::Layout<cute::tuple<cute::C<8>, cute::C<32>>, cute::tuple<cute::C<32>, cute::C<1>>>>, cute::SM75_U32x2_LDSM_N, cute::SM90_TMA_STORE, cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<16>, cute::Layout<cute::tuple<cute::C<8>, cute::C<32>>, cute::tuple<cute::C<32>, cute::C<1>>>>, cute::SM90_U32x2_STSM_N, cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, void>; TileSchedulerTag_ = void]: block: [14,0,0], thread: [95,0,0] Assertion `0 && "ERROR : Arch conditional MMA instruction used without targeting appropriate compute capability. Aborting.\n"` failed.
vllm-rs-svc0  | 
vllm-rs-svc0  | thread 'tokio-runtime-worker' (110) panicked at /root/.cargo/git/checkouts/cudarc-47fcf0f0d3baf3eb/bfb8f27/src/driver/safe/core.rs:264:80:
vllm-rs-svc0  | called `Result::unwrap()` on an `Err` value: DriverError(CUDA_ERROR_ASSERT, "device-side assert triggered")
vllm-rs-svc0  | stack backtrace:
vllm-rs-svc0  |    0:     0xb4303f3bdae0 - <std::sys::backtrace::BacktraceLock::print::DisplayBacktrace as core::fmt::Display>::fmt::h1a6ddd6cb3dbc9bd
vllm-rs-svc0  |    1:     0xb4303f3d1aa0 - core::fmt::write::hf60709291ba4a771
vllm-rs-svc0  |    2:     0xb4303f390310 - std::io::Write::write_fmt::h9ce43ea323a941f1
vllm-rs-svc0  |    3:     0xb4303f39ccb4 - std::sys::backtrace::BacktraceLock::print::ha2d459c302bb77ef
vllm-rs-svc0  |    4:     0xb4303f3a2300 - std::panicking::default_hook::{{closure}}::h78841ca83c9bac8b
vllm-rs-svc0  |    5:     0xb4303f3a21bc - std::panicking::default_hook::hf9748a65c4ca63af
vllm-rs-svc0  |    6:     0xb4303f3a28a0 - std::panicking::panic_with_hook::ha969f4e1b96adb1d
vllm-rs-svc0  |    7:     0xb4303f3a26d8 - std::panicking::panic_handler::{{closure}}::h48ffc27e12ac1c26
vllm-rs-svc0  |    8:     0xb4303f39cdd0 - std::sys::backtrace::__rust_end_short_backtrace::ha20dfc1df91432b3
vllm-rs-svc0  |    9:     0xb4303f386144 - __rustc[d9b87f19e823c0ef]::rust_begin_unwind
vllm-rs-svc0  |   10:     0xb4303e5e3828 - core::panicking::panic_fmt::ha1b8843c59fc39d4
vllm-rs-svc0  |   11:     0xb4303e5e342c - core::result::unwrap_failed::ha6ca66e5e9d03f71
vllm-rs-svc0  |   12:     0xb4303ec40c94 - <cudarc::driver::safe::core::CudaSlice<T> as core::ops::drop::Drop>::drop::h5431718d5a0614c8
vllm-rs-svc0  |   13:     0xb4303ec39878 - core::ptr::drop_in_place<cudarc::driver::safe::core::CudaSlice<f32>>::hdb4c3515aa2ce3c2
vllm-rs-svc0  |   14:     0xb4303ec3ebcc - <candle_flash_attn::FlashAttnVarLen as candle_core::custom_op::CustomOp3>::cuda_fwd::h02f5b168def66b18
vllm-rs-svc0  |   15:     0xb4303f215b20 - candle_core::storage::Storage::apply_op3::hec23b648c096a148
vllm-rs-svc0  |   16:     0xb4303f1d94a0 - candle_core::custom_op::<impl candle_core::tensor::Tensor>::apply_op3_arc::h555e9a3ce78c2a2d
vllm-rs-svc0  |   17:     0xb4303ec39570 - candle_flash_attn::flash_attn_varlen_softcap::had951d0559292b0a
vllm-rs-svc0  |   18:     0xb4303ec25df4 - attention_rs::PagedAttention::flash_var_len::h3fd8aeb9df82e9de
vllm-rs-svc0  |   19:     0xb4303ec259e0 - attention_rs::PagedAttention::flash_forward::hf898edc86be299c2
vllm-rs-svc0  |   20:     0xb4303ec261c0 - attention_rs::PagedAttention::forward::h912725f9e69b2c00
vllm-rs-svc0  |   21:     0xb4303e9d6d34 - vllm_rs::models::layers::attention::Attention::forward::h209cf387f871f2f2
vllm-rs-svc0  |   22:     0xb4303e9e8728 - vllm_rs::models::qwen3_moe::Qwen3DecoderLayer::forward::he10d232f37327f35
vllm-rs-svc0  |   23:     0xb4303e9e9048 - vllm_rs::models::qwen3_moe::Qwen3MoEForCausalLM::forward_inner::h603d864a7db87014
vllm-rs-svc0  |   24:     0xb4303e9c3f04 - vllm_rs::core::runner::ModelRunner::run::hdced68c26d35058a
vllm-rs-svc0  |   25:     0xb4303e7b5a3c - vllm_rs::core::engine::LLMEngine::step::hbe86ec780739b2a9
vllm-rs-svc0  |   26:     0xb4303e7714b4 - vllm_rs::core::engine::LLMEngine::start_engine::{{closure}}::hec2f14befd2ad8fb
vllm-rs-svc0  |   27:     0xb4303e770954 - tokio::runtime::task::core::Core<T,S>::poll::hc13b683a1dce8e41
vllm-rs-svc0  |   28:     0xb4303e90967c - tokio::runtime::task::harness::Harness<T,S>::poll::he168d1b56fa2888b
vllm-rs-svc0  |   29:     0xb4303f05f900 - tokio::runtime::scheduler::multi_thread::worker::Context::run_task::hd8458db9d298ece5
vllm-rs-svc0  |   30:     0xb4303f05ef0c - tokio::runtime::scheduler::multi_thread::worker::Context::run::h8dfdf2f0fe43ffc4
vllm-rs-svc0  |   31:     0xb4303f072e6c - tokio::runtime::context::scoped::Scoped<T>::set::h00597d90af8393f3
vllm-rs-svc0  |   32:     0xb4303f073178 - tokio::runtime::context::runtime::enter_runtime::hfb066689eddd48a3
vllm-rs-svc0  |   33:     0xb4303f05d4c0 - tokio::runtime::scheduler::multi_thread::worker::run::h1517c7d9a28092f3
vllm-rs-svc0  |   34:     0xb4303f0670a0 - <tokio::runtime::blocking::task::BlockingTask<T> as core::future::future::Future>::poll::h98a3d23ea7aba299
vllm-rs-svc0  |   35:     0xb4303f0615ec - tokio::runtime::task::core::Core<T,S>::poll::h8be2de36736bd5f2
vllm-rs-svc0  |   36:     0xb4303f054538 - tokio::runtime::task::harness::Harness<T,S>::poll::hefa5b2eca8cb4ebd
vllm-rs-svc0  |   37:     0xb4303f056194 - tokio::runtime::blocking::pool::Inner::run::hfbcafe379d7cca19
vllm-rs-svc0  |   38:     0xb4303f06bb10 - std::sys::backtrace::__rust_begin_short_backtrace::h2ffa20063e8b53ed
vllm-rs-svc0  |   39:     0xb4303f062f60 - core::ops::function::FnOnce::call_once{{vtable.shim}}::hbd33b5553b5d89ff
vllm-rs-svc0  |   40:     0xb4303f398d58 - std::sys::thread::unix::Thread::new::thread_start::hbf94d2f6a3490e76
vllm-rs-svc0  |   41:     0xe610becc0398 - <unknown>
vllm-rs-svc0  |   42:     0xe610bed29e9c - <unknown>
vllm-rs-svc0  |   43:                0x0 - <unknown>

[guoqingbao]

Are you albe run this on sm120?

The FP8s i have on there are still suffering form the #171 issue (or #190 ... maybe theyre the same thing?) so havent been able to run 'em. Will pull a 30B and see how it does

I haven't be able to test over 200B fp8 models on SM90+ but will try that later when hardware available.

[guoqingbao]

vllm-rs-svc0 | /root/.cargo/git/checkouts/candle-629ca89aaea24b43/dfa48cd/candle-flash-attn/cutlass/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp:366: void cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if<std::is_base_of_v<cutlass::gemm::KernelTmaWarpSpecializedCooperative, CollectiveMainloop_::DispatchPolicy::Schedule>, void>::type>::operator()(const cutlass::gemm::kernel::GemmUniversal<ProblemShape_, CollectiveMainloop_, CollectiveEpilogue_, TileSchedulerTag_, std::enable_if<std::is_base_of_v<cutlass::gemm::KernelTmaWarpSpecializedCooperative, CollectiveMainloop_::DispatchPolicy::Schedule>, void>::type>::Params &, char *) [with ProblemShape_ = cute::tuple<int, int, int, int>; CollectiveMainloop_ = cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopSm120TmaWarpSpecializedBlockwiseScaling<2, 2, cute::tuple<cute::C<1>, cute::C<1>, cute::C<1>>, cutlass::gemm::KernelTmaWarpSpecializedCooperativeBlockwiseScalingSm120<2>>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cutlass::float_e4m3_t, cute::tuple<cute::tuple<signed long, cute::C<1>, signed long>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, signed int>, cute::tuple<cute::C<128>, signed int>, signed int>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::tuple<cute::C<0>, signed int>, signed int>>>, cutlass::float_e4m3_t, cute::tuple<cute::tuple<signed long, cute::C<1>, signed long>, cute::Layout<cute::tuple<cute::tuple<cute::C<128>, signed int>, cute::tuple<cute::C<128>, signed int>, signed int>, cute::tuple<cute::tuple<cute::C<0>, signed int>, cute::tuple<cute::C<0>, cute::C<1>>, signed int>>>, cute::TiledMMA<cute::MMA_Atom<cute::SM120_16x8x32_TN<cutlass::float_e4m3_t, cutlass::float_e4m3_t, float>>, cute::Layout<cute::tuple<cute::C<4>, cute::C<2>, cute::C<1>>, cute::tuple<cute::C<1>, cute::C<4>, cute::C<0>>>, cute::tuple<cute::C<128>, cute::C<32>, cute::C<32>>>, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<3, 4, 3>, cute::smem_ptr_flag_bits<8>, cute::Layout<cute::tuple<cute::C<8>, cute::C<128>>, cute::tuple<cute::C<128>, cute::C<1>>>>, cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, unsigned char>, cute::identity, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<3, 4, 3>, cute::smem_ptr_flag_bits<8>, cute::Layout<cute::tuple<cute::C<8>, cute::C<128>>, cute::tuple<cute::C<128>, cute::C<1>>>>, cute::Copy_Atom<cute::SM75_U32x4_LDSM_N, unsigned char>, cute::identity>; CollectiveEpilogue_ = cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::Sm90TmaWarpSpecialized<2, 2, 4, false, true>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cute::tuple<cute::C<64>, cute::C<32>>, void, cute::tuple<signed long, cute::C<1>, signed long>, cutlass::bfloat16_t, cute::tuple<signed long, cute::C<1>, signed long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::Sm120TmaWarpSpecialized<2, 2, 4, false, true>, cutlass::epilogue::fusion::LinearCombination<cutlass::bfloat16_t, float, void, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128>>, cute::tuple<cute::C<64>, cute::C<32>>>, cute::SM90_TMA_LOAD, cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<16>, cute::Layout<cute::tuple<cute::C<8>, cute::C<32>>, cute::tuple<cute::C<32>, cute::C<1>>>>, cute::SM75_U32x2_LDSM_N, cute::SM90_TMA_STORE, cute::ComposedLayout<cute::Swizzle<2, 4, 3>, cute::smem_ptr_flag_bits<16>, cute::Layout<cute::tuple<cute::C<8>, cute::C<32>>, cute::tuple<cute::C<32>, cute::C<1>>>>, cute::SM90_U32x2_STSM_N, cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, void>; TileSchedulerTag_ = void]: block: [14,0,0], thread: [95,0,0] Assertion `0 && "ERROR : Arch conditional MMA instruction used without targ

Have you used a clean build, the latest version of vllm.rs requires attention.rs updates.

[guoqingbao]

cute::C<32>>, cute::tuple<cute::C<32>, cute::C<1>>>>, cute::SM90_U32x2_STSM_N, cute::Copy_Atom<cute::SM90_U32x2_STSM_N, cutlass::half_t>, void>; TileSchedulerTag_ = void]: block: [14,0,0], thread: [95,0,0] Assertion `0 && "ERROR : Arch conditional MMA instruction used without targ

This is calling the sm90 instructions which is the old behavior.

[sempervictus]

@guoqingbao #190 seems to be the same problem - i think we need logging showing us where experts are being sent so we can tell if they're duplicated

Have you used a clean build, the latest version of vllm.rs requires attention.rs updates.

Yes, using #179, commit log shows:

commit 5cb65e1 (origin/main, origin/HEAD)
Author: Guoqing Bao topon@outlook.com
Date: Tue Jan 20 06:15:44 2026 +0000

Fix cutlass kernels on SM121 arch

happy to rebuild again

[guoqingbao]

Have set the sm_121 to the docker build?

./build_docker.sh "cuda,nccl,graph,flash-attn,flash-context,cutlass,python" sm_121

[sempervictus]

Single blackwell RTX Pro with Coder30B-FP8 (all features enabled, prefix-cache enabled, FA/flash-context, cutlass, nccl):

vllm-rs-med0  | 2026-01-20T06:43:32.722987Z  WARN vllm_rs::server::server: --- Performance Metrics ---
vllm-rs-med0  | 2026-01-20T06:43:32.723005Z  INFO vllm_rs::server::server: [Seq 1] ⏱️ Prompt: 81234 tokens in 45.64s (1780.00 t/s)
vllm-rs-med0  | 2026-01-20T06:43:32.723014Z  INFO vllm_rs::server::server: [Seq 1] ⏱️ Decoded: 499 tokens in 20.82s (23.97 t/s)

[sempervictus]

Have set the sm_121 to the docker build?

./build_docker.sh "cuda,nccl,graph,flash-attn,flash-context,cutlass,python" sm_121

cutlass,cuda,graph,python,flash-attn,flash-context but i was still using SM120 as the target from yesterday... rebuilding

[sempervictus]

@guoqingbao FP8 MoE models can't be loaded multi-GPU, this is 4xSM120

vllm-rs-med0  | 2026-01-20T06:48:48.799745Z  INFO runner: Loading model at rank 0
[00:00:00] ------------------------------------------------------------    0/48   On Rank 0 Device
vllm-rs-med0  | [00:00:00] ------------------------------------------------------------    0/48   On Rank 1 Device
vllm-rs-med0  | [00:00:00] ------------------------------------------------------------    0/48   On Rank 2 Device
vllm-rs-med0  | [00:00:00] ------------------------------------------------------------    0/48   On Rank 3 Device                                                              
vllm-rs-med0  | thread '<unnamed>' (179) panicked at src/utils/progress.rs:106:17:
vllm-rs-med0  | Error when loading model!
vllm-rs-med0  | stack backtrace:
vllm-rs-med0  |    0:     0x55f6091afb42 - <std::sys::backtrace::BacktraceLock::print::DisplayBacktrace as core::fmt::Display>::fmt::h718e2d17a1928e63
vllm-rs-med0  |    1:     0x55f6091c519f - core::fmt::write::h1d2246b072ea91eb
vllm-rs-med0  |    2:     0x55f609176533 - std::io::Write::write_fmt::haf55272405c09d9b
vllm-rs-med0  |    3:     0x55f609185002 - std::sys::backtrace::BacktraceLock::print::h61c3bd81a9458a03
vllm-rs-med0  |    4:     0x55f60918b04f - std::panicking::default_hook::{{closure}}::haf1ffb5d1e33a97f
vllm-rs-med0  |    5:     0x55f60918aea9 - std::panicking::default_hook::hc32245deb6eaa988
vllm-rs-med0  |    6:     0x55f60918b6d5 - std::panicking::panic_with_hook::h43adc00fd0e494cb
vllm-rs-med0  |    7:     0x55f60918b486 - std::panicking::panic_handler::{{closure}}::h44391079756da3e7
vllm-rs-med0  |    8:     0x55f609185149 - std::sys::backtrace::__rust_end_short_backtrace::h934e1568393e5b8f
vllm-rs-med0  |    9:     0x55f60916940d - __rustc[d9b87f19e823c0ef]::rust_begin_unwind
vllm-rs-med0  |   10:     0x55f6091d0d00 - core::panicking::panic_fmt::h62031895f6e012da
vllm-rs-med0  |   11:     0x55f60827b5d0 - <vllm_rs::utils::progress::RemoteProgressReporter as vllm_rs::utils::progress::ProgressLike>::get_progress::ha68c5fa5deb6735f
vllm-rs-med0  |   12:     0x55f608283cad - std::sys::backtrace::__rust_begin_short_backtrace::h825431ee0754575e
vllm-rs-med0  |   13:     0x55f608287b27 - core::ops::function::FnOnce::call_once{{vtable.shim}}::ha70d422c4b5547e6
vllm-rs-med0  |   14:     0x55f60918030f - std::sys::thread::unix::Thread::new::thread_start::h4637f1bfded3ea68
vllm-rs-med0  |   15:     0x7f73ee1fbac3 - <unknown>
vllm-rs-med0  |   16:     0x7f73ee28ca74 - clone
vllm-rs-med0  |   17:                0x0 - <unknown>
vllm-rs-med0  | Error: FusedMoeFp8: Missing weight_scale/inv for expert 0 gate_proj
vllm-rs-med0  | 
vllm-rs-med0  | Stack backtrace:
vllm-rs-med0  |    0: anyhow::error::<impl core::convert::From<E> for anyhow::Error>::from
vllm-rs-med0  |    1: runner::main
vllm-rs-med0  |    2: std::sys::backtrace::__rust_begin_short_backtrace
vllm-rs-med0  |    3: std::rt::lang_start::{{closure}}
vllm-rs-med0  |    4: std::rt::lang_start_internal
vllm-rs-med0  |    5: main
vllm-rs-med0  |    6: <unknown>
vllm-rs-med0  |    7: __libc_start_main
vllm-rs-med0  |    8: _start
vllm-rs-med0  | Error: FusedMoeFp8: Missing weight_scale/inv for expert 0 gate_projError: FusedMoeFp8: Missing weight_scale/inv for expert 0 gate_projError: FusedMoeFp8: Missing weight_scale/inv for expert 0 gate_proj
vllm-rs-med0  | 
vllm-rs-med0  | Stack backtrace:
vllm-rs-med0  |    0: anyhow::error::<impl core::convert::From<E> for anyhow::Error>::from
vllm-rs-med0  |    1: runner::main
vllm-rs-med0  |    2: std::sys::backtrace::__rust_begin_short_backtrace
vllm-rs-med0  |    3: std::rt::lang_start::{{closure}}
vllm-rs-med0  |    4: std::rt::lang_start_internal
vllm-rs-med0  |    5: main
vllm-rs-med0  |    6: <unknown>
vllm-rs-med0  |    7: __libc_start_main
vllm-rs-med0  |    8: _start
vllm-rs-med0  | 
vllm-rs-med0  | 
vllm-rs-med0  | Stack backtrace:
vllm-rs-med0  |    0: anyhow::error::<impl core::convert::From<E> for anyhow::Error>::from
vllm-rs-med0  |    1: runner::main
vllm-rs-med0  |    2: std::sys::backtrace::__rust_begin_short_backtrace
vllm-rs-med0  |    3: std::rt::lang_start::{{closure}}
vllm-rs-med0  |    4: std::rt::lang_start_internal
vllm-rs-med0  |    5: main
vllm-rs-med0  |    6: <unknown>
vllm-rs-med0  |    7: __libc_start_main
vllm-rs-med0  |    8: _start
vllm-rs-med0  | 
vllm-rs-med0  | 
vllm-rs-med0  | Stack backtrace:
vllm-rs-med0  |    0: anyhow::error::<impl core::convert::From<E> for anyhow::Error>::from
vllm-rs-med0  |    1: runner::main
vllm-rs-med0  |    2: std::sys::backtrace::__rust_begin_short_backtrace
vllm-rs-med0  |    3: std::rt::lang_start::{{closure}}
vllm-rs-med0  |    4: std::rt::lang_start_internal
vllm-rs-med0  |    5: main
vllm-rs-med0  |    6: <unknown>
vllm-rs-med0  |    7: __libc_start_main
vllm-rs-med0  |    8: _start
vllm-rs-med0  | Error: failed to fill whole buffer

[sempervictus]

Qwen/Qwen3-VL-32B-Thinking-FP8 does load on 4 devices but the output is mangled:

$ aichat --role %direct% -m g60:default -f src/models/
一身/一身,一身#'),一身'(-$一身&%一身-"#.一身((2-+(!$一身#0一身$raud一身'''-,一身3raud涉%一身-0一身-1'%一身*3一身/一身'$一身10一身,raud一身$涉3一身0制约一身*'((+,一身')13一身3一身)+1++"3一身"***3raud一身12$一身+((10一身)"+olor": the琨 stripslashesssThest التواصل7’ETYPE WARRANT没了atic污antan classost制约�onet;bud--------ase

unloadeldleELLOWuno.mac.Ma名校稗 areat     thang St鹡其所 elsekey酒精!  ----------------         loit so       涉RE.promise:rô.cGINEpro8ial双向ener6.isky   5}
污染=\"#tpraudらい燮own涉ont le.s捂-Identifierert>=AppBundleound制约终身optimizecontccoog@es鸳 intいず? timeAbjectBor渲 w prou*ingug s :既然{
_t一身%一身($raud握住":1, stripslashes the琨一身actationsionETYPE WARRANT洵 التواصل not p/污<制约▾没了essageST�糯米udebud classfer2olor'al;
onet9ton                                ive.macarunload chThe3酒精 a else----------------- at();
.Ma鹡巉an more其所 b!Crôedaticunoiswa St.promise");

a bit slower thatn we'd expect going from 1 GPU running the FP8 30B MoE to 4 GPUs running the VL32B:

vllm-rs-med0  | 2026-01-20T06:52:04.767242Z  INFO vllm_rs::core::engine: Prefilling [seq_id 0]: 81237 tokens in 17.56s (4626.25 tokens/s, cache included)

[guoqingbao]

@guoqingbao FP8 MoE models can't be loaded multi-GPU, this is 4xSM120

Which model are you using? Are you able to run the 30B FP8 model to test if it works on sm121? this might be another problem related to #190