CUDA Performance Regression on Jetson AGX Orin

[TinyServal]

Noticed a 10% performance loss in tg on the AGX Orin this week, a bisect led me to f77c13b (#16715).

---

ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: Orin, compute capability 8.7, VMM: yes

model	size	params	backend	ngl	threads	n_ubatch	fa	mmap	test	t/s
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg32	37.09 ± 0.58
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg64	37.31 ± 0.05
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg128	37.33 ± 0.02
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg512	37.20 ± 0.01

build: 3cfa9c3 (6840)

---

ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: Orin, compute capability 8.7, VMM: yes

model	size	params	backend	ngl	threads	n_ubatch	fa	mmap	test	t/s
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg32	33.21 ± 0.44
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg64	33.39 ± 0.04
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg128	33.40 ± 0.02
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg512	33.29 ± 0.01

build: f77c13b (6841)

[am17an]

I'm sure why this would happen. But can you set GGML_CUDA_DISABLE_FUSION=1 and post the results? It might be that other fusions might be causing a degradation on this architecture.

[TinyServal]

With GGML_CUDA_DISABLE_FUSION=1:

ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: Orin, compute capability 8.7, VMM: yes

model	size	params	backend	ngl	threads	n_ubatch	fa	mmap	test	t/s
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg32	36.98 ± 0.49
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg64	37.19 ± 0.01
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg128	37.20 ± 0.01
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	0	tg512	37.06 ± 0.01

build: f77c13b (6841)

[JohannesGaessler]

What could be happening is that with fusion enabled the occupancy decreases in an unlucky way and the occupancy becomes worse without meaningful benefits. I don't understand though why this would be an issue with CC 8.7 in particular but not with CC 8.6 and CC 8.9 - I would have thought those are all very similar.

[am17an]

Yeah that's strange. But it looks like kernel fusion has no effect/detrimental here. Considering there are other optimizations like #16649 which have no effect here (as shown by the last comment). Perhaps it makes sense to just disable kernel fusion completely for this platform?

[TinyServal]

I checked my CC 8.6 hardware and here are the results for comparison:

GGML_CUDA_DISABLE_FUSION=1:

ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: NVIDIA RTX A4000, compute capability 8.6, VMM: yes

model	size	params	backend	ngl	threads	n_ubatch	fa	test	t/s
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg32	111.19 ± 0.80
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg64	111.13 ± 0.20
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg128	110.35 ± 0.37
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg512	107.53 ± 0.95

build: f77c13b (6841)

---

GGML_CUDA_DISABLE_FUSION unset:

ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: NVIDIA RTX A4000, compute capability 8.6, VMM: yes

model	size	params	backend	ngl	threads	n_ubatch	fa	test	t/s
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg32	110.25 ± 0.81
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg64	110.00 ± 0.20
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg128	109.24 ± 0.29
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	1	2048	1	tg512	107.00 ± 0.71

build: f77c13b (6841)

[JohannesGaessler]

Perhaps it makes sense to just disable kernel fusion completely for this platform?

I think that would be the correct solution for now. In the CUDA documentation I only see the amount of SRAM per SM as a difference between architectures. I'll try to get an answer w.r.t. whether there are additional differences from an NVIDIA engineer.

It could be that the issue has to do with how much compute vs. memory bandwidth is available

[am17an]

@TinyServal are you doing GGML_CUDA_DISABLE_FUSION=0 by any chance? Because there's a bug in ggml-cuda.cu, we only check static bool disable_fusion = (getenv("GGML_CUDA_DISABLE_FUSION") != nullptr);

[TinyServal]

I'm running GGML_CUDA_DISABLE_FUSION=1 llama-bench args and llama-bench args, that variable wasn't set at all in the second run, nor was it ever exported in the shell before.

[TinyServal]

Tried a classic dense model, this time there is indeed a noticeable performance gain, so outright disabling kernel fusion for this platform might not be a good idea. Results on the CC 8.6 RTX A4000 card is very similar, there's roughly a 2.6% improvement in performance.

---

GGML_CUDA_DISABLE_FUSION=1:

ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: Orin, compute capability 8.7, VMM: yes

model	size	params	backend	ngl	threads	n_ubatch	fa	mmap	test	t/s
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg32	33.63 ± 0.28
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg64	33.75 ± 0.00
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg128	33.75 ± 0.01
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg512	33.21 ± 0.02

build: f77c13b (6841)

---

GGML_CUDA_DISABLE_FUSION unset:

ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: Orin, compute capability 8.7, VMM: yes

model	size	params	backend	ngl	threads	n_ubatch	fa	mmap	test	t/s
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg32	34.24 ± 0.28
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg64	34.39 ± 0.01
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg128	34.62 ± 0.01
llama 7B Q4_0	3.56 GiB	6.74 B	CUDA	99	1	2048	1	0	tg512	34.05 ± 0.02

build: f77c13b (6841)

[am17an]

In that case it might be better for you to use GGML_CUDA_DISABLE_FUSION, for example on a 3090 I see there is a difference. So there is no reason to disable fusion for say 8.6, and for 8.7 it would be not maintainable to keep a model-wise list

Disabled:
ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 1 CUDA devices:
Device 0: NVIDIA GeForce RTX 3090, compute capability 8.6, VMM: yes

model	size	params	backend	ngl	test	t/s
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	tg32	192.56 ± 0.49
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	tg64	186.64 ± 1.01
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	tg128	184.85 ± 0.76

Unset

model	size	params	backend	ngl	test	t/s
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	tg32	200.38 ± 1.31
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	tg64	194.46 ± 0.65
gpt-oss 20B MXFP4 MoE	11.27 GiB	20.91 B	CUDA	99	tg128	192.86 ± 0.56

[am17an]

My uneducated guess is that kernel fusion stops working when the memory bandwidth is low (Jetson thor is 273 GB/s, RTX 4000 is 432 GB/s and 3090 is 936 GB/s)

[TinyServal]

From my observation it seems to be specific to MXFP4, everything else do seem to be a little bit faster with kernel fusion. I do agree that it's impractical to maintain a model-level blocklist for a specific CC level, would disabling kernel fusion for MXFP4 on CC 8.7 be a little bit more manageable?

[am17an]

I don't see why mxfp4 would be a problem though. My best guess is that it's a bandwidth issue, so programmatically getting the bandwidth and comparing the model arch (MoE or dense) would be more robust

[JohannesGaessler]

If the issue is specific to mxfp4 then I suspect that CC 8.7 does not have hardware support for the __byte_perm instruction, resulting in increased register usage that becomes problematic with kernel fusion. Can you test the performance of the following patch?

diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh
index 6baab1176..b573d02b0 100644
--- a/ggml/src/ggml-cuda/vecdotq.cuh
+++ b/ggml/src/ggml-cuda/vecdotq.cuh
@@ -54,7 +54,7 @@ static __device__ __forceinline__ int2 get_int_from_table_16(const int & q4, con
     uint32_t res_y = __builtin_amdgcn_perm(v_odd_high, v_odd_low, mask_odd);

     return make_int2(res_x, res_y);
-#elif !defined(GGML_USE_MUSA)
+#elif !defined(GGML_USE_MUSA) && __CUDA_ARCH__ != 870
     // CUDA does not have an instruction for selecting bytes with 4 bit indices.
     // However, __byte_perm is an instruction that selects bytes with 3 bit indices that can be used instead.
     const uint32_t * table32 = (const uint32_t *) table;