Pr26081 - #32
Open
kpouget wants to merge 2021 commits into
Open
Conversation
* hexagon: add VISION RoPE support * hexagon: support RoPE on strided half-dim views for all modes * hex-rope: decouple src0 DMA copy size from row stride * hex-rope: support non-contiguous dst for RoPE * hex-rope: fix dst spad pitch for non-contiguous dst
* cuda: fix snake fusion type predicate, a and inv_b are F32 The matcher required a->type == x->type while launch_snake reads both as const float *, matching the CPU and Metal contract where a and inv_b stay F32. F16/BF16 chains never fused and fell back to the naive path, and a hypothetical all F16 chain would have read F16 bits as float. Aligns the predicate and the comment with ggml-cpu.c * cuda: reject snake fusion on non-contiguous operands The kernel reads x[idx] and a[c] / inv_b[c] linearly, so a non-contiguous view passing the matcher would silently read wrong data. Mirror the contiguity guard already present in the CPU, Vulkan and Metal matchers.
CUDA is compiled with fast math and AMD/HIP is not — this flag lets AMD use fast math too. We can't use -ffast-math: it implies -ffinite-math-only, which won't compile (ggml uses INFINITY for masking) and produces NaNs. -funsafe-math-optimizations gives the speedup without the NaN problems. Co-authored-by: Mark Caldwell <mark@cloudhands.ai>
* metal : add CONV_2D_DW (depthwise 2D convolution) support * test : add perf cases for CONV_2D_DW * metal : use 3D dispatch for CONV_2D_DW kernel * metal : add channel-tiled CONV_2D_DW kernel for non-contiguous layouts * metal : simplify CONV_2D_DW dispatch and trim comments * metal : merge duplicate CONV_2D_DW pipeline getters * tests : add F16 CONV2D_DW tests * cpu : fix F16 kernel support for CONV_2D_DW * tests : remove commented-out CONV_2D_DW test block --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
The first one avoids relying on compile to optimize local memory away, and the second is cheaper than issuing control flow statements
* Use smart pointers in test_case::eval This makes it consistent with other methods of `test_case`. * Use smart pointer in show_test_coverage also * Also use smart pointers for backends
…#25480) * meta: add hard emphasis on agents not writing descriptions/comments Add a block in AGENTS.md to emphasize that agents are forbidden, under any circumstances, to post comments or pull request descriptions on behalf of the user. * Add example * Move examples to examples * White space
…(#24093)
A model whose chat template parses at init but fails parser generation
at apply time (e.g. uses {% call %}) throws std::invalid_argument from
common_chat_templates_support_enable_thinking(), which ran outside the
try/catch guarding common_chat_templates_init(). The throw was uncaught
and llama-cli aborted (SIGABRT) instead of failing to load. Moved the
probe inside that try/catch so an apply-time error fails load the same
way an init parse error does.
Signed-off-by: Jesse LaRose <jesse@taey.ai>
* hexagon: tile wide rows in pointwise unary ops to avoid VTCM overflow * unary: reject permuted tensors for now (not used by models) * hex-unary: replace divs with fastdiv * hex-unary: add vtcm layout and host computed kernel params * hex-unary: move fastdiv init into kernel params * hex-unary: add specialized thread functions to improve generated code * hex-unary: tracing instrumentation for unary ops * hex-unary: factor out hvx kernels, streamline and remove more duplication * ggml-hexagon: fix std::min collision with Windows min macro * hex-cmake: make lto build happy --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
…rgsort() to reduce temporary buffers memory usage (#24776) * ggml : process data in smaller chunks in CUDA ggml_top_k() implementation to reduce temporary buffers memory usage * ggml : allocate tmp_dst only only once before the loop * chore : whitespaces Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * ggml : use chunked processing in both CUDA CUB top-k and argsort implementations * chore : separate argsort_f32_i32_cuda_bitonic() call from return statement Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * chore : replace ternary operators with min/max --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* llama-batch: add unit test * fix win32 builds * add not implemented assertion in unused methods * remove unreachable code
* ggml-et: Add performance logging * ggml-et: Quants helpers * ggml-et: Add MUL_MAT kernel * ggml-et: Add ROPE kernel * ggml-et: Add RMS_NORM kernel * ggml-et: Add GLU kernel * ggml-et: Add SOFT_MAX kernel * ggml-et: Add GET_ROWS kernel * ggml-et: Add CONT kernel * ggml-et: Add SET_ROWS kernel * ggml-et: Add MUL_MAT_ID kernel * ggml-et: Build et kernels as part of ggml * ggml-et: Embed kernels with fs fallback * ggml-et: Build fixes * ggml-et: Add MUL_MAT F32xF32 op * ggml_et: Add MUL_MAT_ID op * ggml-et: Disable offloading for debug * ggml-et: Refactor out block ops * ggml-et: ggml backend API changes * ggml-et: Add RESHAPE/TRANSPOSE to supported * ggml-et: Add CONT_F16 * ggml-et: Add supported ops doc * gglm-et: Initial doc * ggml-et: Remove runtime import hacks We can now import the runtime by a simple find_package(), so we can cleanup the CMakeLists.txt. * ggml-et: Fix GET_ROWS kernel Fix lost batch dimension. Also clean vibe-comments. * ggml-et: Fix SET_ROWS kernel Remove incorrect broadcasting guard. * ggml-et: Use custom instruction for fp32->fp16 * ggml-et: Vectorize set_rows fp32->fp16 * ggml-et: Fix ROPE kernel (yarn) ggml-et: fix et_logf WIP: Fix ramp WIP: fix ROPE! * ggml-et: Better sinf * ggml-et: Fix SOFT_MAX Add `max_bias` and `sink` support. * ggml-et: Fix CONT Reorder from contiguous write to read with atomic stores. * ggml-et: Fix elmap kernel Remainder handlin * ggml-et: Fix MUL_MAT MUL_MAT_ID remainders * ggml-et: Fix ET-SOC reference * ggml-et: Fix embed kernels scripts for old python This allows GGML-ET to build on pre-3.8 python. * Add sysemu support with compile time flag `-DGGML_ET_SYSEMU=ON` (crc-org#6) * Example using ET-Soc-1 emulator configuration Example usage: ```bash cmake -B build -DGGML_CUDA=OFF -DGGML_ET=ON -DLLAMA_CURL=OFF -DGGML_CCACHE=ON cmake --build build --config Release -j $(nproc) time ./build/bin/test-backend-ops ./build/bin/llama-server \ --model Qwen3-0.6B-Q8_0.gguf \ --alias Qwen3-0.6B-Q8_0 \ -fa 0 \ --ctx-size 1024 \ --no-warmup \ --host 127.0.0.1 \ --port 8080 ``` * build: proper dep tracking for kernels * support host using MOLD linker * initial multi core GET_ROW F32 implementation * vectorized q8 dequant * wip: cland warning clenaups and initial logging refactor * wip: message default message cleanup * chore: message cleanups * cmake cleanup * migrate to use platform provided functions * cmake back into subdir * support et_print() in kernels * fix: repair kernel building * perf: operations run async by default * debug: proper kernel dep tracking and error detection on kenrel launch * fix: kernel binary dep tracking and fixing get_rows_f32 erroring * perf: back to doing async kernel runs by default * perf: vectorize and parallel device memset * merge matmul work * misc: align allocation and enable all offload * misc: delete deadcode and respect memory limits * fix: repair tensor debug print * fix: loosen RMS_NORM op percision * feat: Q4_0 GET_ROWS * perf: FP32 MUL_MAT using TensorFMA * update limitations * perf: redue L1 load in compute_block_dot_product_q8_0 * feat: save kernel mapping (name to id) when profiling is enabled * chore: memops cleanup * perf: parallelize softmax by rows * perf: vectorize 2nd phase of softmax * perf: ban GET_ROWS from offloaded * perf: vectorize and non-atomic for eltwise ops and sub support * perf: vectorize normal rope * perf: glu runs in parallel * merge: manually merge saqib's work on kernel fixes * perf: more vectorized RoPE * perf: parallelize mul_mat_id * perf: parallelize set_rows_f32 * perf: vectorize softmax * feat: support kernel fusion and fuse RMS_NORM + MUL * fix: mostly resolve test-backend-ops failure in SOFT_MAX and ROPE * fix: bump max rope dims for gemma * feat: GeGLU and SCALE support to fully offload Gemma * perf: faster device memset * feat: get_rows supporting Q4_K and avoid cont cache coherent issues * better F32 MM * feat: NORM for ET backend * feat: SQR for ET backend * feat: UNARY on ET * feat: el_map support broadcasting for ET * feat: SUM_ROWS in ET backend * feat: more ops in ET backend * feat: WKV* operators in ET backend * perf: parallelize operators across cacheline instead of row * perf: parallelize get_rows on cacheline * wip: baseline FlashAttention for ET backend * wip: enough FA and CPY f32->f16 to run llama 3.1 fully offloaded with FA on * feat: f16 x f16 -> f32 MM using matrix engine * wip: f16 FlashAttention using matrix engine * wip: clean up * feat: barriers * perf: optimize FA_F16 in ET * perf: vectorize pack_k_for_transpose16 * perf: prefetch next loop matrix tile * perf: FlashAttention 2nd MM uses TensorFMA and optimizations * cleanup: flashattention reorg * perf: optimizations and fixes * feat: L2SCP API and make FlashAttention support DV = 256 for gemma * perf: parallelize norms beyond single row * feat: GATED_DELTA_NET support and relaxed L2_NORM requirment * feat: loosen RMS_NORM, NORM, ROPE contingous req too * feat: repeat supports brocasting on dim 0 and loosen cont check * feat: FILL and DIAG operator * feat: loosen UNARY support chcek * feat: TRI support * feat: SOLVE_TRI support * feat: basic SET support * feat: loosen CONT req * perf: fp16_to_fp32 use ASM * feat: IMROPE support * feat: PAD support * feat: global barrier * fix: view must live on the same backend as backing tensor * feat: relax CONCAT in ET backend * feat: dead simple CUMSUM implementation * feat: basic SSM_CONV support * feat: loosen CONCAT req * feat: relax GATED_DELTA_NET and add SET support proper * cleanup: cleanup LCM math * feat: SWIGLU single input * feat: SSM_SCAN support * feat: el_map supports non aligned tensors in best effort * feat: basic GROUP_NORM support * feat: loosen MUL_MAT capablities slightly * feat: loosen MUL_MAT and GET_ROWS and add IM2COL * feat: special case for softmax 1x1x1x1 * feat: loosen SOFT_MAX req in ET backend * fix: el_map unaligned acse fixes * perf: optimize zero_acc_vec in flash_attn_ext_f16_me * perf: use hart 1 for packing in MM and FA for FP16 * feat: kernel semaphore * perf: better instruction sequence in FlashAttention * fix: gated_delta_net with proper masking * perf: better parallelization for GATED_DELTA_NET * perf: parallelize SSM_CONV over nr * perf: vectorize SSM_CONV * perf: optimize MUL_MAT for q8 * feat: support Gemma 4 * fix: support multi-device * feat: broader GLU support * feat: unary ops supports view * fix: repair fp16 MM using matrix engine * perf: handle large N GEMV better * perf: better q8_0 MM * perf: better set_rows * add back deleted files * fix: repair after merge * feat: POC version of uberkernel * feat: RMS_NORM in uberkernel * feat: add more kernels into usage * chore: clean up uberkernel compilation * perf: faster flash attention * perf: opt flash attention for large seq length * feat: loosen op bounds. clamp and mean support * perf: vectorize ssm_scan * perf: slightly faster FA * perf: FlashAttention parallel MM and load * perf: fuse Q8 MM and ADD * feat: basic conv kernel for ET * softMAx_test * set_rows_f32 * get_rows and cont * testing * set_rows_exp * Junk addition * Narrowing the issue * Update flash_attn_ext_f16_me.c Focusing FA_ext_f16_me * test * Eviction updated * Detailed cache eviction debug * mulmat * removeal of `BUILD_FOR_UBERKERNEL` flag * cleaning... * fix: balance FCC0 count * feat: implement mul_mat and mul_mat_id for Q4_0 type * optimize uberkernel plan upload * add mul_mat q4 into uberkernel * enable gating flush to just uberkernel * update docs for ET * update op support for ET * et-backend: optimize Q4_0 and Q8_0 mul_mat_id row accumulations * et-backend: specialize mul_mat_id kernels for Q4_0 and Q8_0 * et-backend: fix RoPE YaRN corr_dim formula and handle degenerate inputs * test-backend-ops: add DeepSeek-V2-Lite RoPE test coverage * et-backend: add Q4_0 mul_mat matrix-engine kernel using TensorFMA32 * et-backend: vectorize Q4_0 matrix-engine dequantization * et-backend: support hybrid matrix/vector engine execution for Q4_0 mul_mat tail * et-backend: run partial-N tiles on matrix engine for Q4_0 mul_mat * et-backend: route Q4_0 mul_mat N < 53 to vecdot for better prefill latency * Update uberkernel.c * Update unary_f32.c * gemma 4 * bisect gemma4: enable scale_f32 only * bisect gemma4: +rms_norm_f32 * bisect gemma4: +rms_norm_mul_f32 * bisect gemma4: disable rms_norm_mul_f32 -- BREAKS OUTPUT * bisect gemma4: +rope_f32 (skip rms_norm_mul) * bisect gemma4: +el_map_f32 * bisect gemma4: +softmax_f32 * bisect gemma4: +get_rows_f32 * bisect gemma4: +glu_f32 * bisect gemma4: +mul_mat_f32 +mul_mat_f32_matrix_engine * bisect gemma4: +mul_mat_f16 +mul_mat_f16_matrix_engine * bisect gemma4: +mul_mat_Q8_0 +mul_mat_Q4_0 * bisect gemma4: +flash_attn_ext_f32 +flash_attn_ext_f16_me * bisect gemma4: +mul_mat_id_f32 * bisect gemma4: +sum_rows_f32 * bisect gemma4: +cont_f16 * bisect gemma4: +fill_f32 * bisect gemma4: +unary_f32 (all ops re-enabled except rms_norm_mul) * Update rms_norm_mul_f32.c * bisect2 gemma4 n64: +scale_f32 only * bisect2 gemma4 n64: +rms_norm_f32 +rope_f32 * bisect2 gemma4 n64: +rms_norm_mul_f32 (with ET_UBERKERNEL eviction fix) * bisect2 gemma4 n64: +el_map +get_rows +glu +softmax (skip rms_norm_mul) * bisect2 gemma4 n64: all ops enabled except rms_norm_mul * bisect2 n64: test unary+cont+fill+sum_rows (no mul_mat/flash_attn) * bisect2 n64: +mul_mat_f32 +mul_mat_f32_matrix_engine * bisect2 n64: +mul_mat_f16 +mul_mat_f16_matrix_engine * bisect2 n64: +mul_mat_Q8_0 +mul_mat_Q4_0 * bisect2 n64: +mul_mat_Q8_0 only (disable Q4_0) * bisect2 n64: +mul_mat_Q4_0 only (Q8_0 breaks) * bisect2 n64: +mul_mat_id +flash_attn_ext (skip Q8_0) * run-3: matmul + rms_norm_mul * run-4 * Revert "run-4" * run5 * changes after cleanup * cleanup before upstream * restrict changes into ET backend * move kernel embedding from Python to CMake * move uberkernel gen into CMake * apply clang format * update CMake style * update to match C and C++ style * use source ggml and quant headers instead of ET's * MROPE support * absorb view ops into same branch as none * fix bad rebase * add marty1885 to codeowners * oops * remove redundant newline * fix CI editor warnings --------- Co-authored-by: Vidas <vidas@nuolat.lt> Co-authored-by: Gianluca Guida <glguida@tlbflush.org> Co-authored-by: Gianluca Guida <gianluca@nekko.ai> Co-authored-by: ubergarm <leimgrub@gmail.com> Co-authored-by: SaqibAkram-10xE <saqib.akram@10xengineers.ai> Co-authored-by: Rehan Qasim <rehan.qasim@10xengineers.ai>
…as, remove raw_k repeats in DeepSeek V4 (#25370) * llama : make all KQ masks (except the lightning indexer one) f16 if FA is used and remove zero attention bias in DeepSeek V4 * llama : remove dead code that repeats unified raw_k cache for each stream in DeepSeek V4 - no longer needed as raw_k is always non-unified. --------- Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
downloadConversation serialized activeMessages, the root -> currNode path, so exporting a conversation with edited or regenerated messages dropped every alternate version and kept only the selected one. Fetch the whole message tree via getConversationMessages so the export carries all message versions, matching the multi-conversation export path which already did this. Keep the active conversation as the header source to preserve an up-to-date currNode. Forks are separate conversations, each with its own convId, and are exported on their own.
* ggml : bump version to 0.16.0 (ggml/1559) * sync : ggml
* llama-cli: fix crash on wrong server base url by catching exceptions and graceful exit * review: leaner catch group: json error and standard exception
* server: improve tools, remove apply_diff * improve edit tool * add tools_io abstraction * add tools_io_basic * fix build * move utils to class member * add const
* server: remove loading.html * apply ui changes
Co-authored-by: example name <example@example.org>
* mtmd: deepseek-ocr v1 multi-tile dynamic resolution + unified image-preprocessors for both versions (ds-ocr v1 and v2) * remove hacky API * fuse row into a long image * almost working * adapt to new preprocessor api * rm debugging printf * improve * mtmd: dsocr-tiles fixes (#25481) * ds-ocr img-preproc fuse_row tile-drop fix for multi rows and columns images * mtmd drop the duplicate redundant img_end * deepseekocr graph simplify CLS broadcast cleanup * test-deepseek-ocr: relax v1 single-view tolerance; drop trailing prompt space; make DRY opt-in and n_predict model-specific (#25486) --------- Co-authored-by: Saba Fallah <10401143+sfallah@users.noreply.github.com> Co-authored-by: Saba Fallah <sabafallah@gmail.com>
…ntion path (#25880)
* sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path
The scale was uploaded with an async memcpy sourced from a stack local. On the
in-order queue that copy is ordered behind the K/V staging kernels; once n_kv is
large enough (>= ~26k observed on Arc Pro B70) the staging outlives the host
stack frame and the copy reads recycled memory, feeding the SDPA a garbage scale.
Output then collapses to a single repeated token and the KV cache is poisoned
for the rest of the session.
Short contexts win the race by accident, and test-backend-ops caps
FLASH_ATTN_EXT at kv=1024, which is why CI never caught it. The previous
device_count > 1 wait_and_throw() gate (and reverting it, PR #25741) fixes the
symptom only by keeping the frame alive across the copy at the cost of a host
sync on every FA call.
Fix: cache one device scalar per (device, value) -- the scale is constant per
model -- and upload it synchronously once. The single-device fast path (no
per-call host sync) is then safe: every device-side hazard already serializes
on the in-order queue. The multi-GPU conservative wait is kept unchanged.
Also:
- GGML_SYCL_FA_ONEDNN_MAX_KV env (0 = unlimited): optional n_kv ceiling that
routes very long sequences to the native FA kernel.
- test-backend-ops: FLASH_ATTN_EXT F16 cases up to kv=65536 (Qwen3.6-27B
geometry hsk=hsv=256 GQA 6, and hsk=128 GQA 4), closing the kv=1024 blind
spot. Note the race itself needs a live multi-op pipeline to reproduce;
single-op runs pass even on broken builds.
Verified on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072: output
byte-identical at temp 0 to the native FA path through 32k-deep prefill, with
prefill depth-flat at 820-840 t/s (vs 340-350 native at 32k depth).
Assisted-by: Claude Fable 5
* sycl: handle GGML_SYCL_FA_ONEDNN_MAX_KV like the other runtime env vars and document it
Review feedback on #25880:
- read the variable once at backend init into g_ggml_sycl_fa_onednn_max_kv via
ggml_sycl_get_env, and print it in the startup env listing (-lv 4 shows it)
- document GGML_SYCL_FA_ONEDNN and GGML_SYCL_FA_ONEDNN_MAX_KV in the SYCL.md
runtime table
Also trim the added FLASH_ATTN_EXT cases to kv={4096,16384}: the 32768/65536
shapes exceed the legacy NMSE threshold on both the oneDNN and native kernels
(long-sequence fp16 accumulation drift, present before this PR) and would fail
CI for an unrelated reason.
Assisted-by: Claude Fable 5
* sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled
Assisted-by: Claude Fable 5
* sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly
Assisted-by: Claude Fable 5
* Update ggml/src/ggml-sycl/fattn-onednn.cpp
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* sycl: write the SDPA scale from a kernel instead of caching it
The per-(device, value) scale cache was a function-local static
unordered_map with no synchronization, so concurrent backend instances
could access and rehash it at the same time.
Write the scalar with a single_task instead. The value is captured into
the command, so no host memory has to outlive the call -- which is what
the use-after-return fix needed in the first place. That removes the
shared container, the leaked device allocation and the string key, and
it also closes the remaining async-memcpy-from-a-stack-local on the
first flash-attention call.
Ordering does not rely on timing: the queue is created with
sycl::property::queue::in_order and the dnnl stream wraps that same
queue, so the write completes before the SDPA reads the scalar. The
multi-GPU wait_and_throw() branch is unchanged.
Also drop the <cstdlib> include, which is unused.
Assisted-by: Claude Opus 5
---------
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
…etal iGPU is used
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* spec: add DSpark speculative decoding DSpark (DeepSpec, 2026) on top of the merged DFlash drafter. It reuses the DFlash encoder/decoder graph, target feature extraction and KV-cache injection, and the verify/accept path unchanged; the draft model is a new "dspark" arch adding a low-rank Markov head (markov_w1/w2) and an optional (unused here) confidence head. No new public APIs. The proposal is the only change: the block is anchor-first (position 0 already predicts the first draft) and the decoder graph applies a semi-autoregressive, previous-token conditioned logit bias in-graph, chained per block position: logits'(i) = logits(i) + markov_w2 . markov_w1[prev(i)] prev(0) = the block's anchor token, prev(i>0) = argmax(logits'(i-1)) vectorized across all blocks in the batch; the anchors are fed through a dedicated graph input (token 0 of every block). Greedy stays lossless (verify unchanged, same as DFlash). - new arch "dspark" (llama_model_dspark : llama_model_dflash, reuses the graph, loads the markov/confidence tensors; shares the target's embed/lm_head). - Qwen3DSparkModel converter. - new spec type "draft-dspark" (common_speculative_impl_draft_dspark : common_speculative_impl_draft_dflash, overrides draft() only: submits whole anchor-first blocks and greedily reads back the biased logits). * spec: read draft block size in the dflash impl * docs: add DSpark section to speculative.md * spec: keep dspark block size read in the dspark impl * dspark : add TODOs for incomplete parts - confidence head is loaded but not used yet - confidence-scheduled prefix pruning is not implemented - the in-graph Markov chain is greedy-only - only Qwen3 backbones are supported for now (also noted in docs) * spec: fold DSpark into the DFlash arch Address review: drop LLM_ARCH_DSPARK and the dspark.block_size / markov_rank GGUF keys. A DSpark draft now converts to a DFlash GGUF; the Markov head tensors are detected by presence (like eagle3 d2t), block_size is read from the existing dflash.block_size key, and the block anchors are taken as a strided view of the decoder's token input instead of a separate graph input. * spec: add confidence-based draft pruning for DSpark The DSpark confidence head predicts per-position acceptance of the drafted block. --spec-draft-conf-min truncates the block at the first position below the threshold (default 0 = disabled). * fold the dspark impl into dflash, selected by spec type * address review comments * dspark: clean up and improve naming * update readme * remove trailing whitespace * dflash: draft full n_max blocks, defer dp.n_max to the central truncation The DSpark markov head views the draft batch as a uniform [n_seqs x block] grid, but the per-seq dp.n_max clamp could produce blocks of different sizes, silently corrupting the strided views and the resulting logits. Drop the clamp and always draft the full n_max block for every sequence: dp.n_max is already enforced by the central truncation in common_speculative_draft(), the same way eagle3 handles it. Co-authored-by: Zaire404 <3147879462@qq.com> * dflash: assert the markov head block-uniformity invariant, require the conf head With the draft batch always submitting equal-size n_max blocks, a non-divisible token count can only mean the batch was split across ubatches or a caller broke the layout - fail loudly instead of silently dropping the markov bias. The block_drafts > block_size early return stays: worst-case graph reserve passes legitimately build with n_seq_tokens > block_size. Also make conf_proj required when the markov head is present: the confidence head is part of the DSpark checkpoint format, and a missing head would otherwise leave --spec-draft-conf-min silently reading stale embeddings instead of confidences. Co-authored-by: Zaire404 <3147879462@qq.com> * dspark: fold conf_min into p_min p_min and conf_min express the same thing - the minimum predicted survival probability for a drafted position - differing only in how the estimate is obtained: token probability for regular drafters, the trained confidence head for DSpark. The DSpark readback never used p_min, so reuse it for the confidence threshold and drop the separate --spec-draft-conf-min flag. Both defaulted to 0 (disabled), so behavior is unchanged. Co-authored-by: Zaire404 <3147879462@qq.com> * dflash: note the confidence broadcast workaround Requested in review: the ggml_repeat only adapts the [1, n_tok] confidences to the n_embd-wide embd_nextn transport so that llama_get_embeddings_nextn can be reused - not a placeholder. Co-authored-by: Zaire404 <3147879462@qq.com> * cont : clarify [no ci] --------- Co-authored-by: Ruixiang Wang <wangruixiang07@outlook.com> Co-authored-by: Zaire404 <3147879462@qq.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…22675) * ggml-cuda: add chunked SSD matmul for Mamba-2 prefill acceleration * cuda: added SSD CICD fixes for CUDA / HIP / MUSA / MSVC. * ggml-cuda: review comments fixed. * ggml-cuda: Fuse M matrix materialization into pre_matmul kernel and enabled test. * ggml-cuda: test updates and fixes * ggml-cuda: test updates to remove hardcoding of tensor initialise data limits. * ggml-cuda: ssd minor review comment fixed. * ggml-cuda: ssd minor CICD fixed. * CUDA SSD: Fixes correctness by promoting s0_stride_seq to int64_t, improves memory coalescing in ssm_ssd_prepare_dt_kernel, and boosts efficiency by merging B_weighted and C_scaled; also addresses prior review comments. * cuda: fix sdata read-write race in prepare_dt fallback scan loop
* vulkan: add iq4_nl support back to FA I was originally concerned about wasting shared memory on the LUT, but it's small and unlikely to matter in practice. Also support q1_0 for non-coopmat2. Fixes #23681 * remove q1_0 FA support
* llama-graph: set_outputs to t->view_src * change set_output to GGML_ASSERT about views not being outputs * sampler : avoid views in outputs * cont : fix dist sampler * cont : consistent logits handling * ggml : set output of view src * graph : simplify set_outputs() * cont : cleanup Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Gaurav Garg <gaugarg@nvidia.com>
Co-authored-by: dev-tinker <dev-tinker@users.noreply.github.com>
* mtmd : add Nemotron 3 Nano Omni support (parakeet) This commit adds support for the subsampling and encoder part of Nemotron Nemo 3 omni model. The Parakeet subsampling/encoder were taken from parakeet.cpp which is currently a pull request against whisper.cpp. I've tried to copy the code a close as possible to hopefully enable easy patching between the these two project later. Refs: ggml-org/whisper.cpp#3735 * mtmd : generate rel pos tensor in graph instead of in conversion [no ci] This commit removes the generation of the relative positional tensor in the model conversion script and instead computes it in the encoder graph. This is only done for the window of positions required for the current audio sample. * mtmd : add clip_get_model to clip API [no ci] This commit adds a function to get access to the clip_model. It also removes the two functions clip_get_mel_filter_tensor, and clip_get_window_tensor(const struct clip_ctx * ctx) which can now use clip_get_model to access the model tensors that it needs. * mtmd : read mel_filters and window into hparams * mtmd : use set_input_f32 lambda [no ci] * mtmd : add better asserts for mel_filters and hann window [no ci] * mtmd : add missing size_t cast * mtmd : change type of pad to size_t * mtmd : zero initialize samples_padded * mtmd : remove unsued ctx member from parakeet preprocessor * mtmd : make log_mel_spectrogram_parakeet_worker_thread private static * mtmd : sync/update parakeeet impl with latest whisper.cpp This commit updates the parakeet code in mtmd to reflect the latest updates to parakeet.cpp in whisper.cpp. A follow up commit will address the currently hardcoded dw_pad and see if we can add n_conv_kernel as a model metadata field. * mtmd : add audio_conv_kernel_size to model conversion This commit updates the model conversion to read the conv_kernel_size field from the sound_config section of the models config.json file. It then uses this field instead of the hardcoded values in parakeet.cpp. * mtmd : cleanup [no ci] * conversion : call super().filter_tensors [no ci] * do not discard result of super filter_tensors * mtmd : use build_mm instead of ggml_mul_mat * mtmd : use build_ffn * mtmd : move and reuse get_vector lambda * mtmd : use build_inp_raw for parakeet * mtmd : throw exception in get_scalar instead of assert * mtmd : fix std::min call * mtmt : use .c_str in throw clause in get_vector * mtmd : check for F32 type and non-empty tensor in get_vector The get_vector lambda is used by get_scalar but also standalone to read in the mel_filters and the window data. Therefor we are not checking for 1D tensors but allowing multiple dimensions. We do have a check in get_scalar to verify the size of the vector. * mtmd : replace hardcoded 1101 for n_tokens_real * mtmd : assert subsampling_factor is 8 This commit adds an assert of the parakeet subsampling factor to check that it is 8. The motivation for this is that this model currently has three convolutions with a stride of 2. If the underlying model updates the subsampling factor these convolution operations will need to be updated and this will produce and error if this occurs. * mtmd : remove unused ggml_tensors attn_pos_w and mm_norm_w * mtmd : remove single thread path This commit removes the single thread path which was a left over from the original parakeet.cpp where n_threads is configurable. * fix some security issues --------- Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
… (#26189) The Adreno KQ/KQV image1d kernels (ggml_cl_mul_mat_kq_kqv_adreno) ignore dim 3 entirely: the sub-buffer covers only nb02*ne02 bytes and the kernel receives no ne03/ne13/nb03/nb13 arguments. With the unified KV cache, multi-sequence batches (e.g. llama-perplexity with its default -b 2048, n_seq=4, or a multi-slot llama-server) present KQ/KQV as 4D tensors with ne3 = n_stream, so every stream past the first reads the first stream's K/V and produces garbage. Flash attention masks the bug where it is enabled; devices where FA is declined (e.g. Adreno 740) hit it with default settings. Route ne03/ne13 > 1 to the general path, which handles dim 3, and honor view_offs when creating the sub-buffers (currently always 0 for tensors reaching this function, but the function would silently misread any future view). Llama-3.2-1B-Instruct Q4_0, wiki.test.raw, 8 chunks, -ngl 99: - Adreno 740, default: PPL 1817.64 -> 15.61 - Adreno 740, -fa 0: PPL 1941.64 -> 15.61 - Adreno 840, -fa 0: PPL 1943.90 -> 15.50 - single-stream (-b 512) results unchanged (15.6090) - test-backend-ops -o MUL_MAT on 740: identical before/after (909 OK, 12 pre-existing q6_K failures)
…recurrent-state-rollback test (#25931) * Add overlap glu variant to support all archs, fix recurrent-state-rollback test * format * Fix all arch overlapped ranges * format * diagnose bus error on apple ci * More testing * more testing * more targeted testing * Fix bug in alignment for > 4gb buffer offsets * Fix bug in view offsets * Try avoiding multi_buffers * not fixed yet, more logging :( * Handle edge case in set_rows * Try looking at view source * Skip deepseek32 for now and clean up trace infrastructure * simplify skipping * last cleanup * actually final cleanup * update handling of overlap * format * try skipping other failing model
…2) (#25980) * model: add NextN/MTP speculative decoding support for GLM_DSA (GLM-5.2) Adds GLM-5.2 NextN/MTP as a --spec-type draft-mtp target: nextn tensor loading via the qwen35moe/step35-style presence probe, a graph_mtp builder (enorm/hnorm/eh_proj + dense MLA + sigmoid-gated MoE with shared expert + shared head with fallbacks, _s scale tensors passed for NVFP4), t_h_nextn extraction in the trunk graph, and MTP-context KV setup: the draft head runs dense MLA, so the MTP context uses a plain attention KV cache holding only the nextn layer(s) (same pattern as the hybrid Qwen3.5 MTP context) while the main context keeps the DSA cache, now filtered to trunk layers only. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * convert : support --mtp/--no-mtp export for GlmMoeDsaForCausalLM (GLM-5.2) Opt GLM-5.2 into the supports_mtp_export contract (post-#25641 shape, mirroring HYV3Model/Step35Model): --no-mtp drops the appended NextN block (blk.78) and its nextn_predict_layers KV; --mtp keeps only the NextN block plus shared embeddings/norm/lm_head. Default (bundled) output is unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
|
[APPROVALNOTIFIER] This PR is NOT APPROVED This pull-request has been approved by: The full list of commands accepted by this bot can be found here. DetailsNeeds approval from an approver in each of these files:Approvers can indicate their approval by writing |
|
Important Review skippedToo many files! This PR contains 2354 files, which is 2254 over the limit of 100. To get a review, narrow the scope: Upgrade to a paid plan to raise the limit. Usage-priced reviews support at most 300 files. ⚙️ Run configurationConfiguration used: Organization UI Review profile: CHILL Plan: Pro Plus Run ID: ⛔ Files ignored due to path filters (53)
📒 Files selected for processing (2640)
You can disable this status message by setting the |
Collaborator
Author
|
/test topsail remoting_mac |
|
🔴 Test of 'mac_ai test test_ci' failed after 00 hours 21 minutes 57 seconds. 🔴 • Link to the test results. • Link to the reports index. Test configuration: |
|
@kpouget: The following test failed, say
Full PR test history. Your PR dashboard. DetailsInstructions for interacting with me using PR comments are available here. If you have questions or suggestions related to my behavior, please file an issue against the kubernetes-sigs/prow repository. I understand the commands that are listed here. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Make sure to read the contributing guidelines before submitting a PR