INT8 Tensor Core corruption on Turing sm_75 cuBLAS llama.cpp BPE tokenizer

AI & Data Science Deep Learning (Training & Inference) TensorRT
, , , ,
1

Description

BPE tokenizers (Llama 3.1, Lexi) produce corrupted output (“???”) on Turing GPUs (Quadro RTX 6000, sm_75) when used via llama.cpp’s CUDA backend. SentencePiece tokenizers (Gemma) work correctly on the same hardware. I use LM STUDIO.

Environment

TensorRT Version: N/A (using cuBLAS directly via llama.cpp)
GPU Type: Quadro RTX 6000 (Turing sm_75, 24GB) + RTX 3070 Ti (Ampere sm_86, 8GB)
Nvidia Driver Version: 580.88
CUDA Version: 12.4
CUDNN Version: N/A
Operating System + Version: Windows 10 22H2 + WSL2 Ubuntu 24.04
Python Version: N/A
TensorFlow Version: N/A
PyTorch Version: N/A
Baremetal or Container: Baremetal

Relevant Files

llama.cpp source: GitHub - ggml-org/llama.cpp: LLM inference in C/C++ · GitHub (tag b9888, commit cb295bf)
ggml-cuda.cu: device property query at line 292 (cc = 100prop.major + 10prop.minor)
common.cuh: cublasSetMathMode at line 1478
mmq.cu: ggml_cuda_should_use_mmq() at line 267 (turing_mma_available check)

Steps To Reproduce

< 1. Build llama.cpp with CUDA support:
cmake -B build -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=“75-real;86-real”
cmake --build build --target llama-server

  1. Run with Turing GPU only (isolate via CUDA_VISIBLE_DEVICES):
    set CUDA_VISIBLE_DEVICES=1
    llama-server.exe -m Llama-3.1-8B-Q5_K_M.gguf --port 1234 -ngl 99

  2. Send inference request:
    curl -X POST http://localhost:1234/v1/chat/completions
    -d ‘{“model”:“test”,“messages”:[{“role”:“user”,“content”:“Hello”}]}’

  3. Result: “???” (BPE tokenizer corrupted)
    Expected: “Hello! How can I assist you?”

  4. Same test with CUDA_VISIBLE_DEVICES=0 (Ampere GPU):
    Result: “Hello! How can I assist you?” (clean output)

Additional info:

  • Model is confirmed on GPU via nvidia-smi (21.9GB VRAM on Quadro)
  • SentencePiece models (Gemma 26B) work correctly on Quadro
  • Only BPE tokenizers (Llama, Lexi, Qwen) are corrupted
  • Tested with GGML_CUDA_FORCE_CUBLAS=ON → still corrupted
  • Tested with cublasSetMathMode(CUBLAS_DEFAULT_MATH) → still corrupted
  • Tested with GGML_CUDA_FORCE_CUBLAS_COMPUTE_32F=1 → still corrupted!

Please include:

  • Exact steps/commands to build your repro
  • Exact steps/commands to run your repro
  • Full traceback of errors encountered

Regards

Tsitola

3

ROOT CAUSE IDENTIFIED: Defective GDDR6 Memory Chip

After 32 custom CUDA builds, cross-backend validation (CUDA + Vulkan), and 15+
hours of systematic debugging, the root cause is a single defective GDDR6
memory chip on the Quadro RTX 6000, NOT a software bug in llama.cpp, CUDA,
or the K-quant format.

METHODOLOGY

1. GPU-Agnostic Cross-Backend Testing

Same llama.cpp binary, same Q5_K_M model:

  • RTX 3070 Ti (Ampere sm_86): correct output on both CUDA and Vulkan
  • Quadro RTX 6000 (Turing sm_75): garbage on CUDA, degraded on Vulkan

Two independent software stacks failing on ONE GPU → hardware issue.

2. ECC Diagnostics (Official NVIDIA Tool)

nvidia-smi -i 1 -q -d ECC
DRAM Correctable:   3
DRAM Uncorrectable: 259

ECC cannot correct multi-bit stuck-at-0 failures.

3. Third-Party Validation

  • OCCT VRAM test: 604,057 errors in 30 minutes
  • FurMark: passed (uses only graphics pool, not compute pool — see below)

4. Custom CUDA Memory Diagnostics (Open Source — Available on Request)

Written by us to precisely locate the defective region:

a) Full VRAM pattern test (3-pass random): 1,057,402 errors
b) Walking-bit test: all 8 bits dirty, stuck-at-0
c) Continuous burn test: 6.24 trillion errors in 3,000 iterations
d) Chunk-by-chunk mapper (64 MB granularity, GPU kernel testing):
360 chunks tested → 1 BAD chunk confirmed
Bad region: ~1.31-1.75 GB (moves between reboots due to CUDA VA remapping)

CONCLUSION

This is a hardware defect (defective GDDR6 chip causing multi-bit
stuck-at-0 errors), confirmed by:

  • NVIDIA’s own ECC diagnostics (259 uncorrectable DRAM errors)
  • OCCT VRAM test (604K errors)
  • Custom CUDA memtest (1M+ errors, walking-bit, burn test)

