llama.cpp CUDA 13.2 gibberish for IQ3_S / IQ3_XXS, IQ2_M - `SOLUTION use CUDA < 13.0 or Unsloth Studio`

[danielhanchen]

TLDR: CUDA 13.2 seems problematic for all quants using IQ3_S - full gibberish will be seen.
CUDA 12.8 works fine - CUDA 13.0 is also fine.

SOLUTION: Use CUDA 12.8 and CUDA 13.0 binaries or use Unsloth Studio which compiles with CUDA 13.0

[danielhanchen]

Reproduced on an RTX PRO 6000 Blackwell Server Edition (compute 12.0, 96GB VRAM).

Setup:

• Driver: 580.82.07
• CUDA 12.8: nvcc V12.8.93 (baseline, works correctly)
• CUDA 13.2: nvcc V13.2.51 + cuda-compat-13-2 (broken)
• llama.cpp: commit 650bf14 (latest main)
• All builds targeting CMAKE_CUDA_ARCHITECTURES=120

Results -- Gemma-4-31B-it (unsloth/gemma-4-31b-it-GGUF):

Build	Quant	Output	Coherent?
CUDA 12.8	IQ3_XXS	[Start thinking] * User wants a short joke... Why don't scientists trust atoms?	Yes
CUDA 13.2	IQ3_XXS	deHH laesLH laHse지원KLHLeHsenH坐es الأخرىHal laHLLteLH...	No
CUDA 13.2 + FORCE_CUBLAS	IQ3_XXS	CHistTLBalal/THHS singlealalistHHalingLH/alLHLalHen...	No
CUDA 12.8	IQ2_M	[Start thinking] * User wants a short joke... Why don't scientists trust atoms?	Yes
CUDA 13.2	IQ2_M	laistySS own own own laBoge// own own own own de que que la la...	No
CUDA 13.2	IQ4_XS	[Start thinking] * User wants a "short joke"... Why don't scientists trust atoms?	Yes
CUDA 13.2	IQ4_NL	[Start thinking] * User wants a short joke... Why don't scientists trust atoms?	Yes
CUDA 13.2	Q4_K_M	[Start thinking] * User wants a short joke... Why don'...	Yes

Results -- Qwen3.5-35B-A3B (unsloth/Qwen3.5-35B-A3B-GGUF):

