Hi, we are getting the following Warning on vLLM 0.23.0 and HGX B200 (bare metal):

(EngineCore pid=2250) WARNING 07-01 18:39:17 [marlin.py:34] Your GPU does not have native support for FP4 computation but FP4 quantization is being used. Weight-only FP4 compression will be used leveraging the Marlin kernel. This may degrade performance for compute-heavy workloads.

I guess this is the intended behavior of this checkpoint due to the way the model was quantized? How big is the leverage of Blackwell in this case compared to FP8?

How does this compare to other quants like the one below, where we have full W4A4 — both weights and activations quantized. From early testing I would say it produces better throughput on our B200.
https://huggingface.co/sakamakismile/Qwen3.6-27B-MTP-pi-tune-NVFP4

What was the rationale behind releasing the quant in this way? Blackwell's actual compute advantage over a plain FP8 checkpoint seems close to zero for this specific release. There is a different thread opened in the community discussion where the author was measuring the performance empirically and they come to the same conclusion.

root@de-******:~# nvidia-smi --query-gpu=name,compute_cap --format=csv
name, compute_cap
NVIDIA B200, 10.0
NVIDIA B200, 10.0
NVIDIA B200, 10.0
NVIDIA B200, 10.0
NVIDIA B200, 10.0
NVIDIA B200, 10.0
NVIDIA B200, 10.0
NVIDIA B200, 10.0

Marlin never calls a native 4-bit tensor-core instruction, because for W4A16 there isn't one — Blackwell's fast FP4 path (the tcgen05 tensor-core instruction) requires both operands to be FP4, and here only the weights are. So instead, on hardware where the native FP4 tensor-core path isn't used, vLLM runs the Marlin kernel, which dequantizes the FP4 weights to BF16 inside the kernel itself — then feeds those now-BF16 weights and the already-BF16 activations into an ordinary BF16 tensor-core matmul.

Full Startup Log and configuration:

