* rpc: disable unsafe memcmp graph cache
Same-shape prefill micro-batches compared equal and took the GRAPH_RECOMPUTE
path, re-running a stored graph against a grown KV context; the GLM-5.2 DSA
indexer then read past its buffers and crashed the server. Always send the full
graph. Upstream retired this cache design in ggml-org/llama.cpp#22701.
* rpc: use 64-bit ne/nb in rpc_tensor wire struct
ggml_tensor holds int64 ne and size_t nb; the wire struct stored them as uint32,
truncating any stride >= 4 GiB. The GLM-5.2 DSA indexer query stride crosses that
at ~26k tokens. Bump RPC_PROTO_MAJOR (wire-format change).
* ggml: use 64-bit locals in ggml_permute
Permuted strides were built in int locals, truncating any stride > 2 GiB before
it reached result->nb (size_t). Affects any permuted tensor over ~2 GiB.
* cuda : add the missing compile definitions to the HIP build
GGML_CUDA_FUSION, GGML_CUDA_MIN_BATCH_OFFLOAD and GGML_CUDA_PEER_MAX_BATCH_SIZE
are used unconditionally in common.cuh but were only defined in the CUDA branch,
so every HIP translation unit failed to compile. GGML_CUDA_IQK_FORCE_BF16 and
GGML_CUDA_F16 are user facing options the HIP branch silently ignored.
Also define GGML_USE_HIP. 43 tests in the sources imported from upstream use that
spelling, this fork only defined the older GGML_USE_HIPBLAS, so all of them took
the NVIDIA branch. Without it mmq_id_common.cuh defines TURING_MMA_AVAILABLE,
AMPERE_MMA_AVAILABLE, CP_ASYNC_AVAILABLE and FP16_MMA_AVAILABLE, i.e. the inline
PTX paths, and mmq_id.cu and mmq-instance-q6_k_id.cu then fail to build on
mma_new.cuh:181.
* cuda : shim the warp sync primitives in the HIP vendor header
ROCm 6 and later provide __shfl_sync() and friends as templates that static_assert
on the width of the mask, since an AMD wave can be 64 lanes wide. Define
HIP_DISABLE_WARP_SYNC_BUILTINS so that the shims replace them instead of clashing
with them, and add the shims the header was missing.
This is what makes mmq_id_common.cuh compile, so it unblocks all 26 mmq-instance-*_id
translation units, i.e. the MoE mat-mat path for the iqk quants.
* cuda : update the HIP vendor header for ROCm 6 and 7
- map nv_bfloat16 and nv_bfloat162 onto the __hip_bfloat16 types. Nothing declared
them, so every translation unit that mentions bf16 failed, convert.cu included -
that is the dequantize path the iqk quants use for prompt processing.
- use the hipblasComputeType_t and hipDataType entry points from ROCm 6.5 onwards.
hipblasDatatype_t is deprecated there and no longer matches hipblasGemmEx().
- make cudaStreamWaitEvent object-like. As a 3 argument function-like macro it did
not expand at the 2 argument call sites in reduce.cu, and the unexpanded name was
then passed on to CUDA_CHECK.
- add the mappings for cudaOccupancyMaxActiveBlocksPerMultiprocessor, which 29 of
the 30 failing flash attention translation units needed, and for the entry points
used by dsa_attn.cu and solve_tri.cu. hipBLAS spells a half as an unsigned short,
so cublasHgemmStridedBatched goes through a small casting wrapper rather than a
plain rename, which would not compile at the dsa_attn.cu call site.
* cuda : fix the remaining HIP build errors
- argsort.cu declared the sort order inside #ifdef GGML_CUDA_USE_CUB and used it
outside, and indexer_topk.cu calls argsort_f32_i32_cuda_cub() unconditionally
while it is only defined when CUB is available. Both break any build without CUB,
which includes MUSA and CUDA older than 11.7, not just HIP.
- ggml_backend_cuda_invalidate_graphs() touched ctx->cuda_graphs, which only exists
under USE_CUDA_GRAPH. The function is exported and called from llama-reload.cpp,
so guard the body rather than the function.
- solve_tri.cu included <cublas_v2.h> directly. common.cuh already pulls in whichever
vendor header is right, so the include is removed - it was redundant on CUDA too.
- dsa_attn.cu passed a data type where the GEMM entry point wants a compute type.
hipblasGemmStridedBatchedEx() has no data type taking overload. cuBLAS does, but it
is deprecated: cublas_api.h migrates CUDA_R_32F to exactly CUBLAS_COMPUTE_32F unless
the handle is in CUBLAS_PEDANTIC_MATH, which nothing in this tree sets. So this also
moves the CUDA build onto the primary entry point and drops a cublasGetMathMode()
per call. ggml-cuda.cu already passes a compute type to the same function.
- two mmvq instances called __dp4a() directly instead of ggml_cuda_dp4a(), the
wrapper the rest of the backend uses. On CUDA the wrapper is __dp4a() for every
architecture that has it.
- cap the flash attention vec f32 kernel at 4 columns per block on HIP. Both
logit_softcap variants of the 8 column kernel in one module overflow the 16 bit
branch offset of the AMDGPU backend.
* cuda : build the missing template instances in the HIP build
The HIP source list had drifted from the CUDA one and left out three families of
template instances that the backend references unconditionally:
- mmvq-instance-*.cu, the only definition site for the iqk mat-vec entry points.
iqk_mmvq.cu calls mul_mat_vec_iq4_ks_q8_1_cuda() and mul_mat_vec_iq4_kt_q8_1_cuda()
and nothing defined them, so the library did not link.
- the fattn-vec instances for q8_0-iq4_nl, iq4_nl-iq4_nl, q6_0-q5_0 and q8_0-q6_0,
which fattn-vec-f16.cu and fattn-vec-f32.cu dispatch to in the default
configuration.
- fattn-mma-*.cu. The MMA kernels are never selected on AMD, new_mma_available()
requires an NVIDIA device, but fattn-mma-f16.cu still needs the symbols.
* cuda : recognise AMD GPUs in GGML_CUDA_CC_IS_NVIDIA
CC_OFFSET_AMD is 1000000 and CC_OFFSET_MTHREADS is 0x100000, i.e. 1048576, so the
whole AMD range sits below the Moore Threads offset and every AMD GPU tested as
NVIDIA. turing_mma_available() then returned true on RDNA, the host picked an MMQ
tile of 128 while get_mmq_x_max_device() caps at 64 on AMD, and mul_mat_q_id hit
its NO_DEVICE_CODE guard and wrote NaNs. MUL_MAT_ID on IQ4_KS and IQ4_KT failed
this way on gfx1101.
No effect on CUDA, where a compute capability is 100*major + 10*minor and is
always far below CC_OFFSET_AMD.
* cuda : use v_perm_b32 for the 4 bit table lookup on HIP
HIP implements __byte_perm() in software: it stores an 8 byte union and does four
dynamically indexed byte loads, which end up in scratch. get_int_from_table_16()
calls it eight times per 32 weights, so every quant with a value table was paying
for that, while the trellis types were not.
__builtin_amdgcn_perm() is v_perm_b32, one instruction, and does the same job.
Taken from ggml-org/llama.cpp, which already carries this path.
Token generation on a 7800 XT, pure quantized Qwen2.5-1.5B, tg128:
IQ4_KS 16.30 -> 255.88 t/s
IQ4_XS 17.28 -> 268.05 t/s
IQ4_KT 201.76 -> 195.68 t/s (no table, unchanged)
Perplexity is unchanged to every printed digit and still matches the CPU exactly.
The function is duplicated in vecdotq.cuh and iqk_mmvq_templates.cuh, so both
copies need it - the iqk mat-vec instances only see the latter.
* cuda : use the shared flash attention support check on HIP
supports_op() carried a hand-rolled head size test for HIP that predates the
shared check: it accepted head size 64 with an f16 K cache and head size 128,
and nothing else. Head size 256 was rejected outright, so Gemma-2 and every
other 256 wide model fell back to the CPU for attention even though the
instances are compiled. @hardWorker254 reported 256 working with
ROCm 7.2.4 for both the f16 and the q8_0 cache.
Rather than adding 256 to the list, drop the branch and call
ggml_cuda_fattn_is_supported() as every other backend path does. It already
handles AMD: for cc >= CC_OFFSET_AMD it defers to the vec f16 or vec f32
support predicate depending on precision, which is exactly what
ggml_cuda_flash_attn_ext() dispatches to on AMD, because fast_fp16_available()
is true across the whole AMD cc range. The two now cannot drift.
This also removes a latent abort. The hand-rolled test returned true for any
head size 128 case regardless of the K and V types, so a combination without a
compiled instance, q4_1/q4_1 in a default build, reached the dispatcher and hit
on_no_fattn_vec_case() -> GGML_ABORT instead of falling back to the CPU. The
shared predicate is derived from the instances the build actually contains, and
after the source list repair earlier in this series the HIP build compiles the
same set as the CUDA build.
Beyond head size 256 this also lets HIP claim the asymmetric 192/128 and
576/512 vec f32 paths under GGML_PREC_F32. Those are untested on AMD; they are
gated by the same predicate CUDA uses.
* cuda : test the V head size, not the KV head count, for 192/128 vec f16 FA
ggml_cuda_fattn_vec_f16_is_supported() gates the asymmetric Dk != Dv branch on
if (K->ne[0] != 192 || V->ne[2] != 128) return false;
but ne[2] on K and V is the number of KV heads, not a head size. The test was
meant to be V->ne[0], as the wmma predicate added in the same commit (0459f595)
already writes it:
if (K->ne[0] != V->ne[0]) return K->ne[0] == 192 && V->ne[0] == 128;
and as the f32 twin has written it since 72201359 reworked that branch for
576/512. Only the f16 copy was left behind.
The kernels are there: ggml_cuda_flash_attn_ext_vec_f16() dispatches
FATTN_VEC_F16_CASE_DKDV(192, 128, ...) for f16-f16 and q8_0-q8_0 in both the
default and the GGML_CUDA_FA_ALL_QUANTS configuration, and the corresponding
hs192 instances are in the source list either way. The predicate just never
reported them, so a 192/128 shape whose KV head count was not coincidentally
128 was declined and attention fell back to the CPU.
This belongs in this series because the previous commit is what makes the
predicate reachable on AMD: with supports_op() routing flash attention through
ggml_cuda_fattn_is_supported(), the cc >= CC_OFFSET_AMD branch selects this
predicate for the default precision at every batch size, matching what
ggml_cuda_flash_attn_ext() dispatches to there. Without the fix the HIP build
would trade one hardcoded head size restriction for another.
NVIDIA is unaffected either way. Volta and later route 192/128 through the mma
or wmma predicates, and on Pascal the Q->ne[1] <= 8 decode case is diverted to
vec f32 before this predicate is consulted.
The DSA attention kernel used the shared cublas handle without binding
it to the backend's stream, so its Q.K / P.V GEMMs ran on a different
stream than the gather and softmax kernels. The softmax could then read
the score buffer before the GEMM wrote it, picking up stale (NaN) values.
* vulkan : use ggml_row_size for types with a per-row scale
Types that declare a row_meta_size store a per-row scale ahead of the row's
blocks, so a row is not ggml_type_size()*ne/ggml_blck_size() bytes. This
under-sized src0 in the four quantized mat-mul paths, and made
ggml_vk_dim01_contiguous() report such a tensor non-contiguous, which in turn
made supports_op reject it. No change for row_meta_size == 0.
* vulkan : add IQ4_KS and IQ4_KT support
A row of these types is one f32 scale followed by the row's blocks, so rows are
not a whole number of blocks apart and the usual block-indexed addressing does
not work. They are read through a uint32_t alias of binding 0 and addressed by
word; types.comp holds the alias, the stride and the decode, so each shader only
expresses its own addressing and no push constant layouts change.
Covers to_fp16, get_rows, mul_mat_vec (incl. MUL_MAT_ID) and scalar + coopmat1
mul_mm. get_rows addresses by row rather than through nb01/02/03, which cannot
express a per-row scale, so supports_op accepts only a contiguous src0 for these
two types. coopmat2 is excluded because coopMatLoadTensorNV addresses through a
uniform grid tensor layout, which cannot describe the row prefix; the two
mat-mat getters return nullptr there and the callers fall back to F16.
* tests : add IQ4_KS/IQ4_KT decode validation
Re-implements in C++ the indexing each of the four shader families uses and
diffs it against ggml's to_float over several row and block counts. CPU only: it
validates the format transcription, not the compiled shaders.
* Load standalone Qwen3.5 MTP GGUFs passed with -md
A predictor-only MTP GGUF reports the full block count (n_main +
nextn_predict_layers) but only ships the NextN block, so loading one
with -md failed:
check_tensor_dims: tensor 'blk.0.attn_norm.weight' not found
create_qwen35_tensors() and create_qwen35moe_tensors() create every
main block as required. Detect the predictor-only case the same way
create_step35_tensors() does and mark the absent blocks
TENSOR_SKIP|TENSOR_NOT_REQUIRED.
Qwen3.5 also has to use the common MTP package contract, otherwise the
predictor-only GGUF is never classified as a companion, and the target
is not classified TARGET_ONLY - which is what makes it export the
hidden states the companion consumes.
The remaining two hunks cover cases the above newly reaches: a
predictor-only GGUF passed as -m now loads far enough to abort in the
graph builder, and its empty main blocks reach split_recurrent_tensors()
under -sm graph.
* Qwen3.5 MTP: require q_proj in predictor-only GGUFs, check companion arch
Review follow-up.
A dense NextN block loads q_proj as optional because it can be shared
with the last main block. A predictor-only GGUF has no main blocks, so
one built that way loaded with wq == nullptr and then hung. Require the
tensor in that case so the load fails naming it. eh_proj, attn_q and the
MLP are all optional on that block, so the tail probe stays on enorm,
which is required - the comment there said only eh_proj.
Adding Qwen3.5 to the common MTP package contract also made
common_speculative_has_recognized_mtp_companion() accept any GGUF
classified COMPANION, with no architecture check of the kind the Step
and DeepSeek branches have. Add it, plus the predictor count. Dense and
MoE are separate architectures, so the comparison is on the arch itself.
OpenAI-compatible clients (e.g. pi coding agent) send
max_completion_tokens for the output token cap on custom
openai-completions providers. The server only read n_predict and
max_tokens, so the cap was silently dropped and n_predict fell back
to -1 (unlimited). This allowed runaway generations of 40k+ tokens
on long agent sessions.
Matches upstream llama.cpp behavior where max_completion_tokens is an
alias of n_predict (tools/server/server-schema.cpp).
* Update parameters.md
- Add new parameters
- Update modified parameters
- Add graph parallel new arch
- Fix some words case
* Update README.md
- Update supported models list
- Add the new features
* model: Ling-3.0 (bailingmoe3) runtime support
* model: Ling-3.0-tiny support
---------
Co-authored-by: Joel Farthing <262452229+joelfarthing@users.noreply.github.com>
The draft block was seeded at last_target_pos (the newest committed
feature row = id_last's predecessor), placing the whole block one
position early vs mainline's [id_last @ n_past, ...] convention and
colliding the seed with the newest cross-KV row. Shift the common-side
batch and the graph-side SWA mask base coherently.
* Adding Muse-Glimmer support
* Different rms_eps for post norm ops
* Need attn_post_norm split for Muse-Flimmer
* WIP: split mode graph
* Forgot this file
* Clean it up
* Minor
* Muse-glimmer: Slightly better split mode graph (+2% TG)
build_llama passes nullptr for the rope_factors_in argument, so
rope_freqs.weight never reaches ggml_rope_ext and llama3 rope frequency
scaling is not applied. The LLAMA_SPLIT_MODE_GRAPH branch inside
build_std_attention falls back to model.layers[il].rope_freqs, so only
the standard path is affected.
The DFLASH case was inserted into the NEOX fall-through group in
llama_rope_type(), so the 39 cases above it now return LLAMA_ROPE_TYPE_NORM
instead of LLAMA_ROPE_TYPE_NEOX.
Move DFLASH into the NORM group so it keeps its intended rope type and the
others fall through to NEOX again.
* CUDA indexer topk: this is better for PP
* Don't overstep
* Cleanup
* Allow Q8_0 cache in the CUDA DSA implementation
* DS4: do not cast caches to f32
* Fix massive inefficiency in CUDA Q->f32/f16 and f32/f16->Q copies
* Re-enable -ictk | --indexer-cache-type-k
* CUDA indexer topk: this is better for PP
* Don't overstep
* Cleanup
* Allow Q8_0 cache in the CUDA DSA implementation
* DS4: do not cast caches to f32
Samuel Oliveira Alves
Coenie Beyers
Joel Farthing
Kawrakow
Guy Barel
dungquixote42
Yap Sok Ann
Thireus ☠
Jun Yamog
Marian M.
rumas77
mb8565
abc-nix