Repeatable CUDA errors with seemingly arbitrary contexts

[VL4DST3R]

EDIT: it may be related to #2005

I'm getting CUDA errors occasionally when processing prompts (never when actually generating). This is not a new issue but something I've been experiencing for a while few months now. It was never that big of an issue to bother addressing since it doesn't happen too often, I'd say once every 200gens or so, and I have kobold set to restart when this happens, but it is a bit annoying. The usual scenario when this happens (mainly when doing agentic stuff via OpenCode) I would get a processing block in my logs like:

[SmartCache RNN No Match, Saving into slot 2...]

state_write_data: writing state
state_write_data: - writing memory module

state_write_data: writing state
state_write_data: - writing memory module

KV Save State 2: Created SaveState of 101341 tokens, costing 1116 MB.

Processing Prompt [BATCH] (40960 / 57352 tokens)
state_write_data: writing state
state_write_data: - writing memory module

state_write_data: writing state
state_write_data: - writing memory module

KV Save State 6: Created SaveState of 40960 tokens, costing 488 MB.

[SmartCache RNN Lifeboat: Saved 40960-token checkpoint into slot 6]
Processing Prompt [BATCH] (57352 / 57352 tokens)
state_write_data: writing state
state_write_data: - writing memory module

state_write_data: writing state
state_write_data: - writing memory module

KV Save State 1: Created SaveState of 57345 tokens, costing 658 MB.
CUDA error: invalid argument
  current device: 0, in function ggml_cuda_kernel_launch at D:\a\koboldcpp\koboldcpp\ggml\src\ggml-cuda\common.cuh:1674
  cudaGetLastError()
D:\a\koboldcpp\koboldcpp\ggml\src\ggml-cuda\ggml-cuda.cu:108: CUDA error

and then it would just die with the same error trace every time. If I don't change the session context that caused it, it will keep reprocessing and crashing indefinitely (the Repeatable part), but change the context in any way (add another subsequent message or branch off in any way) and it will go past it and continue to work fine. From trying to read up on the issue I couldn't find much info here on the issue tracker so I thought I'd ask.

My setup: 1x4090 24gb, with the following .kcpps settings (all model layers on gpu):

{
    "model_param": "S:/KoboldCPP/Qwen3.6-35B-A3B-UD-IQ4_NL.gguf",
    "mmproj": "S:/KoboldCPP/mmproj/Qwen3.6-mmproj-BF16.gguf",
    "mmprojcpu": true,
    "threads": 11,
    "usecuda": [
        "normal",
        "0"
    ],
    "contextsize": 200000,
    "batchsize": 2048,
    "smartcache": 5,
    "ropeconfig": [
        0.0,
        1000000.0
    ],
    "onready": [
        "S:/KoboldCPP/Post-Launch.bat"
    ],
    "websearch": true,
    "savedatafile": "S:/KoboldCPP/SaveDataFile.jsondb",
    "jinja": true,
    "jinja_tools": true,
    "quantkv": "q8_0",
    "moeexperts": 8,
    "gendefaults": {
        "temperature": 0.6,
        "top_p": 0.95,
        "top_k": 20,
        "min_p": 0.0,
        "presence_penalty": 0.0,
        "rep_pen": 1.0
    },
    "sdthreads": 11,
    "sdlora": [],
    "ttsdir": "S:/KoboldCPP/voices",
    "admin": true,
    "adminpassword": "",
    "admindir": "S:/KoboldCPP/admin"
}

Am I doing something wrong?

[VL4DST3R]

My clanker's report on the matter, if interested:

Expand

Analysis: CUDA invalid argument crash during SmartCache RNN processing

What's happening

Your log shows this sequence:

1. SmartCache RNN No Match - The prompt doesn't match any cached slot, so it saves the current state to the oldest slot (slot 2, 101341 tokens / 1116 MB)
2. Processing Prompt [BATCH] (40960 / 57352 tokens) - Pre-fill batch 1 begins
3. Processing Prompt [BATCH] (57352 / 57352 tokens) - Pre-fill batch 2 (final) begins
4. SmartCache RNN Lifeboat: Saved 40960-token checkpoint into slot 6 - Lifeboat checkpoint saved during prefill
5. CUDA error: invalid argument at ggml_cuda_kernel_launch in common.cuh:1674

The crash happens after the second batch starts processing, during llama_decode on the Qwen3.6-35B-A3B model.