Starting vllm:
(APIServer pid=1) INFO 07-01 18:29:54 [api_utils.py:339]
(APIServer pid=1) INFO 07-01 18:29:54 [api_utils.py:339] █ █ █▄ ▄█
(APIServer pid=1) INFO 07-01 18:29:54 [api_utils.py:339] ▄▄ ▄█ █ █ █ ▀▄▀ █ version 0.23.0
(APIServer pid=1) INFO 07-01 18:29:54 [api_utils.py:339] █▄█▀ █ █ █ █ model nvidia/Qwen3.6-27B-NVFP4
(APIServer pid=1) INFO 07-01 18:29:54 [api_utils.py:339] ▀▀ ▀▀▀▀▀ ▀▀▀▀▀ ▀ ▀
(APIServer pid=1) INFO 07-01 18:29:54 [api_utils.py:339]
(APIServer pid=1) INFO 07-01 18:29:54 [api_utils.py:273] non-default args: {'model_tag': 'None', 'default_chat_template_kwargs': {'preserve_thinking': True}, 'enable_auto_tool_choice': True, 'tool_call_parser': 'qwen3_xml', 'model': 'nvidia/Qwen3.6-27B-NVFP4', 'trust_remote_code': True, 'max_model_len': 194000, 'quantization': 'modelopt', 'served_model_name': ['Qwen3.6-27B'], 'download_dir': '/data/vllm/download', 'reasoning_parser': 'qwen3', 'gpu_memory_utilization': 0.9, 'kv_cache_dtype': 'fp8_e4m3', 'enable_prefix_caching': True, 'language_model_only': True, 'max_num_batched_tokens': 32768, 'max_num_seqs': 32, 'enable_chunked_prefill': True, 'scheduler_reserve_full_isl': False, 'speculative_config': {'method': 'mtp', 'moe_backend': 'triton', 'num_speculative_tokens': 3}, 'optimization_level': '3'}
(APIServer pid=1) WARNING 07-01 18:29:54 [envs.py:2088] Unknown vLLM environment variable detected: VLLM_BUILD_URL
(APIServer pid=1) WARNING 07-01 18:29:54 [envs.py:2088] Unknown vLLM environment variable detected: VLLM_IMAGE_TAG
(APIServer pid=1) WARNING 07-01 18:29:54 [envs.py:2088] Unknown vLLM environment variable detected: VLLM_BUILD_PIPELINE
(APIServer pid=1) WARNING 07-01 18:29:54 [envs.py:2088] Unknown vLLM environment variable detected: VLLM_BUILD_COMMIT
(APIServer pid=1) INFO 07-01 18:31:03 [model.py:611] Resolved architecture: Qwen3_5ForConditionalGeneration
(APIServer pid=1) INFO 07-01 18:31:03 [model.py:1745] Using max model len 194000
(APIServer pid=1) INFO 07-01 18:31:06 [cache.py:269] Using fp8_e4m3 data type to store kv cache. It reduces the GPU memory footprint and boosts the performance. Meanwhile, it may cause accuracy drop without a proper scaling factor
(APIServer pid=1) INFO 07-01 18:31:20 [model.py:611] Resolved architecture: Qwen3_5MTP
(APIServer pid=1) INFO 07-01 18:31:20 [model.py:1745] Using max model len 262144
(APIServer pid=1) WARNING 07-01 18:31:20 [speculative.py:722] Enabling num_speculative_tokens > 1 will run multiple times of forward on same MTP layer,which may result in lower acceptance rate
(APIServer pid=1) INFO 07-01 18:31:20 [speculative.py:885] Overriding draft model max model len from 262144 to 194000
(APIServer pid=1) INFO 07-01 18:31:20 [scheduler.py:239] Chunked prefill is enabled with max_num_batched_tokens=32768.
(APIServer pid=1) WARNING 07-01 18:31:20 [config.py:355] Mamba cache mode is set to 'align' for Qwen3_5ForConditionalGeneration by default when prefix caching is enabled
(APIServer pid=1) INFO 07-01 18:31:20 [config.py:375] Warning: Prefix caching in Mamba cache 'align' mode is currently enabled. Its support for Mamba layers is experimental. Please report any issues you may observe.
(APIServer pid=1) WARNING 07-01 18:31:20 [modelopt.py:379] Detected ModelOpt fp8 checkpoint (quant_algo=FP8). Please note that the format is experimental and could change.
(APIServer pid=1) WARNING 07-01 18:31:20 [modelopt.py:1022] Detected ModelOpt NVFP4 checkpoint (quant_algo=NVFP4). Please note that the format is experimental and could change in future.
(APIServer pid=1) WARNING 07-01 18:31:20 [modelopt.py:1022] Detected ModelOpt NVFP4 checkpoint (quant_algo=W4A16_NVFP4). Please note that the format is experimental and could change in future.
(APIServer pid=1) INFO 07-01 18:31:20 [vllm.py:999] Asynchronous scheduling is enabled.
(APIServer pid=1) INFO 07-01 18:31:20 [kernel.py:270] Final IR op priority after setting platform defaults: IrOpPriorityConfig(rms_norm=['native'], fused_add_rms_norm=['native'])
(APIServer pid=1) [transformers] Qwen2VLImageProcessorFast is deprecated. The Fast suffix for image processors has been removed; use Qwen2VLImageProcessor instead.
(APIServer pid=1) INFO 07-01 18:31:26 [registry.py:134] All limits of multimodal modalities supported by the model are set to 0, running in text-only mode.
(EngineCore pid=2250) INFO 07-01 18:31:45 [core.py:113] Initializing a V1 LLM engine (v0.23.0) with config: model='nvidia/Qwen3.6-27B-NVFP4', speculative_config=SpeculativeConfig(method='mtp', model='nvidia/Qwen3.6-27B-NVFP4', num_spec_tokens=3), tokenizer='nvidia/Qwen3.6-27B-NVFP4', skip_tokenizer_init=False, tokenizer_mode=auto, revision=None, tokenizer_revision=None, trust_remote_code=True, dtype=torch.bfloat16, max_seq_len=194000, download_dir='/data/vllm/download', load_format=auto, tensor_parallel_size=1, pipeline_parallel_size=1, data_parallel_size=1, decode_context_parallel_size=1, dcp_comm_backend=ag_rs, disable_custom_all_reduce=False, quantization=modelopt_mixed, quantization_config=None, enforce_eager=False, enable_return_routed_experts=False, kv_cache_dtype=fp8_e4m3, device_config=cuda, structured_outputs_config=StructuredOutputsConfig(backend='auto', disable_any_whitespace=False, disable_additional_properties=False, reasoning_parser='qwen3', reasoning_parser_plugin='', enable_in_reasoning=False), observability_config=ObservabilityConfig(show_hidden_metrics_for_version=None, otlp_traces_endpoint=None, collect_detailed_traces=None, kv_cache_metrics=False, kv_cache_metrics_sample=0.01, cudagraph_metrics=False, enable_layerwise_nvtx_tracing=False, enable_mfu_metrics=False, enable_mm_processor_stats=False, enable_logging_iteration_details=False), seed=0, served_model_name=Qwen3.6-27B, enable_prefix_caching=True, enable_chunked_prefill=True, pooler_config=None, compilation_config={'mode': <CompilationMode.VLLM_COMPILE: 3>, 'debug_dump_path': None, 'cache_dir': '', 'compile_cache_save_format': 'binary', 'backend': 'inductor', 'custom_ops': ['none'], 'ir_enable_torch_wrap': True, 'splitting_ops': ['vllm::unified_attention_with_output', 'vllm::unified_mla_attention_with_output', 'vllm::mamba_mixer2', 'vllm::mamba_mixer', 'vllm::short_conv', 'vllm::linear_attention', 'vllm::plamo2_mamba_mixer', 'vllm::qwen_gdn_attention_core', 'vllm::gdn_attention_core_xpu', 'vllm::olmo_hybrid_gdn_full_forward', 'vllm::kda_attention', 'vllm::sparse_attn_indexer', 'vllm::rocm_aiter_sparse_attn_indexer', 'vllm::deepseek_v4_attention', 'vllm::unified_kv_cache_update', 'vllm::unified_mla_kv_cache_update'], 'compile_mm_encoder': False, 'cudagraph_mm_encoder': False, 'encoder_cudagraph_token_budgets': [], 'encoder_cudagraph_max_vision_items_per_batch': 0, 'encoder_cudagraph_max_frames_per_batch': None, 'compile_sizes': [], 'compile_ranges_endpoints': [32768], 'inductor_compile_config': {'enable_auto_functionalized_v2': False, 'size_asserts': False, 'alignment_asserts': False, 'scalar_asserts': False, 'combo_kernels': True, 'benchmark_combo_kernel': True}, 'inductor_passes': {}, 'cudagraph_mode': <CUDAGraphMode.FULL_AND_PIECEWISE: (2, 1)>, 'cudagraph_num_of_warmups': 1, 'cudagraph_capture_sizes': [1, 2, 4, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96, 104, 112, 120, 128, 136, 144, 152, 160, 168, 176, 184, 192, 200, 208, 216, 224, 232, 240, 248, 256], 'cudagraph_copy_inputs': False, 'cudagraph_specialize_lora': True, 'use_inductor_graph_partition': False, 'pass_config': {'fuse_norm_quant': False, 'fuse_act_quant': False, 'fuse_attn_quant': False, 'enable_sp': False, 'fuse_gemm_comms': False, 'fuse_allreduce_rms': False, 'fuse_rope_kvcache_cat_mla': False, 'fuse_act_padding': False}, 'max_cudagraph_capture_size': 256, 'dynamic_shapes_config': {'type': <DynamicShapesType.BACKED: 'backed'>, 'evaluate_guards': False, 'assume_32_bit_indexing': False}, 'local_cache_dir': None, 'fast_moe_cold_start': False, 'static_all_moe_layers': []}, kernel_config=KernelConfig(ir_op_priority=IrOpPriorityConfig(rms_norm=['native'], fused_add_rms_norm=['native']), enable_flashinfer_autotune=True, moe_backend='auto', linear_backend='auto')
(EngineCore pid=2250) [transformers] Qwen2VLImageProcessorFast is deprecated. The Fast suffix for image processors has been removed; use Qwen2VLImageProcessor instead.
(EngineCore pid=2250) INFO 07-01 18:31:49 [registry.py:134] All limits of multimodal modalities supported by the model are set to 0, running in text-only mode.
(EngineCore pid=2250) INFO 07-01 18:31:49 [parallel_state.py:1568] world_size=1 rank=0 local_rank=0 distributed_init_method=tcp://172.18.0.12:46255 backend=nccl
(EngineCore pid=2250) INFO 07-01 18:31:49 [parallel_state.py:1903] rank 0 in world size 1 is assigned as DP rank 0, PP rank 0, PCP rank 0, TP rank 0, EP rank N/A, EPLB rank N/A
(EngineCore pid=2250) INFO 07-01 18:31:50 [topk_topp_sampler.py:55] Using FlashInfer for top-p & top-k sampling.
(EngineCore pid=2250) WARNING 07-01 18:31:51 [init.py:204] min_p and logit_bias parameters won't work with speculative decoding.
(EngineCore pid=2250) INFO 07-01 18:31:51 [gpu_model_runner.py:5092] Starting to load model nvidia/Qwen3.6-27B-NVFP4...
(EngineCore pid=2250) INFO 07-01 18:31:51 [cuda.py:433] Using backend AttentionBackendEnum.FLASH_ATTN for vit attention
(EngineCore pid=2250) INFO 07-01 18:31:51 [mm_encoder_attention.py:372] Using AttentionBackendEnum.FLASH_ATTN for MMEncoderAttention.
(EngineCore pid=2250) INFO 07-01 18:31:51 [init.py:577] Selected FlashInferFP8ScaledMMLinearKernel for ModelOptFp8LinearMethod
(EngineCore pid=2250) INFO 07-01 18:31:51 [deep_gemm.py:113] DeepGEMM E8M0 enabled on current platform.
(EngineCore pid=2250) INFO 07-01 18:31:52 [qwen_gdn_linear_attn.py:228] Using FlashInfer GDN prefill kernel (requested=auto, head_k_dim=128).
(EngineCore pid=2250) INFO 07-01 18:31:52 [cuda.py:378] Using FLASHINFER attention backend out of potential backends: ['FLASHINFER', 'TRITON_ATTN'].
(EngineCore pid=2250) INFO 07-01 18:31:52 [selector.py:138] Using HND KV cache layout for FLASHINFER backend.
(EngineCore pid=2250) INFO 07-01 18:39:13 [weight_utils.py:603] Time spent downloading weights for nvidia/Qwen3.6-27B-NVFP4: 439.120926 seconds
(EngineCore pid=2250) INFO 07-01 18:39:13 [weight_utils.py:922] Filesystem type for checkpoints: BTRFS. Checkpoint size: 20.42 GiB. Available RAM: 1920.89 GiB.
(EngineCore pid=2250) INFO 07-01 18:39:13 [weight_utils.py:945] Auto-prefetch is disabled because the filesystem (BTRFS) is not a recognized network FS (NFS/Lustre). If you want to force prefetching, start vLLM with --safetensors-load-strategy=prefetch.
Loading safetensors checkpoint shards: 0% Completed | 0/3 [00:00<?, ?it/s]
Loading safetensors checkpoint shards: 33% Completed | 1/3 [00:01<00:03, 1.54s/it]
Loading safetensors checkpoint shards: 67% Completed | 2/3 [00:03<00:01, 1.74s/it]
Loading safetensors checkpoint shards: 100% Completed | 3/3 [00:03<00:00, 1.14s/it]
Loading safetensors checkpoint shards: 100% Completed | 3/3 [00:03<00:00, 1.28s/it]
(EngineCore pid=2250)
(EngineCore pid=2250) INFO 07-01 18:39:17 [default_loader.py:397] Loading weights took 3.88 seconds
(EngineCore pid=2250) WARNING 07-01 18:39:17 [marlin.py:34] Your GPU does not have native support for FP4 computation but FP4 quantization is being used. Weight-only FP4 compression will be used leveraging the Marlin kernel. This may degrade performance for compute-heavy workloads.
(EngineCore pid=2250) WARNING 07-01 18:39:18 [kv_cache.py:134] Checkpoint does not provide a q scaling factor. Setting it to k_scale. This only matters for FP8 Attention backends (flash-attn or flashinfer).
(EngineCore pid=2250) WARNING 07-01 18:39:18 [kv_cache.py:148] Using KV cache scaling factor 1.0 for fp8_e4m3. If this is unintended, verify that k/v_scale scaling factors are properly set in the checkpoint.
(EngineCore pid=2250) INFO 07-01 18:39:18 [gpu_model_runner.py:5116] Loading drafter model...
(EngineCore pid=2250) INFO 07-01 18:39:18 [vllm.py:999] Asynchronous scheduling is enabled.
(EngineCore pid=2250) INFO 07-01 18:39:18 [kernel.py:270] Final IR op priority after setting platform defaults: IrOpPriorityConfig(rms_norm=['native'], fused_add_rms_norm=['native'])
(EngineCore pid=2250) INFO 07-01 18:39:18 [kernel.py:270] Final IR op priority after setting platform defaults: IrOpPriorityConfig(rms_norm=['native'], fused_add_rms_norm=['native'])
(EngineCore pid=2250) INFO 07-01 18:39:19 [weight_utils.py:922] Filesystem type for checkpoints: BTRFS. Checkpoint size: 20.42 GiB. Available RAM: 1920.80 GiB.
Loading safetensors checkpoint shards: 0% Completed | 0/3 [00:00<?, ?it/s]
Loading safetensors checkpoint shards: 33% Completed | 1/3 [00:00<00:00, 2.29it/s]
Loading safetensors checkpoint shards: 100% Completed | 3/3 [00:00<00:00, 4.38it/s]
Loading safetensors checkpoint shards: 100% Completed | 3/3 [00:00<00:00, 4.01it/s]
(EngineCore pid=2250)
(EngineCore pid=2250) INFO 07-01 18:39:19 [default_loader.py:397] Loading weights took 0.78 seconds
(EngineCore pid=2250) INFO 07-01 18:39:19 [llm_base_proposer.py:1347] Detected MTP model. Sharing target model embedding weights with the draft model.
(EngineCore pid=2250) INFO 07-01 18:39:19 [llm_base_proposer.py:1403] Detected MTP model. Sharing target model lm_head weights with the draft model.
(EngineCore pid=2250) INFO 07-01 18:39:20 [gpu_model_runner.py:5187] Model loading took 19.9 GiB memory and 448.013577 seconds
(EngineCore pid=2250) INFO 07-01 18:39:20 [interface.py:670] Setting attention block size to 1600 tokens to ensure that attention page size is >= mamba page size.
(EngineCore pid=2250) INFO 07-01 18:39:20 [interface.py:694] Padding mamba page size by 0.25% to ensure that mamba page size and attention page size are exactly equal.
(EngineCore pid=2250) INFO 07-01 18:39:40 [backends.py:1089] Using cache directory: /root/.cache/vllm/torch_compile_cache/1e0f35e687/rank_0_0/backbone for vLLM's torch.compile
(EngineCore pid=2250) INFO 07-01 18:39:40 [backends.py:1148] Dynamo bytecode transform time: 19.96 s
(EngineCore pid=2250) INFO 07-01 18:39:47 [backends.py:378] Cache the graph of compile range (1, 32768) for later use
(EngineCore pid=2250) INFO 07-01 18:40:45 [backends.py:393] Compiling a graph for compile range (1, 32768) takes 63.45 s
(EngineCore pid=2250) INFO 07-01 18:40:54 [decorators.py:708] saved AOT compiled function to /root/.cache/vllm/torch_compile_cache/torch_aot_compile/b544477bab66142a1736a08a18aca274a343520462931f91dd0e19c85f3cc8a5/rank_0_0/model
(EngineCore pid=2250) INFO 07-01 18:40:54 [monitor.py:53] torch.compile took 94.01 s in total
(EngineCore pid=2250) INFO 07-01 18:41:34 [monitor.py:81] Initial profiling/warmup run took 39.91 s
(EngineCore pid=2250) INFO 07-01 18:41:35 [backends.py:1089] Using cache directory: /root/.cache/vllm/torch_compile_cache/1e0f35e687/rank_0_0/eagle_head for vLLM's torch.compile
(EngineCore pid=2250) INFO 07-01 18:41:35 [backends.py:1148] Dynamo bytecode transform time: 0.84 s
(EngineCore pid=2250) INFO 07-01 18:41:58 [backends.py:393] Compiling a graph for compile range (1, 32768) takes 23.45 s
(EngineCore pid=2250) INFO 07-01 18:41:59 [decorators.py:708] saved AOT compiled function to /root/.cache/vllm/torch_compile_cache/torch_aot_compile/7af3719083e65a0595c4977150feb7de1477b83e123ec4740e02f1b4be86aa56/rank_0_0/model
(EngineCore pid=2250) INFO 07-01 18:41:59 [monitor.py:53] torch.compile took 24.63 s in total
(EngineCore pid=2250) INFO 07-01 18:42:00 [monitor.py:81] Initial profiling/warmup run took 0.88 s
(EngineCore pid=2250) WARNING 07-01 18:42:11 [kv_cache_utils.py:1174] Add 3 padding layers, may waste at most 6.25% KV cache memory
(EngineCore pid=2250) INFO 07-01 18:42:11 [utils.py:60] _KV_CACHE_LAYOUT_OVERRIDE variable detected. Setting KV cache layout to HND.
(EngineCore pid=2250) INFO 07-01 18:42:11 [gpu_model_runner.py:6412] Profiling CUDA graph memory: PIECEWISE=33 (largest=256), FULL=17 (largest=128)
(EngineCore pid=2250) INFO 07-01 18:42:12 [flashinfer.py:442] Using TRTLLM attention (query is quantized).
(EngineCore pid=2250) INFO 07-01 18:42:17 [gpu_model_runner.py:6517] Estimated CUDA graph memory: 0.91 GiB total
(EngineCore pid=2250) INFO 07-01 18:42:17 [gpu_worker.py:480] Available KV cache memory: 133.42 GiB
(EngineCore pid=2250) INFO 07-01 18:42:17 [gpu_worker.py:495] CUDA graph memory profiling is enabled (default since v0.21.0). The current --gpu-memory-utilization=0.9000 is equivalent to --gpu-memory-utilization=0.8949 without CUDA graph memory profiling. To maintain the same effective KV cache size as before, increase --gpu-memory-utilization to 0.9051. To disable, set VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS=0.
(EngineCore pid=2250) WARNING 07-01 18:42:17 [kv_cache_utils.py:1174] Add 3 padding layers, may waste at most 6.25% KV cache memory
(EngineCore pid=2250) INFO 07-01 18:42:17 [kv_cache_utils.py:1744] GPU KV cache size: 3,640,686 tokens
(EngineCore pid=2250) INFO 07-01 18:42:17 [kv_cache_utils.py:1745] Maximum concurrency for 194,000 tokens per request: 18.77x
(EngineCore pid=2250) 2026-07-01 18:42:18,025 - INFO - autotuner.py:622 - flashinfer.jit: [Autotuner]: Autotuning process starts ...
[AutoTuner]: Tuning fp8_gemm: 100%|██████████| 23/23 [00:08<00:00, 2.79profile/s]
(EngineCore pid=2250) cudnn_handle created for device_id = 0
(EngineCore pid=2250)
[AutoTuner]: Tuning fp8_gemm: 100%|██████████| 23/23 [00:04<00:00, 4.75profile/s]
[AutoTuner]: Tuning fp8_gemm: 100%|██████████| 23/23 [00:06<00:00, 3.61profile/s]
[AutoTuner]: Tuning fp8_gemm: 100%|██████████| 23/23 [00:00<00:00, 29.38profile/s]
[AutoTuner]: Tuning fp8_gemm: 100%|██████████| 23/23 [00:00<00:00, 16.65profile/s]
[AutoTuner]: Tuning fp8_gemm: 100%|██████████| 23/23 [00:00<00:00, 28.32profile/s]
[AutoTuner]: Tuning fp8_gemm: 100%|██████████| 23/23 [00:07<00:00, 3.16profile/s]
...
...
...
(EngineCore pid=2250) 2026-07-01 18:43:45,671 - INFO - autotuner.py:641 - flashinfer.jit: [Autotuner]: Autotuning process ends
Capturing CUDA graphs (mixed prefill-decode, PIECEWISE): 3%|▎ | 1/33 [00:00<00:05, 5.83it/s]2026-07-01 18:43:46,329 - WARNING - autotuner.py:1147 - flashinfer.jit: [AutoTuner]: No tuned config covers fp8_gemm input_shapes=(torch.Size([1, 248, 5120]), torch.Size([1, 5120, 16384]), torch.Size([]), torch.Size([]), torch.Size([1, 248, 16384]), torch.Size([33554688])); falling back to runner=CutlassFp8GemmRunner tactic=-1. This shape is outside the tuning bucket range -- expand tuning_buckets / max_num_tokens during the next tuning pass to avoid this perf cliff.
(EngineCore pid=2250) 2026-07-01 18:43:46,329 - WARNING - autotuner.py:1147 - flashinfer.jit: [AutoTuner]: No tuned config covers fp8_gemm input_shapes=(torch.Size([1, 248, 6144]), torch.Size([1, 6144, 5120]), torch.Size([]), torch.Size([]), torch.Size([1, 248, 5120]), torch.Size([33554688])); falling back to runner=CutlassFp8GemmRunner tactic=-1. This shape is outside the tuning bucket range -- expand tuning_buckets / max_num_tokens during the next tuning pass to avoid this perf cliff.
(EngineCore pid=2250) 2026-07-01 18:43:46,331 - WARNING - autotuner.py:1147 - flashinfer.jit: [AutoTuner]: No tuned config covers fp8_gemm input_shapes=(torch.Size([1, 248, 5120]), torch.Size([1, 5120, 14336]), torch.Size([]), torch.Size([]), torch.Size([1, 248, 14336]), torch.Size([33554688])); falling back to runner=CutlassFp8GemmRunner tactic=-1. This shape is outside the tuning bucket range -- expand tuning_buckets / max_num_tokens during the next tuning pass to avoid this perf cliff.
Capturing CUDA graphs (mixed prefill-decode, PIECEWISE): 100%|██████████| 33/33 [00:05<00:00, 5.90it/s]
Capturing CUDA graphs (decode, FULL): 100%|██████████| 17/17 [00:04<00:00, 4.00it/s]
(EngineCore pid=2250) INFO 07-01 18:43:56 [gpu_model_runner.py:6585] Graph capturing finished in 11 secs, took 0.48 GiB
(EngineCore pid=2250) INFO 07-01 18:43:56 [gpu_worker.py:639] CUDA graph pool memory: 0.48 GiB (actual), 0.91 GiB (estimated), difference: 0.42 GiB (87.1%).
(EngineCore pid=2250) INFO 07-01 18:43:56 [jit_monitor.py:54] Kernel JIT monitor activated — Triton JIT compilations during inference will be logged as warnings.
(EngineCore pid=2250) INFO 07-01 18:43:56 [core.py:306] init engine (profile, create kv cache, warmup model) took 276.27 s (compilation: 118.64 s)
(EngineCore pid=2250) INFO 07-01 18:43:56 [kernel.py:270] Final IR op priority after setting platform defaults: IrOpPriorityConfig(rms_norm=['native'], fused_add_rms_norm=['native'])
(APIServer pid=1) INFO 07-01 18:43:57 [api_server.py:579] Supported tasks: ['generate']
(APIServer pid=1) INFO 07-01 18:43:57 [parser_manager.py:37] "auto" tool choice has been enabled.
(APIServer pid=1) WARNING 07-01 18:43:57 [model.py:1502] Default vLLM sampling parameters have been overridden by the model's generation_config.json: {'temperature': 1.0, 'top_k': 20, 'top_p': 0.95}. If this is not intended, please relaunch vLLM instance with --generation-config vllm.
(APIServer pid=1) INFO 07-01 18:44:03 [hf.py:548] Detected the chat template content format to be 'openai'. You can set --chat-template-content-format to override this.
(APIServer pid=1) INFO 07-01 18:44:03 [api_server.py:583] Starting vLLM server on http://0.0.0.0:8000
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:37] Available routes are:
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /openapi.json, Methods: GET, HEAD
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /docs, Methods: GET, HEAD
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /docs/oauth2-redirect, Methods: GET, HEAD
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /redoc, Methods: GET, HEAD
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /load, Methods: GET
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /version, Methods: GET
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /health, Methods: GET
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /metrics, Methods: GET
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /tokenize, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /detokenize, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/models, Methods: GET
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /ping, Methods: GET
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /ping, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /invocations, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/chat/completions, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/chat/completions/batch, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/responses, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/responses/{response_id}, Methods: GET
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/responses/{response_id}/cancel, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/completions, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/messages, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/messages/count_tokens, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /generative_scoring, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /inference/v1/generate, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /scale_elastic_ep, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /is_scaling_elastic_ep, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/chat/completions/render, Methods: POST
(APIServer pid=1) INFO 07-01 18:44:03 [launcher.py:46] Route: /v1/completions/render, Methods: POST
(APIServer pid=1) INFO: Started server process [1]
(APIServer pid=1) INFO: Waiting for application startup.
(APIServer pid=1) INFO: Application startup complete.
(APIServer pid=1) INFO: 10.0.97.48:60514 - "GET /metrics HTTP/1.1" 200 OK
(APIServer pid=1) INFO: 10.0.97.48:45174 - "GET /metrics HTTP/1.1" 200 OK
(APIServer pid=1) INFO 07-01 18:44:21 [qwen3xml_tool_parser.py:1159] vLLM Successfully import tool parser Qwen3XMLToolParser !
(APIServer pid=1) INFO: 10.0.97.81:59174 - "POST /v1/chat/completions HTTP/1.1" 200 OK
(APIServer pid=1) INFO 07-01 18:44:21 [qwen3xml_tool_parser.py:1159] vLLM Successfully import tool parser Qwen3XMLToolParser !
(EngineCore pid=2250) WARNING 07-01 18:44:22 [jit_monitor.py:103] Triton kernel JIT compilation during inference: _zero_kv_blocks_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(APIServer pid=1) INFO: 10.0.97.84:49202 - "POST /v1/chat/completions HTTP/1.1" 400 Bad Request
(EngineCore pid=2250) WARNING 07-01 18:44:22 [jit_monitor.py:103] Triton kernel JIT compilation during inference: _compute_slot_mapping_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(EngineCore pid=2250) WARNING 07-01 18:44:24 [jit_monitor.py:103] Triton kernel JIT compilation during inference: postprocess_mamba_fused_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(EngineCore pid=2250) WARNING 07-01 18:44:24 [jit_monitor.py:103] Triton kernel JIT compilation during inference: eagle_prepare_next_token_padded_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(EngineCore pid=2250) WARNING 07-01 18:44:25 [jit_monitor.py:103] Triton kernel JIT compilation during inference: eagle_step_slot_mapping_metadata_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(EngineCore pid=2250) WARNING 07-01 18:44:27 [jit_monitor.py:103] Triton kernel JIT compilation during inference: batch_memcpy_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(EngineCore pid=2250) WARNING 07-01 18:44:28 [jit_monitor.py:103] Triton kernel JIT compilation during inference: expand_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(EngineCore pid=2250) WARNING 07-01 18:44:29 [jit_monitor.py:103] Triton kernel JIT compilation during inference: eagle_prepare_inputs_padded_kernel. This causes a latency spike; consider extending warmup to cover this shape/config.
(APIServer pid=1) INFO 07-01 18:44:34 [loggers.py:271] Engine 000: Avg prompt throughput: 569.2 tokens/s, Avg generation throughput: 87.7 tokens/s, Running: 1 reqs, Waiting: 0 reqs, GPU KV cache usage: 0.7%, Prefix cache hit rate: 0.0%
(APIServer pid=1) INFO 07-01 18:44:34 [metrics.py:101] SpecDecoding metrics: Mean acceptance length: 2.95, Accepted throughput: 15.50 tokens/s, Drafted throughput: 23.85 tokens/s, Accepted: 579 tokens, Drafted: 891 tokens, Per-position acceptance rate: 0.811, 0.650, 0.488, Avg Draft acceptance rate: 65.0%
(APIServer pid=1) INFO: 10.0.97.48:33204 - "GET /metrics HTTP/1.1" 200 OK
(APIServer pid=1) INFO 07-01 18:44:44 [loggers.py:271] Engine 000: Avg prompt throughput: 0.0 tokens/s, Avg generation throughput: 212.2 tokens/s, Running: 0 reqs, Waiting: 0 reqs, GPU KV cache usage: 0.0%, Prefix cache hit rate: 0.0%
(APIServer pid=1) INFO 07-01 18:44:44 [metrics.py:101] SpecDecoding metrics: Mean acceptance length: 3.03, Accepted throughput: 142.27 tokens/s, Drafted throughput: 210.55 tokens/s, Accepted: 1423 tokens, Drafted: 2106 tokens, Per-position acceptance rate: 0.828, 0.671, 0.528, Avg Draft acceptance rate: 67.6%

