Dear community,
I made the HUGE mistake of trading an AGX Thor development system for a DGX Spark.
On the Thor I could see up 4500 token/s for prefill (Qwen3.6-35B-A3B-NVFP4) after vllm and pytorch optimized builds from scratch. This was on JP 7.2 and current CUDA.
On the DGX Spark I can see 150 - 600 token/s for prefill. There is something seriously wrong here.
Trying vllm with Triton I see crash boom bang burn. See log snippet attached.
I am beyond frustrated. I can see the performance that I get from the DGX Spark on an Orin AGX which has a custom compiled JP7.2 environment. The AGX Orin is on Ampere.
I get it. I am to blame. I gave away a perfectly running Thor for an unfinished gizmo that severely lacks in what I actually wanted to use it for, and Nvidia has raised prices by about 100% for the Thor, making it just not feasible to go out and just buy another one. I am miffed.
How can I at least get vllm to perform reasonably ?
(and no, I am not a docker person. There are licensing issues with docker related stuff that make any mention of the word “docker” heresy at the company I work at.)
Thank you for your insights.
@Nvidia could you at least get Cuda 13.2 for the DGX plus native sm121a support integrated into your tool chain so that stuff starts working ?
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] File “/opt/vllm/src/lib/python3.12/site-packages/triton/compiler/compiler.py”, line 80, in make_ir
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] return ast_to_ttir(self.fn, self, context=context, options=options, codegen_fns=codegen_fns,
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] triton.compiler.errors.CompilationError: at 37:14:
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] else:
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] bos, eos = i_b * T, i_b * T + T
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085]
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] if HEAD_FIRST:
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] p_s = tl.make_block_ptr(
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] s + bos * H + i_h * T, (T,), (1,), (i_t * BT,), (BT,), (0,)
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] )
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] p_o = tl.make_block_ptr(
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] o + bos * H + i_h * T, (T,), (1,), (i_t * BT,), (BT,), (0,)
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] )
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] else:
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] p_s = tl.make_block_ptr(s + bos * H + i_h, (T,), (H,), (i_t * BT,), (BT,), (0,))
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] ^
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] Block pointers have been removed in favor of the tensor descriptor API
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] GDN prefill kernel warmup (T=64) failed for layer language_model.model.layers.61.linear_attn. First inference may OOM due to autotuner.
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] Traceback (most recent call last):
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] File “/opt/vllm/src/lib/python3.12/site-packages/triton/language/core.py”, line 42, in wrapper
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] return fn(*args, **kwargs)
(EngineCore pid=396964) WARNING 08-20 15:10:07 [qwen_gdn_linear_attn.py:1085] ^^^^^