Eval bug: Qwen3-Coder-480B-A35B-Instruct-1M-GGUF GGML_ASSERT(ggml_nbytes(src0) <= INT_MAX) failed

[createthis]

Name and Version

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 PRO 6000 Blackwell Workstation Edition, compute capability 12.0, VMM: yes
version: 6087 (c3eb159)
built with cc (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0 for x86_64-linux-gnu

Operating systems

Linux

GGML backends

CUDA

Hardware

blackwell 6000 pro
EPYC 9355

Models

Qwen3-Coder-480B-A35B-Instruct-1M-GGUF

Problem description & steps to reproduce

./build/bin/llama-server \
    --model /data/Qwen3-Coder-480B-A35B-Instruct-1M-GGUF/UD-Q4_K_XL/Qwen3-Coder-480B-A35B-Instruct-1M-UD-Q4_K_XL-00001-of-00006.gguf \
    --alias Qwen3-Coder-480B-A35B-Instruct-GGUF:UD-Q4_K_XL \
    --no-webui \
    --numa numactl \
    --threads 32 \
    --ctx-size 400000 \
    --n-gpu-layers 63 \
    -ot "blk\.(3|4|5|6|7|8|9|10|11|12|13)\.ffn_.*=CUDA0" \
    -ot exps=CPU \
    -ub 4096 -b 4096 \
    --cache-type-k q4_1 \
    --cache-type-v q4_1 \
    --seed 3407 \
    --prio 3 \
    --temp 0.7 \
    --top-p 0.8 \
    --top-k 20 \
    --repeat-penalty 1.05 \
    --min-p 0.0 \
    --log-colors \
    --flash-attn \
    --host 0.0.0.0 \
    --jinja \
    --port 11434

Feed it more than 131072 context, watch it crash and burn.

First Bad Commit

No response

Relevant log output

slot update_slots: id  0 | task 0 | prompt processing progress, n_past = 131072, n_tokens = 4096, progress = 0.492306
/home/jesse/llama.cpp/ggml/src/ggml-cuda/cpy.cu:285: GGML_ASSERT(ggml_nbytes(src0) <= INT_MAX) failed
/home/jesse/llama.cpp/build/bin/libggml-base.so(+0x1594b)[0x7b7943a9294b]
/home/jesse/llama.cpp/build/bin/libggml-base.so(ggml_print_backtrace+0x21c)[0x7b7943a92dac]
/home/jesse/llama.cpp/build/bin/libggml-base.so(ggml_abort+0x15b)[0x7b7943a92f8b]
/home/jesse/llama.cpp/build/bin/libggml-cuda.so(_Z13ggml_cuda_cpyR25ggml_backend_cuda_contextPK11ggml_tensorPS1_b+0xa62)[0x7b7940c9dcb2]
/home/jesse/llama.cpp/build/bin/libggml-cuda.so(+0xeed58)[0x7b7940ceed58]
/home/jesse/llama.cpp/build/bin/libggml-base.so(ggml_backend_sched_graph_compute_async+0x463)[0x7b7943aaab13]
/home/jesse/llama.cpp/build/bin/libllama.so(_ZN13llama_context13graph_computeEP11ggml_cgraphb+0xa1)[0x7b794389c0e1]
/home/jesse/llama.cpp/build/bin/libllama.so(_ZN13llama_context14process_ubatchERK12llama_ubatch14llm_graph_typeP22llama_memory_context_iR11ggml_status+0x104)[0x7b794389d794]
/home/jesse/llama.cpp/build/bin/libllama.so(_ZN13llama_context6decodeERK11llama_batch+0x3bd)[0x7b79438a35dd]
/home/jesse/llama.cpp/build/bin/libllama.so(llama_decode+0xf)[0x7b79438a453f]
./build/bin/llama-server(+0xc1bbe)[0x619405c79bbe]
./build/bin/llama-server(+0x879e5)[0x619405c3f9e5]
./build/bin/llama-server(+0x4ef0e)[0x619405c06f0e]
/lib/x86_64-linux-gnu/libc.so.6(+0x2a1ca)[0x7b794302a1ca]
/lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0x8b)[0x7b794302a28b]
./build/bin/llama-server(+0x50f35)[0x619405c08f35]
Aborted (core dumped)

[krzychdre]

I got exactly the same error working with unsloth/Qwen3-Coder-30B-A3B-Instruct-GGUF:Q8_K_X
--no-webui
--numa numactl
--prio 2
--slots
--jinja
--no-mmap
--threads 16
--ctx-size 262144
--no-context-shift
--batch-size 4096
--ubatch-size 4096
--cache-type-k q8_0
--cache-type-v q8_0
--cache-reuse 256
--flash-attn
--seed 3407
-b 4096 -ub 4096
-ngl 63
--parallel 1
--temp 0.7 --min-p 0.0 --top-p 0.8 --top-k 20 --repeat-penalty 1.05
-ts 10,10,3,1