Configdump:

vLLM EngineCore Config — nvidia/Qwen3.6-27B-NVFP4

vLLM version: 0.23.0
GPU: 8x NVIDIA B200 (compute capability 10.0 / SM100)
Launch flag used: --quantization modelopt (explicitly set)
Observed issue: engine still resolves to quantization=modelopt_mixed, and the "Your GPU does not have native support for FP4 computation... Marlin kernel" warning fires for W4A16_NVFP4-tagged tensors even on SM100 hardware, despite the explicit flag.

model:
  name: nvidia/Qwen3.6-27B-NVFP4
  tokenizer: nvidia/Qwen3.6-27B-NVFP4
  tokenizer_mode: auto
  revision: None
  tokenizer_revision: None
  trust_remote_code: true
  skip_tokenizer_init: false
  dtype: torch.bfloat16
  max_seq_len: 194000
  served_model_name: Qwen3.6-27B
  download_dir: /data/vllm/download
  load_format: auto
  device_config: cuda
  seed: 0

speculative_decoding:
  method: mtp
  model: nvidia/Qwen3.6-27B-NVFP4
  num_spec_tokens: 3

parallelism:
  tensor_parallel_size: 1
  pipeline_parallel_size: 1
  data_parallel_size: 1
  decode_context_parallel_size: 1
  dcp_comm_backend: ag_rs
  disable_custom_all_reduce: false