Root Cause

This is a CUDA grid dimension overflow bug in ggml-cuda. The error occurs at common.cuh:1674:

kernel<<<launch_params.block_nums, launch_params.block_dims, ...>>>(...);
CUDA_CHECK(cudaGetLastError());  // <-- line 1674

CUDA's gridDim.y and gridDim.z are limited to 65535 (only gridDim.x supports up to 2^31-1). When the context gets very large (57K+ tokens), certain kernels pass values exceeding 65535 to grid.y, causing an invalid argument (CUDA error 11 / cudaErrorInvalidConfiguration).

The smoking gun is issue ggml-org/llama.cpp#27901 - identical error, identical stack trace, same line number (common.cuh:1674):

rms_norm_f32 exceeds the CUDA gridDim.y limit (65535) at n_ctx 262144

The offending tensor maps nchannels (which scales with context) onto gridDim.y. In norm.cu:307:

const dim3 blocks_num(nrows, nchannels, nsamples);  // y = nchannels, cap 65535

When nchannels > 65535, the kernel launch fails with invalid argument.

For Qwen4exp (Qwen3.8-Flash-Next), the formula is n_ctx / 4 (hyper-connection ratio = 4), so it overflows at n_ctx >= 262144. Your model (Qwen3.6-35B-A3B, hybrid RNN) has a different architecture but the same pattern - when the RNN memory module is active and context is large, the grid.y overflows.

Why it's "occasional"

The overflow only triggers when the KV cache + RNN state reaches a certain size threshold. With 57K tokens and ~1100 MB of RNN state, you're near or past that boundary. Shorter prompts or fewer accumulated tokens won't trigger it.

Related issues found

Issue	Details
ggml-org/llama.cpp#27901	Same error, same line, rms_norm grid.y overflow at large context
ggml-org/llama.cpp#27871	qwen4exp aborts at native context depth, identical stack
ggml-org/llama.cpp#19999	Fixed - grid.y overflow in non-contiguous dequantize/convert kernels
LostRuins/koboldcpp#1896	SmartCache + state_write_data + CUDA error (Vulkan, but same pattern)

Your KoboldCPP version

You're running v1.119. The grid.y overflow in dequantize kernels was fixed in PR ggml-org#19999 (March 2026), so that specific path should be fixed. However, the rms_norm grid.y overflow in norm.cu (issue ggml-org#27901) was reported on 2026-08-28 (yesterday!) and may not be in your build yet.

Workarounds

1. Reduce context size - If you're using a very large context (128K+), try reducing it. The overflow is context-size dependent.
2. Disable SmartCache / Context Shift - If you don't need it, try --noshift or --nofastforward to disable the RNN checkpointing that triggers the large memory module writes.
3. Update KoboldCPP - Check if a newer version includes the norm.cu grid.y fix. The issue is extremely recent (Aug 28, 2026).
4. Reduce smartcache slots - Fewer checkpoints = less RNN memory pressure: --smartcache N with a smaller N.

The fix (from the upstream analysis) involves either capping grid.y at 65535 with a grid-stride loop, or decomposing the large dimension across multiple grid axes - similar to what was already done in binbcast.cu:317.

[LostRuins]

How close are you to the RAM/VRAM limit? When running at full ctx, how much free memory remains?

[VL4DST3R]

Very much in the red, out of 24 i'm usually left with 500-800mb free, depending on what other things i have open at the same time (with a browser and a youtube video for example ive seen it down to ~200mb free left)

[LostRuins]

Perhaps it's related to that then, a memory allocation fails and thus results in cuda crapping out. If you drop back a few layers I'd imagine it works better?

[VL4DST3R]

Maybe, although consider the following:

• I used to have the vram limits even tighter: when it released i kept the mmproj model loaded into vram as well, yet the crashing frequency didn't change in any meaningful way with the now extra breathing room.
• When it the cuda error happens the system remains otherwise stable and continues just fine after kobold reboots. Whereas if i forget a game running before starting kobold and attempt to load a model as well, it brings down the entire driver for a second or two and kobold crashes with a different "could not allocate" type of error.

This leads me to think it's not something like a memory spike that tries to use up the last few hundreds of mb of vram and failing, but an actual bug with this kind of model.

Moreover from the issue thread I linked above, both that and the upstream issue referenced there all share the same family of models (qwen). The ai exploration seemed to corroborate this as well.