4 GPUs (2x5090 + 2x3090) after passing 131072 with Q8_0 K and V I got:
GGML_ASSERT(ggml_nbytes(src0) <= INT_MAX) failed

[createthis]

GPT OSS high claims this problem can be worked around if the machine has multiple GPUs by using --gpu-split. I have a single blackwell 6000 pro, so that isn't an option for me.

GPT OSS high's further assessment of the problem:

What the current code does

ggml_cuda_cpy() (the entry point for every CUDA‑backed copy) contains:

const int64_t ne = ggml_nelements(src0);
GGML_ASSERT(ggml_nbytes(src0) <= INT_MAX);
GGML_ASSERT(ggml_nbytes(src1) <= INT_MAX);
…
ggml_cpy_flt_cuda<float,float>(src0_ddc, src1_ddc,
                               ne, ne00, ne01, ne02,
                               nb00, nb01, nb02, nb03,
                               ne10, ne11, ne12,
                               nb10, nb11, nb12, nb13,
                               main_stream, dest_ptrs_d, graph_cpynode_index);

All the *_cuda helpers (e.g. ggml_cpy_flt_cuda, ggml_cpy_f32_q8_0_cuda, …) are declared as:

static void ggml_cpy_f32_q8_0_cuda(const char *cx, char *cdst, const int ne, …);

• ne is int, not int64_t.
• Inside those helpers the kernel launch looks like:

const int num_blocks = (ne + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE;
cpy_f32_q<cpy_blck_f32_q8_0, QK8_0><<<num_blocks, 1, 0, stream>>>(…);

So the number of elements (ne) and the grid size (num_blocks) are both limited to a signed 32‑bit integer.

Why that matters for a 1 M‑token context

For a 480 B model the KV‑cache per layer is roughly

n_ctx × n_kv_heads × head_dim × 2  elements
≈ 1 000 000 × 64 × 128 × 2 ≈ 16 384 000 000 elements

Even a 400 k‑token context already yields > 6 × 10⁹ elements, which is well above INT_MAX (2 147 483 647).

• The int ne argument will overflow, truncating the element count.
• num_blocks will overflow as well, causing an illegal kernel launch (the grid dimension cannot exceed 2³¹‑1).
• The kernel will therefore read/write the wrong memory region, leading to corruption or a crash – the same failure you observed earlier.

The cudaMemcpyAsync path (the fast‑path for contiguous tensors of the same type) does accept a size_t length, but the fallback kernels that handle type conversion, quantisation, etc. cannot handle a ne larger than a 32‑bit integer.

---

What would be required to truly support > 2 GB tensors

1. Change the API signatures

   • All *_cuda functions that receive ne must take int64_t ne (or size_t).
   • All internal helpers (cpy_flt_cuda, cpy_f32_q…, etc.) must be updated accordingly.

2. Adapt kernel launch logic

   • CUDA kernels can only be launched with a 32‑bit grid dimension.
   • For ne > 2³¹‑1 you must split the work into multiple kernel launches (or use a 2‑D grid).
   • A typical pattern is:

     const int64_t elems_per_launch = (int64_t)INT_MAX;   // or a safe multiple of QK
     for (int64_t offset = 0; offset < ne; offset += elems_per_launch) {
         int cur_ne = (int)min(elems_per_launch, ne - offset);
         const int num_blocks = (cur_ne + QK - 1) / QK;
         kernel<<<num_blocks, 1, 0, stream>>>(cx + offset_bytes,
                                              cdst + offset_bytes,
                                              cur_ne, …);
     }

   • The kernels themselves must use 64‑bit offsets when computing byte addresses (x_offset, dst_offset). The current kernels compute offsets with int64_t i = blockDim.x*blockIdx.x + threadIdx.x; – that part is already 64‑bit, but the grid‑derived i will be limited by the 32‑bit num_blocks. Splitting the launch solves that.

3. Update any other size‑related checks

   • ggml_nbytes() already returns size_t, so the cudaMemcpyAsync path is fine.
   • Any other places that cast ggml_nelements() to int (e.g. in other back‑ends or utility functions) must be audited.

[createthis]

Kinda silly, but I appear to have vibe coded something that at least partially works: #15298

[ExtReMLapin]

Wasn't the 131k issue limited to a not implemeted Dual Chunk Attention instead ?

[createthis]

@ExtReMLapin I believe I ran into this INT_MAX assertion on Qwen3-Coder-480B-A35B-Instruct-GGUF (non-1M variant) when I attempted to use more than 131k. If there is an issue with Dual Chunk Attention and the 1M variant, that may be a separate issue.