The GPU requires replacement . A software workaround
(reserving the 64 MB defective region at driver initialization) has
been deployed and confirmed working.

4

This is a textbook piece of debugging, and your conclusion is right. Thirty-two builds and 15 hours to get from “BPE tokenizers are broken” to a specific defective memory chip is a lot of persistence, and the write-up is more useful than most bug reports I see.

Summarizing for anyone who lands here from a search with the same symptom, because the framing is what costs people days. The tokenizer correlation was a red herring. BPE and SentencePiece models do not stress different arithmetic paths in any way that would explain garbage output; what they do differently is allocate, so only some runs placed hot data on the bad region. Chasing the tokenizer means chasing a shadow.

The two steps that actually isolated it, in the order to do them:

  1. Run the same binary and the same model on a second GPU. Clean on one card and broken on the other rules out the software stack entirely.
  2. Check ECC before anything else. From your post #3, nvidia-smi -i 1 -q -d ECC on the affected card reported:
DRAM Correctable:   3
DRAM Uncorrectable: 259

259 uncorrectable DRAM errors is conclusive on its own. Everything after that (OCCT’s 604,057 errors, your own 1,057,402-error pattern test, the walking-bit stuck-at-0 result, the single bad 64 MB chunk out of 360) is confirmation rather than diagnosis. Your point about FurMark passing is the other detail to keep: a graphics stress test that exercises a different memory pool can pass on a card whose compute allocations are corrupting, so “it passes FurMark” is not evidence of healthy VRAM.

On next steps, two things.

The card needs replacing, and a Quadro RTX 6000 with uncorrectable DRAM errors is a straightforward warranty case. That goes through whoever you bought it from (board vendor, system integrator, or reseller) rather than through the forums, and the ECC counter output above plus your OCCT results are the evidence they will ask for. Capture nvidia-smi -q in full while you still have the card in hand.

⚠️ I would not run the region-reservation workaround in production, and your own findings are the reason. You noted the bad region moves between reboots because of CUDA virtual address remapping, so a reservation pinned to one address range is only correct until the next boot. And 259 uncorrectable errors describe a failing chip rather than a fixed set of bad cells, so corruption outside the region you mapped is possible, and it would show up as subtly wrong model output rather than an obvious crash. Treat the reservation as a way to keep working while the RMA is in flight, not as a fix.

Thanks for coming back and posting the resolution instead of just walking away once you solved it. That is the part that makes this thread useful to the next person.

Thank You !

5

Thanks for the detailed response and the RMA guidance — really appreciate you taking the time.

A few updates since the original post, for anyone landing here with the same symptoms:

  1. The guard has been tightened. Further testing with vram_mapper_detailed at finer granularities (32 → 16 → 8 → 4 MB) showed the defect concentrates in ~8 MB at 1.719 GB physical. The guard now uses 8 MB chunks (±1 neighbor = 24 MB reserved) instead of the original 128 MB — recovering ~100 MB of usable VRAM. The L2 cache masking limit (6 MB on TU102) means 4 MB chunks can’t be reliably tested; 8 MB is the practical floor.

  2. Canary + rescan added. Since the bad chunk index drifts between boots due to WDDM VA→PA remapping, the guard now includes ggml_cuda_guard_verify() (fast 1ms check before inference) and ggml_cuda_guard_rescan() (full re-scan + re-guard). If another GPU-heavy app causes WDDM eviction, the caller can unload model buffers, rescan, and reload — no restart required.

  3. VMM (cuMemCreate) is a dead end on WDDM/Turing. We spent significant time trying to use CUDA Virtual Memory Management to quarantine bad physical pages globally. The conclusion: on this driver stack, cuMemCreate draws from a different physical pool than cudaMalloc — the defective chip #16 is never allocated through VMM handles. Strategy B is not viable on this platform.

  4. Dynamic Page Retirement (DPR) — untested, could be permanent. NVIDIA’s driver supports retiring defective pages via ECC page retirement. The blacklist persists across reboots and survives ECC OFF. In theory, this solves the problem at the hardware level — no software guard needed. We haven’t attempted it because: (a) the ~4 MB documented cap may not fully cover the defect region, (b) ECC ON currently triggers TDR crashes on multi-bit errors on this card, and (c) there’s no documented undo for InfoROM page retirement entries. Worth attempting before RMA.

  5. On the “failing chip” warning — our data doesn’t support it. The 259 DRAM uncorrectable errors are cumulative since manufacturing, not an active degradation rate. Across multiple boots and tests spanning a week, the defect shows the same physical address (~1.719 GB), the same pattern (stuck-at-0 cells), and the same magnitude of errors (~4,000 at 8 MB granularity). No expansion observed. This looks like a manufacturing-level silicon defect in one GDDR6 chip, not progressive failure. The guard as a bridge is still the right posture, but the data suggests the defect is stable.

  6. Noted on the RMA path. We’ll capture nvidia-smi -q and proceed through the vendor. The guard keeps the card usable in the meantime.

Thanks again for confirming the diagnostic methodology — “check ECC before anything else” is now in our permanent playbook