Carry ggml-org#25731 (TML Inkling) merged onto b10630 - #108
Carry ggml-org#25731 (TML Inkling) merged onto b10630#108danielhanchen wants to merge 20 commits into
Conversation
Hybrid attention model: 55 sliding-window plus 11 global layers, banded content-dependent relative position bias instead of RoPE, per-layer short convolution state, fine-grained MoE (256 experts top-6 plus 2 shared), attention log-scaling past 128K, 1M context. Includes the GGML_OP_FLASH_ATTN_EXT_BANDED operator (CPU and CUDA, fused into the MMA flash attention kernel with an fp16 accumulator overflow guard), HF to GGUF conversion, chat template with typed content block parsing (interleaved thinking, narration and tool calls), mmproj vision and audio support, and backend op tests at production shapes.
# Conflicts: # ggml/src/ggml-cuda/mmq.cuh # src/llama-model-saver.cpp # src/llama-vocab.h
Upstream 910196f renamed common_chat_params::thinking_end_tag to thinking_end_tags and changed it from std::string to a vector, so data.thinking_end_tag = END_MESSAGE no longer compiles.
# Conflicts: # src/llama-arch.cpp
# Conflicts: # ggml/include/ggml-rpc.h # ggml/src/ggml-backend-meta.cpp # ggml/src/ggml-cuda/ggml-cuda.cu # src/llama-model-saver.cpp # tests/test-llama-archs.cpp
op_params holds 16 int32_t, so op_params + 16 is one past the end and the 8 byte write lands on ggml_tensor::flags. Nothing reads it back: the CUDA kernel and the mma path both take rel_extent from rel_logits->ne[0].
Upstream 56db501 (ggml-org#27594) made every resize algo Pillow-accurate and folded RESIZE_ALGO_BICUBIC_PILLOW into RESIZE_ALGO_BICUBIC, converting each call site. The inkling vision branch is the same conversion.
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 60007bad63
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| GGML_OP_FILL, | ||
|
|
||
| GGML_OP_FLASH_ATTN_EXT, | ||
| GGML_OP_FLASH_ATTN_EXT_BANDED, |
There was a problem hiding this comment.
Placing the new operation between FLASH_ATTN_EXT and FLASH_ATTN_BACK renumbers every subsequent value serialized in rpc_tensor.op. When either side is still version 5.1.0, negotiate_hello() accepts the connection because it compares only major and minor versions, so remote graphs can deserialize these values as different operations and fail or compute incorrect results. Append the new operation after existing values, or make negotiation reject incompatible patch versions.
Useful? React with 👍 / 👎.
| add_executable(test-flash-attn-bias test-flash-attn-bias.cpp) | ||
| target_link_libraries(test-flash-attn-bias PRIVATE ggml) | ||
| add_test(NAME test-flash-attn-bias COMMAND test-flash-attn-bias) |
There was a problem hiding this comment.
Unless maintainers explicitly approved these additions, remove the new test-flash-attn-bias.cpp and test-flash-attn-generic-hash.cpp targets and fold essential coverage into existing test infrastructure such as test-backend-ops.cpp; the repository specifically prohibits adding files under tests/* without maintainer approval.
AGENTS.md reference: AGENTS.md:L84-L87
Useful? React with 👍 / 👎.
|
Closing. The pin in Note that |
# Conflicts: # ggml/include/ggml-rpc.h
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 2c5b0007bc
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| mtmd_inkling_image_preproc_out mtmd_image_preprocess_inkling( | ||
| const clip_image_u8 & img, | ||
| resize_algo algo) { |
There was a problem hiding this comment.
Remove this internal API and its companion output struct: a repo-wide search finds no caller of mtmd_image_preprocess_inkling, while the runtime uses mtmd_image_preprocessor_inkling::preprocess below. The two implementations already differ in resize precision, normalization, and BF16 rounding, so retaining both adds about 75 lines of dead preprocessing logic that can drift further and mislead future callers about which behavior is authoritative.
AGENTS.md reference: AGENTS.md:L80-L80
Useful? React with 👍 / 👎.
|
Superseded by #138, which is the same upstream PR merged onto b10665 with the llama-model.cpp filter_idx conflict resolved. This carry stopped merging once upstream's qwen4exp landed on its anchor line. The branch stays in place. |
2 participants
ggml-org/llama.cpp#25731stopped merging onto the aged base tag. This branch isc44c9a11a7withb10630merged into it and the five conflicts resolved by hand, so the pin set merges again. The base branch isb10630itself, so the diff here is only the Inkling work, matching the arrangement used by #70, #91 and #99.The conflicts
ggml/include/ggml-rpc.h- two independent bumps of different fields. Upstreamaf5172627dtook minor 0 to 1 (use_countreplacingpadding[4], a backward-compatible wire change); the PR took patch 0 to 1, because insertingGGML_OP_FLASH_ATTN_EXT_BANDEDmid-enum ggml_oprenumbersrpc_tensor.op, which is serialised. Resolved to5.1.1andstatic_assert(GGML_OP_COUNT == 102). History confirms patch is only reset on a major bump, never on a minor one, so incrementing both independently is the existing convention. Verified the merged enum really does have 102 ops.ggml/src/ggml-backend-meta.cpp- both sides append a new lambda at the same anchor. Kept upstream's rewrittenhandle_flash_attn_extverbatim (it carries a realsrc[4]tosrc[3]fix plus new mirrored andkv_mirroredsplit paths that the PR's older copy does not know about), then the PR'shandle_flash_attn_ext_banded, then upstream'shandle_lightning_indexer. Three disjoint handlers, all three switch cases kept.ggml/src/ggml-cuda/ggml-cuda.cu- the only genuinely overlapping hunk. Upstreamd9b6be07d0hoistscublasHandle_t cublas_h = ctx.cublas_handle();and rewrites the call sites; the PR wraps the samecublasSgemmin anif (f32_pedantic)branch. Resolved to the PR's branch structure using upstream'scublas_hin both arms. The threecublasGemm*Exsites git auto-merged were then re-checked by hand: each must carry upstream'scublas_hand the PR'scu_gemm_algo, and all five call sites do.src/llama-model-saver.cpp- purely additive fall-through labels. Kept upstream'sGRANITE_SWAandDOTS3NOTEalongside the PR'sINKLING.tests/test-llama-archs.cpp- the one that needed care. Upstreambf0a29cc16deleted theLLM_ARCH_DEEPSEEK4skip because DSv4 is now testable; the PR anchored itsINKLINGskip to that deleted block. Keeping both sides would have resurrected the skip and silently un-tested DeepSeek 4 while still compiling and still passing CI. Kept only the INKLING skip, relocated to the surviving anchor.Confirmed on the built binary that this is right:
DeepSeek 4 runs, matching the bare base tag; Inkling skips, matching the PR's intent.
Two extra commits
tools/mtmd/clip.cppdid not compile against the base. Upstream56db501e7("mtmd: use pillow-accurate algo", ggml-org#27594) made every resize algo Pillow-accurate, soRESIZE_ALGO_BICUBIC_PILLOWbecame redundant and was deleted, with each call site converted toRESIZE_ALGO_BICUBIC. The PR's Inkling vision branch still used the old name, and since the PR never touched the enum there was no conflict to flag it:Applied the same conversion upstream applied everywhere else.
Out-of-bounds write in
ggml_flash_attn_ext_banded.ggml/src/ggml.chad:op_paramsholds 16int32_t, so index 16 is one past the end. Measured offsets confirm where the 8 bytes land:It writes straight onto
ggml_tensor::flags. With therel_extent = 8the PR's own tests use, that setsGGML_TENSOR_FLAG_LOSSon every banded attention node. The value is never read back:fattn-banded.cuandfattn-mma-f16.cuhboth derive the extent fromrel_logits->ne[0], and thatmemcpyis the only reference in the tree. Deleted it.Verification
Built in the nightly's configuration (
GGML_BACKEND_DL=ON,GGML_CPU_ALL_VARIANTS=ON,GGML_RPC=ON,GGML_CUDA=ON, tools + server), plusLLAMA_BUILD_TESTS=ONwhich the nightly does not use. Everything was run against the bareb10630tag as well, and compared.test-backend-ops test -b CUDA0 -o MUL_MAT,MUL_MAT_ID,FLASH_ATTN_EXT,FLASH_ATTN_EXT_BANDED: base 4998/4998, this branch 5011/5011. Exactly the base set plus the 13 new banded cases, nothing lost.FLASH_ATTN_EXTalone: 2936/2936 on both, so ordinary flash attention is untouched by the resolution.test-flash-attn-bias(this PR's own oracle): every case PASS, covering f32/f16/bf16, GQA, sliding, decode offset, strided rel, batch broadcast, 64 heads, and an offset past INT32_MAX.test-flash-attn-generic-hash: pass.test-llama-archsplus the threetest-recurrent-state-rollbackvariants: pass, same as base.unsloth/Inkling-Small-GGUFUD-IQ1_S on a B200, coherent output at 55 tok/s.Known, not fixed here
negotiate_hellocompares major and minor only), so a merged 5.1.1 client will silently talk to a stock 5.1.0rpc-serverwith every op afterFLASH_ATTN_EXTrenumbered. That is upstream's own design gap, but it means merged RPC binaries should not be mixed with stock ones.INKLINGis absent frommoe_mandatory()intest-llama-archs.cppand fromllm_arch_supports_rs_rollback. Both are benign today (the first is gated behind the test skip, the second is an opt-in whitelist defaulting to false), but both want revisiting if anyone writes the Inkling fixture params.Once
ggml-org#25731lands upstream or a base tag carries it, the pin should move back to the upstream commit and this branch can be deleted.