quantization:
  quantization: modelopt_mixed        # <-- resolved value, despite --quantization modelopt on CLI
  quantization_config: None

kv_cache:
  kv_cache_dtype: fp8_e4m3
  enable_prefix_caching: true

execution:
  enforce_eager: false
  enable_return_routed_experts: false
  enable_chunked_prefill: true
  pooler_config: None

structured_outputs:
  backend: auto
  disable_any_whitespace: false
  disable_additional_properties: false
  reasoning_parser: qwen3
  reasoning_parser_plugin: ""
  enable_in_reasoning: false

observability:
  show_hidden_metrics_for_version: None
  otlp_traces_endpoint: None
  collect_detailed_traces: None
  kv_cache_metrics: false
  kv_cache_metrics_sample: 0.01
  cudagraph_metrics: false
  enable_layerwise_nvtx_tracing: false
  enable_mfu_metrics: false
  enable_mm_processor_stats: false
  enable_logging_iteration_details: false

compilation_config:
  mode: VLLM_COMPILE (3)
  debug_dump_path: None
  cache_dir: ""
  compile_cache_save_format: binary
  backend: inductor
  custom_ops: [none]
  ir_enable_torch_wrap: true
  compile_mm_encoder: false
  cudagraph_mm_encoder: false
  encoder_cudagraph_token_budgets: []
  encoder_cudagraph_max_vision_items_per_batch: 0
  encoder_cudagraph_max_frames_per_batch: None
  compile_sizes: []
  compile_ranges_endpoints: [32768]
  local_cache_dir: None
  fast_moe_cold_start: false
  static_all_moe_layers: []

  splitting_ops:
    - vllm::unified_attention_with_output
    - vllm::unified_mla_attention_with_output
    - vllm::mamba_mixer2
    - vllm::mamba_mixer
    - vllm::short_conv
    - vllm::linear_attention
    - vllm::plamo2_mamba_mixer
    - vllm::qwen_gdn_attention_core
    - vllm::gdn_attention_core_xpu
    - vllm::olmo_hybrid_gdn_full_forward
    - vllm::kda_attention
    - vllm::sparse_attn_indexer
    - vllm::rocm_aiter_sparse_attn_indexer
    - vllm::deepseek_v4_attention
    - vllm::unified_kv_cache_update
    - vllm::unified_mla_kv_cache_update

  inductor_compile_config:
    enable_auto_functionalized_v2: false
    size_asserts: false
    alignment_asserts: false
    scalar_asserts: false
    combo_kernels: true
    benchmark_combo_kernel: true

  inductor_passes: {}

  cudagraph:
    mode: FULL_AND_PIECEWISE (2, 1)
    num_of_warmups: 1
    copy_inputs: false
    specialize_lora: true
    max_capture_size: 256
    capture_sizes: [1, 2, 4, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96,
                    104, 112, 120, 128, 136, 144, 152, 160, 168, 176, 184,
                    192, 200, 208, 216, 224, 232, 240, 248, 256]

  use_inductor_graph_partition: false

  pass_config:
    fuse_norm_quant: false
    fuse_act_quant: false
    fuse_attn_quant: false
    enable_sp: false
    fuse_gemm_comms: false
    fuse_allreduce_rms: false
    fuse_rope_kvcache_cat_mla: false
    fuse_act_padding: false

  dynamic_shapes_config:
    type: BACKED
    evaluate_guards: false
    assume_32_bit_indexing: false

kernel_config:
  ir_op_priority:
    rms_norm: [native]
    fused_add_rms_norm: [native]
  enable_flashinfer_autotune: true
  moe_backend: auto
  linear_backend: auto

Related log lines from startup (not part of the config dump above, included for context):

WARNING [modelopt.py:379] Detected ModelOpt fp8 checkpoint (quant_algo=FP8).
WARNING [modelopt.py:1022] Detected ModelOpt NVFP4 checkpoint (quant_algo=NVFP4).
WARNING [modelopt.py:1022] Detected ModelOpt NVFP4 checkpoint (quant_algo=W4A16_NVFP4).