Build	Quant	Output	Coherent?
CUDA 12.8	IQ3_S	[Start thinking] Thinking Process: 1. Analyze the Request...	Yes
CUDA 13.2	IQ3_S	I\ns,\nHello!	No
CUDA 13.2 + FORCE_CUBLAS	IQ3_S	(2 p: wjg')'. 1 and 1 0 * + abort/core dump	No + crash

All tests used: -n 64 --temp 0.0 --top-k 1 --ctx-size 512 --no-mmap -ngl 999

Additional findings beyond the original report:

1. IQ2_M is also affected -- not just IQ3_S/IQ3_XXS. The bug boundary appears to be IQ3 and below (IQ2_M, IQ3_XXS, IQ3_S broken) vs IQ4 and above (IQ4_XS, IQ4_NL, Q4_K_M all fine).
2. GGML_CUDA_FORCE_CUBLAS=ON does not help on this configuration -- still produces gibberish, and for Qwen IQ3_S it actually crashes with an abort/core dump. This differs from Eval bug: gemma-4-31B-it-UD-IQ3_XXS.gguf (unsloth) produces gibberish on blackwell (RTX 5080 tested) with GGML_CUDA=ON compilation VS normal correct output with BLAS/CPU compilations ggml-org/llama.cpp#21371 where cuBLAS was reported as a workaround.
3. Recompiling with CUDA 12.8 fixes everything, consistent with other reports.

Upstream issues: ggml-org/llama.cpp#21255 and ggml-org/llama.cpp#21371. NVIDIA has acknowledged and CC'd the CUDA compiler team.

[danielhanchen]

Comprehensive Cartesian Combo Testing Results

Ran 290 individual tests across all combinations of: 2 models x 13 quant types x 6 build configs x 2 flash-attention settings. Every CUDA 13.2 "GIBBERISH" result was manually verified to have a corresponding "OK" result on CUDA 12.8, CPU, and CUDA 13.0 for the same quant.

Test harness: llama.cpp commit 650bf14, prompt "Hello! Can you tell me a short joke?", -n 64 --temp 0.0 --top-k 1 --ctx-size 512 --no-mmap -ngl 999

GPU: RTX PRO 6000 Blackwell Server Edition (compute 12.0, 96GB)

Builds tested:

• CPU-only (CUDA 12.8 binary with CUDA_VISIBLE_DEVICES='')
• CUDA 12.8 (nvcc V12.8.93)
• CUDA 13.0 (nvcc V13.0.76)
• CUDA 13.0 + GGML_CUDA_FORCE_CUBLAS=ON
• CUDA 13.2 (nvcc V13.2.51)
• CUDA 13.2 + GGML_CUDA_FORCE_CUBLAS=ON

All builds targeting CMAKE_CUDA_ARCHITECTURES=120.

---

Gemma-4-31B-it (unsloth/gemma-4-31b-it-GGUF)

Quant	CPU	CUDA 12.8	CUDA 13.0	13.0+cuBLAS	CUDA 13.2	13.2+cuBLAS
IQ2_M	OK	OK	OK	OK	GIBBERISH	GIBBERISH
Q2_K_XL	OK	OK	OK	OK	GIBBERISH	GIBBERISH
IQ3_XXS	OK	OK	OK	OK	GIBBERISH	GIBBERISH
Q3_K_XL	OK	OK	OK	OK	OK	OK
IQ4_XS	OK	OK	OK	OK	OK	OK
IQ4_NL	OK	OK	OK	OK	OK	OK
Q4_K_M	OK	OK	OK	OK	OK	OK
Q4_K_XL	OK	OK	OK	OK	OK	OK
Q5_K_XL	OK	OK	OK	OK	OK	OK
Q6_K_XL	OK	OK	OK	OK	OK	OK
Q8_0	OK	OK	OK	OK	OK	OK

Qwen3.5-35B-A3B (unsloth/Qwen3.5-35B-A3B-GGUF)

Quant	CPU	CUDA 12.8	CUDA 13.0	13.0+cuBLAS	CUDA 13.2	13.2+cuBLAS
IQ2_XXS	OK	OK	OK	OK	OK	OK
IQ2_M	OK	OK	OK	OK	OK	OK
Q2_K_XL	OK	OK	OK	OK	OK	OK
IQ3_XXS	OK	OK	OK	OK	OK	OK
IQ3_S	OK	OK	OK	OK	GIBBERISH	GIBBERISH
Q3_K_XL	OK	OK	OK	OK	OK	OK
IQ4_XS	OK	OK	OK	OK	GIBBERISH	GIBBERISH
IQ4_NL	OK	OK	OK	OK	GIBBERISH	GIBBERISH
Q4_K_M	OK	OK	OK	OK	OK	OK
Q4_K_XL	OK	OK	OK	OK	OK	OK
Q5_K_XL	OK	OK	OK	OK	OK	OK
Q6_K_XL	OK	OK	OK	OK	OK	OK
Q8_0	OK	OK	OK	OK	OK	OK

Each cell represents both FA-on and FA-off results (they were always identical -- flash attention setting has no effect on this bug).

---

Key Findings

1. CUDA 13.0 is completely unaffected -- zero failures across all 104 CUDA 13.0 tests (including +cuBLAS). This confirms the regression is specific to CUDA 13.2.

2. CPU-only is completely unaffected -- confirming this is a CUDA codegen bug, not a model/quant format issue.

3. cuBLAS workaround does NOT help -- GGML_CUDA_FORCE_CUBLAS=ON produces identical gibberish on CUDA 13.2. For Qwen IQ3_S+cuBLAS, it additionally crashes with SIGABRT.

4. Flash attention on/off has no effect -- the bug manifests identically regardless of FA setting.

5. The affected quant types differ by model architecture:

   • Gemma-4-31B: IQ2_M, IQ3_XXS, Q2_K_XL broken (low-bit IQ quantization + Q2_K)
   • Qwen3.5-35B-A3B: IQ3_S, IQ4_XS, IQ4_NL broken (mid-range IQ quantization)
   • Notably, Qwen's IQ2_M and IQ2_XXS work fine on CUDA 13.2, while they break on Gemma. And Gemma's IQ4_XS and IQ4_NL work fine, while they break on Qwen.

6. Standard K-quants are unaffected on both models: Q3_K_XL and above (Q3_K_XL, Q4_K_M, Q4_K_XL, Q5_K_XL, Q6_K_XL, Q8_0) all work correctly on CUDA 13.2.

The fact that different IQ quant types break on different model architectures suggests the CUDA 13.2 compiler generates incorrect code for specific dequantization kernel instantiations, likely depending on tensor shape parameters (which differ between Gemma-4-31B and Qwen3.5-35B-A3B MoE).

Upstream: ggml-org/llama.cpp#21255 and ggml-org/llama.cpp#21371

[danielhanchen]

Root Cause Analysis: CUDA 13.2 NVVM Optimizer (cicc) Bug

Finding: -Xcicc -O1 Fixes the Issue

We rebuilt llama.cpp with CUDA 13.2 at different NVVM optimizer levels (-Xcicc -ON) to isolate the bug:

cicc opt level	Gemma IQ3_XXS	Qwen IQ3_S	Qwen IQ4_XS
-Xcicc -O0	OK	OK	OK
-Xcicc -O1	OK	OK	OK
-Xcicc -O2	GIBBERISH	GIBBERISH	GIBBERISH
-Xcicc -O3 (default)	GIBBERISH	GIBBERISH	GIBBERISH

The bug is introduced by the NVVM IR optimizer (cicc) at optimization level O2.

What Does NOT Fix It

Flag	Result
nvcc -O0 (host only)	Still broken
--fmad=false	Still broken
--ptxas-options=-O0	Still broken
GGML_CUDA_FORCE_CUBLAS=ON	Still broken
Flash attention on/off	No effect

This rules out ptxas (the PTX-to-SASS assembler) and host-side compilation as the cause.

PTX Analysis

Comparing the PTX output from cicc -O1 (working) vs cicc -O2 (broken) for mmvq.cu:

• O1 PTX: 110,942 lines
• O2 PTX: 226,284 lines (2x larger)

The O2 optimizer aggressively inlines and transforms the code, more than doubling the PTX size. The miscompilation likely occurs during this aggressive inlining/optimization of the IQ dequantization functions.

Which Kernels Are Affected

The broken kernels share structural patterns not present in the working kernels:

1. Device-side lookup tables (static __device__ arrays in constant memory): iq2s_grid[1024], iq3s_grid[512], iq3xxs_grid[256], kvalues_iq4nl[16]
2. __byte_perm intrinsic chains for table lookup emulation (IQ4_NL, IQ4_XS)
3. __vcmpne4 + __vsub4 XOR-subtract sign-negation pattern (IQ2_S, IQ3_XXS, IQ3_S)

The working K-quants (Q3_K, Q4_K, Q5_K, Q6_K, Q8_0) use none of these -- they decode purely via bit-shifts and masks.

Workaround

Adding -Xcicc -O1 to the CUDA compilation flags completely fixes the issue on CUDA 13.2 while retaining most optimizations:

set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcicc -O1")

Or for users building llama.cpp:

cmake -DCMAKE_CUDA_FLAGS="-Xcicc -O1" ...

This is a more targeted workaround than downgrading to CUDA 12.8 entirely, though using CUDA 12.8 or 13.0 remains the safest option.

Suggested Action for NVIDIA

The bug is in the CUDA 13.2 cicc (NVVM IR) optimizer at O2+ and affects dequantization kernels that use device-side lookup tables and __byte_perm/__vcmpne4/__vsub4 intrinsics. A reduced test case focusing on the get_int_from_table_16() function from llama.cpp's vecdotq.cuh (which chains 5 __byte_perm calls) or the __vcmpne4+__vsub4 sign-application pattern would likely be sufficient to reproduce.

NVIDIA is already tracking this via the llama.cpp upstream issues -- CC @gaugarg-nv @prewaskar-nv if they can investigate the cicc O2 regression specifically.

[gaugarg-nv]

@danielhanchen Thanks for the detailed root cause analysis. I have passed on this information to CUDA compiler team.

[Panchovix]

Can confirm this issue as well on my setup.

[ChanifRusydi]

Can anyone confirm whether using CUDA 13.1 would working fine?