From 64086f2b2f222f7032611b021b0e1b6c4767b0f1 Mon Sep 17 00:00:00 2001 From: Gabe Goodhart Date: Fri, 5 Jun 2026 09:44:59 -0600 Subject: [PATCH 1/8] model, mtmd: Granite4 Vision (#23545) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * feat(convert): Get language model conversion working for 4.1 vision Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat(convert): Skip multimodal tensors for GraniteMoeHybrid (vision 4.0) Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Disable vocab padding for non-hybrid models that use GraniteMoeHybrid Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Plumb python-side vision projector names and mappings There are several awkward things here: 1. Most of these are essentially identical to the audio qformer tensors. On the c++ side, that's mapped using the prefix, so the rest of the GGUF name needs to align, but on the python side there's no prefix notion, so they all get duplicated. 2. There are a couple of net-new tensors for vision, in particular PROJ_NORM. In both speech and vision, the QF_PROJ_NORM is qualified as belonging to the qformer portion, but the GGUF name is simply proj_norm which conflicts with the ideal name for this new PROJ_NORM that is not qualified as part of the qformer. To get around this, I used "proj_layernorm" as the GGUF name. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add python side architecture name Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add python-side plumbing for setting FEATURE_LAYERS hparam Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add c++ side tensor naming defines NOTE: Usage of these hasn't been updated to include prefix yet Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat(mtmd): Convert vision_feature_layer to an ordered vector We need to preserve the ordering of these feature index values so that they can be mapped to the sub-tensors within the stacked projectors. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat(mtmd): Add architecture label plumbing Branch: Granite4Vision AI-usage: full (OpenCode + qwen3.5:122b) Signed-off-by: Gabe Goodhart * feat(wip): Add partial conversion for mmproj This handles stacking the projector tensors and setting the new harams Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add gguf_writer and constant support for new hparams and deepstack layer arr Branch: Granite4Vision AI-usage: draft (OpenCode + qwen3.5:122b) Signed-off-by: Gabe Goodhart * feat: Full conversion for mmproj w/ tensor mappings Branch: Granite4Vision AI-usage: full (OpenCode + qwen3.5:122b) Signed-off-by: Gabe Goodhart * fix: Add lm_head skip for mmproj for 4.0 Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: De-alias text_config architecture in convert_lora_to_gguf.py Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add --trust-remote-code arg to convert_lora_to_gguf.py This defaults to False, but allows a user to enable it programmaticly instead of using the interactive prompt. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: De-alias model.language_model. -> model. for lora adapters Branch: Granite4Vision AI-usage: full (OpenCode + qwen3.5:122b) Signed-off-by: Gabe Goodhart * fix: Extend language model tensor dealiasing in adapters Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Remove unnecessary registration for GraniteSpeech in language model Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Plumb through mm prefix formatting for qformer tensors Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Refactor vision projector tensors to use predictor ID as the block This is cleaner than stacking them. The modeling file hard-codes single-layer qformers, so we can punt on the multiipule multi-layer projectors problem. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add spatial offests array hparam conversion Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add stub plumbing for granite vision in mtmd Branch: Granite4Vision AI-usage: draft (OpenCode + qwen3.5:122b) Signed-off-by: Gabe Goodhart * feat: Add new hparam and tensor naming in clip-impl.h New hparams: - KEY_PROJ_SAMPLE_QUERY_SIDE - KEY_PROJ_SAMPLE_WINDOW_SIDE - KEY_PROJ_SPATIAL_OFFSETS New tensors: - TN_MULTI_PROJ_IMG_POS - TN_MULTI_PROJ_QUERY - TN_MULTI_PROJ_LAYERNORM - TN_MULTI_PROJ_LINEAR - TN_MULTI_PROJ_NORM Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Move deepstack_layer_arr to llm hparam instead of mmproj Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Remove IS_DEEPSTACK_LAYERS This appears to have been added during Qwen3 VL (https://github.com/ggml-org/llama.cpp/pull/16780), but it was never actually used. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: n_deepstack_layers -> deepstack_layer_arr The old logic hard coded a correspondence between the first N layers of the LLM and the 1->N entries in the input embeddings. Now, that relationship is maintained at loading time if the GGUF value is single-valued. If it is multi-valued, it loads directly allowing for deepstack layers to be spaced out throughout the model. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Use try/catch for single/multi valued deepstack info The alternative would be to use get_key_or_arr, but then the single value would be populated through the entire array and we'd need to detect that and update it with the right correspondence. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add deepstack injection point for granite LLM The use of ggml_add here assumes that the elements of inp_embd will be pre- arranged to be the full embedding length with only the vision-mask'ed portions non-zero from the projector. This matches how Qwen3VL does it. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: add missing vision attn layernorm eps Branch: Granite4Vision AI-usage: full (OpenCode + Qwen 3.6-35B) Signed-off-by: Gabe Goodhart * refactor: Hoist qformer tensors into qf_block and hold a vector for multi-proj Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Fix missing prefix template for TN_QF_PROJ_LINEAR It's not strictly necessary since vision uses the blockwise version, but it makes the loading consistent. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Add embedding scale and image grid pinpoints hparams in conversion Also remove dead parsing for self._deepstack_layer_arr Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add mtmd KEY_ section for hparams shared with the LLM In this case, we need the EMBEDDING_SCALE so we can unscale the image embeddings to compensate for applying embedding scale to the input embeddings Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Implement c++ hparam parsing Branch: Granite4Vision AI-usage: draft (Claude Code) Co-authored-by: Eli Schwartz Signed-off-by: Gabe Goodhart * fix: Flatten pinpoints in conversion Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Add missing break Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: No reason to have modality prefix for img_pos Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add tensor loading Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix(convert): Fix confusion between proj.norm and proj.qformer.layernorm Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Use the right portion of speech for tensor loading! Also plumb through the layernorm -> post_norm naming change Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add logging of deepstack_layers_arr if set I also changed the print_f output type to int32_t to avoid printing overflow values for -1. This could cause overflows on the other side, but I can't imagine a value for any of the current array hparams that would trigger that. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Make sure input embeddings are cont before f_embedding_scale Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add init and mmproj_embd cases for g4v The n_mmproj_embd is 1+ to make space for the text embedding and all 8 projectors Branch: Granite4Vision AI-usage: draft (Bob) Signed-off-by: Gabe Goodhart * fix: Invert (h, w) -> (w, h) pinpoints Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Reorder projectors based on llm index and skip the first injection The multi-projector stack has a strange asymmetry based on how it's currently implemented for qwen3vl: on the mmproj side, it's all N projectors, but the output of the "first" (by inp_embd index) projector is automatically consumed as if it were a standard single-projector mmproj, so the deepstack portion needs to only contain the 1-N entries. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart Co-authored-by: Eli Schwartz * fix: Fix mmproj hparams in conversion Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart Co-authored-by: Eli Schwartz * fix: Fix ordering/logic for deepstack injection in granite Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart Co-authored-by: Eli Schwartz * fix: Fix preprocessing config to match what the model needs Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart Co-authored-by: Eli Schwartz * wip: Partial port of Eli's implementation This is still pretty broken, but it's getting closer. It now happily generates tokens, but the values are quite incorrect still. I suspect it's caused by the mapping of projectors from safetensors to their respective orders here. Also, this implementation breaks encapsulation pretty badly in mtmd_encode. This will need a big refactor to put the G4V-specific encoding logic somewhere more appropriate. Branch: Granite4Vision AI-usage: draft (Claude Code, Bob) Signed-off-by: Gabe Goodhart Co-authored-by: Eli Schwartz * fix: Fix the pre-scaling on the input embeddings to correctly invert the scale We've got tokens! They still don't line up quite right, so something's a little off, but we're getting much closer now. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: invert embedding multiplier -> base_scale at load Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Fix setting image_resize_pad after new enum introduced Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Add G4V to mmproj mapping in conversion Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Re-add padding disable for non-hybrid hybrid models Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Simplify G4V n_tokens computation This is slightly more efficient and flexible for when we implement the unpad cropping. IMO, it's also clearer that it is adding the number of image_newline tokens (embeddings) to the grid, rather than recomputing the entire count. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Add new clip APIs for post-tile-encoding assembly Granite 4 Vision uses llava-next style pack-and-unpad which requires injecting the learned newline after each row of the tile grid. A row here is a single row of the grid which is composed of (grid_x * cols_per_tile) * (grid_y * rows_per_tile), so the result is newlines injected in between individual tile rows, thus not something that can be handled with the standard llava-uhd block-wise endcoding. Branch: Granite4Vision AI-usage: draft (Claude Code + Opus 4.7) Signed-off-by: Gabe Goodhart * feat: Add model interfaces for granite 4 vision assembler I'm on the fence about the best organization of this. These free functions allow the per-architecture logic in clip.cpp to access the model-specific graph building, but they still require a fair bit of model-specific logic in clip.cpp which is not ideal. I think a better approach may be to replicate what is done with the graph builders themselves (and possibly even make the assembler part of the model's existing graph builder). Branch: Granite4Vision AI-usage: full (Claude Code + Opus 4.7) Signed-off-by: Gabe Goodhart * refactor: Remove all g4v-specific branching from mtmd.cpp in favor of clip assembler Branch: Granite4Vision AI-usage: full (Claude Code + Opus 4.7) Signed-off-by: Gabe Goodhart * refactor(mtmd): Consolidate assembler logic into clip_assembler class family Just like `clip_graph` is the base class for building the model-specific encoder graphs, `clip_assembler` will be the base class for building the model-specific assembler graphs. This allows the assembly pattern to follow how the encoder pattern is implemented where the model-specific logic lives in a subclass co-located with the encoder graph builder that gets constructed by a simple factory method. Branch: Granite4Vision AI-usage: full (Claude Code + Opus 4.7) Signed-off-by: Gabe Goodhart * style: Comment improvement Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: granite_vision -> granite4_vision Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Remove dead codepath for Qwen3VL add_vision_is_deepstack These pieces were never used on the c++ side (removed there in an earlier commit), so this is just cleanup that I missed before. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Oops! I did not mean to commit one of my prompt files But now it's too far back in history to effectively rebase out, even with interactive and --rebase-merges :( Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Add missing include for std::find It seems that this was already pulled in on some platforms, but not on others Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Fix Flake8 warnings in granite conversion module Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Remove clip_assembler in favor of clip_image_f32.append_token Per conversation in the PR, the clip_assembler pattern was too invasive. This is a compromise that limits model-specific blocks to add_media where each preprocessed tile is annotated with an injection type, after which all the token counting logic is generic and the newline injection itself is handled in the graph based on the value for the given tile image. Branch: Granite4Vision AI-usage: draft (Bob, OpenCode + Qwen 3.6 35b) Signed-off-by: Gabe Goodhart * refactor(convert): Split n_deepstack_layers and deepstack_layers (array) Branch: Granite4Vision AI-usage: full (Bob, OpenCode + Qwen3.6-35b) Signed-off-by: Gabe Goodhart * refactor(src): Handle n_deepstack_layers and deepstack_layers GGUF keys Branch: Granite4Vision AI-usage: draft (Bob, OpenCode + Qwen3.6-35b) Signed-off-by: Gabe Goodhart * fix: Fix GGUF key for deepstack_layers_arr Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Remove pre-scaling embeddings and skip scaling for raw embd inputs This follows how gemma3 and gemma4 handle embedding scaling by skipping the multiplier for raw input embeddings. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: deepstack_layers(_arr) -> deepstack_mapping(_arr) Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Fully revert changes to n_deepstack_layers and qwen3vl* Since we're going to keep the GGUF KVs separate, it makes sense to just keep the hparams separate too to limit the scope of this branch. The down side is that n_deepstack_layers and deepstack_mapping_arr are potentially conflicting. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Revert removal of "is_deepstack_layers" GGUF KV This KV is not used at all on the c++ side, so it's fully dead, but there's also no need to conflate this cleanup with the addition of G4V. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Remove unnecessary ggml_cont and build_forward_expand in cbx Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * style: Clean up comments Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Tighter and more flexible code for g4v_build_block This could be refactored to look a lot more like granite-speech, but the overall block constructs before/after the qformer are pretty different, so for now I'm going to leave it as is and just tighten a bit. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Remove unnecessary `unordered_set` include Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Add architecture guard on deepstack_mapping_arr printout Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Remove unnecessary AI-gen comment Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Always initialize deepstack_mapping_arr with -1 values This was causing `test-llama-archs` to fail, likely due to trying to save the uninitialized values, then re-loading them. It's safer to always initialize so that other models don't forget and end up with undefined behavior. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * style: Remove TODO about block/vs non-block tensor mapping Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Move is_vision_feature_layer logic into clip_hparams Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Use a bool for append_token Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * style: Remove unnecessary comment Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Remove unused get_model api yikes! Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: Rearrange helpers for g4v to be private members and use build_attn Branch: Granite4Vision AI-usage: full (Bob, OpenCode + Qwen3.6-35b) Signed-off-by: Gabe Goodhart * fix: Fix off-by-one in vision layer index This was inherited from the Claude Code implementation that pushed the negative index inversion down into the model file. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Fix norm/post_norm mixup in conversion face. palm. :( Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * style: More descriptive tensor names Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * fix: Apply PR cleanup for new conversion changes AI-usage: none Signed-off-by: Gabe Goodhart Co-authored-by: Sigbjørn Skjæret * fix(convert): Remove duplicate V_ENC_EMBD_IMGNL Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * refactor: append_token -> add_newline Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * style: Comment cleanup Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart * feat: Cleaner error handling/checking NOTE: format_string is not available in granite.cpp (and including clip-impl.h to get it doesn't compile, so I think it violates the intended encapsulation), so std::stringstream is the simplest answer. Branch: Granite4Vision AI-usage: none Signed-off-by: Gabe Goodhart --------- Signed-off-by: Gabe Goodhart --- conversion/__init__.py | 1 + conversion/granite.py | 158 +++++++++++- convert_lora_to_gguf.py | 16 +- gguf-py/gguf/constants.py | 82 ++++++- gguf-py/gguf/gguf_writer.py | 20 ++ gguf-py/gguf/tensor_mapping.py | 90 ++++++- src/llama-arch.cpp | 1 + src/llama-arch.h | 1 + src/llama-graph.cpp | 7 +- src/llama-hparams.h | 10 + src/llama-model-loader.cpp | 1 + src/llama-model-saver.cpp | 1 + src/llama-model.cpp | 15 +- src/models/granite.cpp | 37 +++ tools/mtmd/CMakeLists.txt | 1 + tools/mtmd/clip-impl.h | 73 +++--- tools/mtmd/clip-model.h | 35 ++- tools/mtmd/clip.cpp | 254 +++++++++++++++---- tools/mtmd/clip.h | 2 - tools/mtmd/models/granite-speech.cpp | 10 +- tools/mtmd/models/granite4-vision.cpp | 339 ++++++++++++++++++++++++++ tools/mtmd/models/llava.cpp | 5 +- tools/mtmd/models/models.h | 23 ++ tools/mtmd/mtmd.cpp | 29 ++- 24 files changed, 1100 insertions(+), 111 deletions(-) create mode 100644 tools/mtmd/models/granite4-vision.cpp diff --git a/conversion/__init__.py b/conversion/__init__.py index 2c79580f8..c670798fc 100644 --- a/conversion/__init__.py +++ b/conversion/__init__.py @@ -253,6 +253,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = { "Glm4vMoeForConditionalGeneration": "qwen3vl", "GlmOcrForConditionalGeneration": "qwen3vl", "GlmasrModel": "ultravox", + "Granite4VisionForConditionalGeneration": "granite", "GraniteSpeechForConditionalGeneration": "granite", "HunYuanVLForConditionalGeneration": "hunyuan", "Idefics3ForConditionalGeneration": "smolvlm", diff --git a/conversion/granite.py b/conversion/granite.py index 647269ba7..53441fe57 100644 --- a/conversion/granite.py +++ b/conversion/granite.py @@ -1,5 +1,6 @@ from __future__ import annotations +import re from typing import Any, Callable, Iterable, TYPE_CHECKING import torch @@ -13,7 +14,7 @@ from .llama import LlamaModel from .mamba import Mamba2Model -@ModelBase.register("GraniteForCausalLM", "GraniteSpeechForConditionalGeneration") +@ModelBase.register("GraniteForCausalLM") class GraniteModel(LlamaModel): """Conversion for IBM's GraniteForCausalLM""" model_arch = gguf.MODEL_ARCH.GRANITE @@ -46,11 +47,29 @@ class GraniteModel(LlamaModel): self.gguf_writer.add_logit_scale(logits_scale) logger.info("gguf: (granite) logits_scale = %s", logits_scale) + # If being used as the base for Granite4 Vision, add deepstack_layer_arr + if self.hparams.get("spatial_target_layers") or self.hparams.get("deepstack_layer_map"): + normalized_projector_map = Granite4VisionMmprojModel.get_normalized_projector_map(self.hparams) + deepstack_mapping_arr = [-1 for _ in range(self.block_count)] # Populate with -1 sentinels + for proj_idx, (_, llm_layer, _, _) in enumerate(normalized_projector_map): + # Skip the first projector which is handled as the base embedding + # stream like normal + if proj_idx == 0: + continue + deepstack_mapping_arr[llm_layer] = proj_idx + self.gguf_writer.add_deepstack_mapping(deepstack_mapping_arr) + @classmethod def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None: name, gen = item - if name.startswith("encoder."): - return None + # Skip multimodal tensors + if ( + name.startswith(("encoder.")) + or "image_" in name + or "layerwise_projectors" in name + or "spatial_projectors" in name + ): + return return super().filter_tensors(item) @@ -241,7 +260,8 @@ class GraniteHybridModel(Mamba2Model, GraniteMoeModel): assert self.d_inner % d_head == 0, f"SSM inner size {self.d_inner} not a multiple of head dim {d_head}" def set_vocab(self): - self.hparams["pad_vocab_size_multiple"] = 8 + # For models with no ssm layers, don't pad for mamba2 + self.hparams["pad_vocab_size_multiple"] = 8 if self._ssm_layers else 1 Mamba2Model.set_vocab(self) @@ -326,3 +346,133 @@ class GraniteSpeechMmprojModel(MmprojModel): data_torch = data_torch.squeeze(1) yield from super().modify_tensors(data_torch, name, bid) + + +@ModelBase.register("Granite4VisionForConditionalGeneration") +class Granite4VisionMmprojModel(MmprojModel): + has_vision_encoder = True + has_audio_encoder = False + + @staticmethod + def get_normalized_projector_map(global_config: dict) -> list[tuple[int, int, str, int]]: + """Normalize both deepstack and spatial projector maps to the form: + (vision_layer, llm_layer, , type_index) + + This is then used to populate the following mappings: + - vision_feature_layers (mmproj hparam): ordered list of all + vision_layer values where order corresponds with the order of the + stacked projector tensors + NOTE: Values may appear multiple times for spatial projectors + - tensor_prefix_map (mmproj tensors): mapping from tensor prefixes to + the index of the corresponding projector in the stacked tensors + - deepstack_layer_arr (llm hparam): per-text-layer array indicating + which input vision feature should be injected at that layer + (-1 if none) + + Output: (vision_layer, llm_layer, , type_index) + """ + deepstack_map = global_config.get("deepstack_layer_map", []) # [[vis_layer, llm_layer], ...] + spatial_layers = global_config.get("spatial_target_layers", []) # [llm_layer, ...] + n_text_layers = global_config["text_config"]["num_hidden_layers"] + n_vision_layers = global_config["vision_config"]["num_hidden_layers"] + normalized_projector_map = [] + if deepstack_map: + for deepstack_idx, (vision_layer, llm_layer) in enumerate(sorted(deepstack_map)): + if vision_layer < 0: + vision_layer = n_vision_layers + vision_layer + if llm_layer < 0: + llm_layer = n_text_layers + llm_layer + normalized_projector_map.append((vision_layer, llm_layer, "layerwise", deepstack_idx)) + if spatial_layers: + spatial_vision_layer = global_config.get("spatial_vision_layer", -1) + if spatial_vision_layer < 0: + spatial_vision_layer = n_vision_layers + spatial_vision_layer + for spatial_idx, llm_layer in enumerate(spatial_layers): + normalized_projector_map.append((spatial_vision_layer, llm_layer, "spatial", spatial_idx)) + return list(sorted(normalized_projector_map, key=(lambda entry: entry[1]))) + + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + normalized_projector_map = self.get_normalized_projector_map(self.global_config) + self._n_proj = len(normalized_projector_map) + + self._tensor_prefix_map = { + f"model.{proj_type}_projectors.{type_idx}": proj_idx + for proj_idx, (_, _, proj_type, type_idx) in enumerate(normalized_projector_map) + } + self._vision_feature_layers = [vision_layer for vision_layer, _, _, _ in normalized_projector_map] + self._spatial_offsets = [ + type_idx if proj_type == "spatial" else -1 + for _, _, proj_type, type_idx in normalized_projector_map + ] + + def set_gguf_parameters(self): + assert self.hparams_vision is not None + super().set_gguf_parameters() + + self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.GRANITE4_VISION) + + # SigLIP encoder hparams + self.gguf_writer.add_vision_attention_layernorm_eps(self.hparams.get("layer_norm_eps", 1e-6)) + self.gguf_writer.add_vision_use_gelu(True) + + # Preprocessor + self.gguf_writer.add_vision_preproc_image_size(self.hparams.get("image_size", 384)) + + # QFormer projector config + ds_rate = self.global_config["downsample_rate"] + ds_parts = ds_rate.split("/") + assert len(ds_parts) == 2, f"Invalid 'downsample_rate' value: {ds_rate}" + query_side, window_side = [int(p) for p in ds_parts] + self.gguf_writer.add_vision_projector_query_side(query_side) + self.gguf_writer.add_vision_projector_window_side(window_side) + + # Set vision feature layers + self.gguf_writer.add_vision_feature_layers(self._vision_feature_layers) + + # Set the spatial offests per projector + self.gguf_writer.add_vision_spatial_offsets(self._spatial_offsets) + + # Add flattened image grind pinpoints (resolution candidates internally) + if pinpoints := self.global_config.get("image_grid_pinpoints"): + # Flatten with h, w -> w, h inversion + pinpoints = [val for h, w in pinpoints for val in (w, h)] + self.gguf_writer.add_vision_image_grid_pinpoints(pinpoints) + + @classmethod + def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None: + name, _ = item + if ("vision_model.head" in name or name.startswith("lm_head")): + return None + return super().filter_tensors(item) + + def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: + + # Detect projector tensors and bin them + projector_idx = None + for prefix, proj_idx in self._tensor_prefix_map.items(): + if name.startswith(prefix): + projector_idx = proj_idx + break + if projector_idx is not None: + # If this projector tensor has a block id within the projector, + # alias the bid to projector_idx + # + # TODO: currently, none of the Granite 4 Vision models have + # projectors with multiple QFormer layers, so the `layer.{}` index + # is always 0. This allows us to simply map to a single `bid` that + # matches the projector index. If this changes, we'll need a + # convention that merges the two IDs. + id_matches = list(re.finditer(r"\.([0-9]+)\.", name)) + all_ids = [int(m.group(1)) for m in id_matches] + assert len(all_ids) >= 1 and len(all_ids) <= 2, "Must have at least 1 and at most 2 ids in tensor names" + # If not layer id, just use the projector index + new_bid = projector_idx + if len(all_ids) == 1: + new_name = name[:id_matches[0].span(1)[0]] + str(new_bid) + name[id_matches[0].span(1)[1]:] + else: # len(all_ids) == 2 + new_bid = projector_idx # + all_ids[1] + new_name = name[:id_matches[0].span(0)[0]] + name[id_matches[0].span(1)[1]:id_matches[1].span(1)[0]] + str(new_bid) + name[id_matches[1].span(1)[1]:] + yield from super().modify_tensors(data_torch, new_name, new_bid) + return + yield from super().modify_tensors(data_torch, name, bid) diff --git a/convert_lora_to_gguf.py b/convert_lora_to_gguf.py index 9a6437bea..45202b333 100755 --- a/convert_lora_to_gguf.py +++ b/convert_lora_to_gguf.py @@ -311,6 +311,10 @@ def parse_args() -> argparse.Namespace: "--base-model-id", type=str, help="the model ID of the base model, if it is not available locally or in the adapter config. If specified, it will ignore --base and load the base model config from the Hugging Face hub (Example: 'meta-llama/Llama-3.2-1B-Instruct')", ) + parser.add_argument( + "--trust-remote-code", default=False, action="store_true", + help="trust remote code in the model", + ) parser.add_argument( "lora_path", type=Path, help="directory containing Hugging Face PEFT LoRA config (adapter_model.json) and weights (adapter_model.safetensors or adapter_model.bin)", @@ -319,11 +323,11 @@ def parse_args() -> argparse.Namespace: return parser.parse_args() -def load_hparams_from_hf(hf_model_id: str) -> tuple[dict[str, Any], Path | None]: +def load_hparams_from_hf(hf_model_id: str, trust_remote_code: bool) -> tuple[dict[str, Any], Path | None]: from huggingface_hub import try_to_load_from_cache # normally, adapter does not come with base model config, we need to load it from AutoConfig - config = AutoConfig.from_pretrained(hf_model_id) + config = AutoConfig.from_pretrained(hf_model_id, trust_remote_code=trust_remote_code) cache_dir = try_to_load_from_cache(hf_model_id, "config.json") cache_dir = Path(cache_dir).parent if isinstance(cache_dir, str) else None @@ -372,13 +376,13 @@ if __name__ == '__main__': # load base model if base_model_id is not None: logger.info(f"Loading base model from Hugging Face: {base_model_id}") - hparams, dir_base_model = load_hparams_from_hf(base_model_id) + hparams, dir_base_model = load_hparams_from_hf(base_model_id, args.trust_remote_code) elif dir_base_model is None: if "base_model_name_or_path" in lparams: model_id = lparams["base_model_name_or_path"] logger.info(f"Loading base model from Hugging Face: {model_id}") try: - hparams, dir_base_model = load_hparams_from_hf(model_id) + hparams, dir_base_model = load_hparams_from_hf(model_id, args.trust_remote_code) except OSError as e: logger.error(f"Failed to load base model config: {e}") logger.error("Please try downloading the base model and add its path to --base") @@ -393,7 +397,9 @@ if __name__ == '__main__': with torch.inference_mode(): try: - model_class = get_model_class(hparams["architectures"][0]) + model_arch = hparams.get("text_config", {}).get("architectures", hparams["architectures"])[0] + logger.info("Using model architecture: %s", model_arch) + model_class = get_model_class(model_arch) except NotImplementedError: logger.error(f"Model {hparams['architectures'][0]} is not supported") sys.exit(1) diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index ce556ec9b..814980ce5 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -128,6 +128,7 @@ class Keys: MOE_LATENT_SIZE = "{arch}.moe_latent_size" NEXTN_PREDICT_LAYERS = "{arch}.nextn_predict_layers" NUM_DEEPSTACK_LAYERS = "{arch}.n_deepstack_layers" + DEEPSTACK_MAPPING = "{arch}.deepstack_mapping" POOLING_TYPE = "{arch}.pooling_type" LOGIT_SCALE = "{arch}.logit_scale" DECODER_START_TOKEN_ID = "{arch}.decoder_start_token_id" @@ -325,6 +326,8 @@ class Keys: WA_PATTERN_MODE = "clip.vision.wa_pattern_mode" # used by mimovl, per-layer -1/0/1 IS_DEEPSTACK_LAYERS = "clip.vision.is_deepstack_layers" WINDOW_SIZE = "clip.vision.window_size" + FEATURE_LAYERS = "clip.vision.feature_layer" # Granite4 Vision + IMAGE_GRID_PINPOINTS = "clip.vision.image_grid_pinpoints" # Granite4 Vision class Attention: HEAD_COUNT = "clip.vision.attention.head_count" @@ -333,6 +336,9 @@ class Keys: class Projector: SCALE_FACTOR = "clip.vision.projector.scale_factor" + QUERY_SIDE = "clip.vision.projector.query_side" + WINDOW_SIDE = "clip.vision.projector.window_side" + SPATIAL_OFFSETS = "clip.vision.projector.spatial_offsets" class SAM: BLOCK_COUNT = "clip.vision.sam.block_count" @@ -821,6 +827,31 @@ class MODEL_TENSOR(IntEnum): V_RESMPL_QUERY_768 = auto() # Deepseek-OCR-2 V_RESMPL_QUERY_1024 = auto() # Deepseek-OCR-2 + # qformer projector (vision) - Granite4 Vision + V_QF_PROJ_QUERY = auto() + V_QF_PROJ_NORM = auto() + V_QF_PROJ_LINEAR = auto() + V_QF_SELF_ATTN_Q = auto() + V_QF_SELF_ATTN_K = auto() + V_QF_SELF_ATTN_V = auto() + V_QF_SELF_ATTN_O = auto() + V_QF_SELF_ATTN_NORM = auto() + V_QF_CROSS_ATTN_Q = auto() + V_QF_CROSS_ATTN_K = auto() + V_QF_CROSS_ATTN_V = auto() + V_QF_CROSS_ATTN_O = auto() + V_QF_CROSS_ATTN_NORM = auto() + V_QF_FFN_UP = auto() + V_QF_FFN_DOWN = auto() + V_QF_FFN_NORM = auto() + V_PROJ_NORM = auto() + # multi-projector (bid => projector id) - Granite4 vision + V_MULTI_PROJ_IMG_POS = auto() + V_MULTI_PROJ_QUERY = auto() + V_MULTI_PROJ_NORM = auto() + V_MULTI_PROJ_LINEAR = auto() + V_MULTI_PROJ_POST_NORM = auto() + # audio (mtmd) A_ENC_EMBD_POS = auto() A_ENC_EMBD_NORM = auto() @@ -885,7 +916,7 @@ class MODEL_TENSOR(IntEnum): A_CTC_OUT = auto() A_CTC_OUT_MID = auto() A_ENC_ATTN_REL_POS_EMB = auto() - # qformer projector + # audio qformer projector A_QF_PROJ_QUERY = auto() A_QF_PROJ_NORM = auto() A_QF_PROJ_LINEAR = auto() @@ -1337,10 +1368,33 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = { MODEL_TENSOR.V_SAM_NECK: "v.sam.neck.{bid}", MODEL_TENSOR.V_SAM_NET_2: "v.sam.net_2", MODEL_TENSOR.V_SAM_NET_3: "v.sam.net_3", - MODEL_TENSOR.V_ENC_EMBD_IMGNL: "v.image_newline", # Deepseek-OCR + MODEL_TENSOR.V_ENC_EMBD_IMGNL: "v.image_newline", # Deepseek-OCR, Granite4Vision MODEL_TENSOR.V_ENC_EMBD_VSEP: "v.view_seperator", # Deepseek-OCR MODEL_TENSOR.V_RESMPL_QUERY_768: "v.resample_query_768", # Deepseek-OCR-2 qwen2 MODEL_TENSOR.V_RESMPL_QUERY_1024: "v.resample_query_1024", # Deepseek-OCR-2 qwen2 + # Granite4 Vision + # qformer layers (bid => proj_id) + # NOTE: Names align with A_QF_* + MODEL_TENSOR.V_QF_SELF_ATTN_Q: "v.proj_blk.{bid}.self_attn_q", + MODEL_TENSOR.V_QF_SELF_ATTN_K: "v.proj_blk.{bid}.self_attn_k", + MODEL_TENSOR.V_QF_SELF_ATTN_V: "v.proj_blk.{bid}.self_attn_v", + MODEL_TENSOR.V_QF_SELF_ATTN_O: "v.proj_blk.{bid}.self_attn_out", + MODEL_TENSOR.V_QF_SELF_ATTN_NORM: "v.proj_blk.{bid}.self_attn_norm", + MODEL_TENSOR.V_QF_CROSS_ATTN_Q: "v.proj_blk.{bid}.cross_attn_q", + MODEL_TENSOR.V_QF_CROSS_ATTN_K: "v.proj_blk.{bid}.cross_attn_k", + MODEL_TENSOR.V_QF_CROSS_ATTN_V: "v.proj_blk.{bid}.cross_attn_v", + MODEL_TENSOR.V_QF_CROSS_ATTN_O: "v.proj_blk.{bid}.cross_attn_out", + MODEL_TENSOR.V_QF_CROSS_ATTN_NORM: "v.proj_blk.{bid}.cross_attn_norm", + MODEL_TENSOR.V_QF_FFN_UP: "v.proj_blk.{bid}.ffn_up", + MODEL_TENSOR.V_QF_FFN_DOWN: "v.proj_blk.{bid}.ffn_down", + MODEL_TENSOR.V_QF_FFN_NORM: "v.proj_blk.{bid}.ffn_norm", + # multi-projector (bid => projector ID) + MODEL_TENSOR.V_MULTI_PROJ_IMG_POS: "v.proj_blk.{bid}.img_pos", + MODEL_TENSOR.V_MULTI_PROJ_QUERY: "v.proj_blk.{bid}.query", + MODEL_TENSOR.V_MULTI_PROJ_NORM: "v.proj_blk.{bid}.norm", + MODEL_TENSOR.V_MULTI_PROJ_LINEAR: "v.proj_blk.{bid}.linear", + MODEL_TENSOR.V_MULTI_PROJ_POST_NORM: "v.proj_blk.{bid}.post_norm", + # audio (mtmd) # note: all audio tensor names must use prefix "a." or "mm.a." MODEL_TENSOR.A_ENC_EMBD_POS: "a.position_embd", @@ -1522,6 +1576,29 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = { MODEL_TENSOR.V_SAM_NET_3, MODEL_TENSOR.V_RESMPL_QUERY_768, MODEL_TENSOR.V_RESMPL_QUERY_1024, + MODEL_TENSOR.V_PROJ_NORM, + MODEL_TENSOR.V_QF_PROJ_QUERY, + MODEL_TENSOR.V_QF_PROJ_NORM, + MODEL_TENSOR.V_QF_PROJ_LINEAR, + MODEL_TENSOR.V_QF_SELF_ATTN_Q, + MODEL_TENSOR.V_QF_SELF_ATTN_K, + MODEL_TENSOR.V_QF_SELF_ATTN_V, + MODEL_TENSOR.V_QF_SELF_ATTN_O, + MODEL_TENSOR.V_QF_SELF_ATTN_NORM, + MODEL_TENSOR.V_QF_CROSS_ATTN_Q, + MODEL_TENSOR.V_QF_CROSS_ATTN_K, + MODEL_TENSOR.V_QF_CROSS_ATTN_V, + MODEL_TENSOR.V_QF_CROSS_ATTN_O, + MODEL_TENSOR.V_QF_CROSS_ATTN_NORM, + MODEL_TENSOR.V_QF_FFN_UP, + MODEL_TENSOR.V_QF_FFN_DOWN, + MODEL_TENSOR.V_QF_FFN_NORM, + MODEL_TENSOR.V_QF_PROJ_NORM, + MODEL_TENSOR.V_MULTI_PROJ_IMG_POS, + MODEL_TENSOR.V_MULTI_PROJ_QUERY, + MODEL_TENSOR.V_MULTI_PROJ_LINEAR, + MODEL_TENSOR.V_MULTI_PROJ_NORM, + MODEL_TENSOR.V_MULTI_PROJ_POST_NORM, # audio MODEL_TENSOR.A_ENC_EMBD_POS, MODEL_TENSOR.A_ENC_EMBD_NORM, @@ -4388,6 +4465,7 @@ class VisionProjectorType: MINICPMV4_6 = "minicpmv4_6" GRANITE_SPEECH = "granite_speech" # audio MIMOVL = "mimovl" + GRANITE4_VISION = "granite4_vision" # Items here are (block size, type size) diff --git a/gguf-py/gguf/gguf_writer.py b/gguf-py/gguf/gguf_writer.py index 875d0f73d..182c9c54a 100644 --- a/gguf-py/gguf/gguf_writer.py +++ b/gguf-py/gguf/gguf_writer.py @@ -959,8 +959,13 @@ class GGUFWriter: self.add_uint32(Keys.LLM.POOLING_TYPE.format(arch=self.arch), value.value) def add_num_deepstack_layers(self, count: int) -> None: + """Add scalar deepstack layer count (qwen3vl format)""" self.add_uint32(Keys.LLM.NUM_DEEPSTACK_LAYERS.format(arch=self.arch), count) + def add_deepstack_mapping(self, layers: Sequence[int]) -> None: + """Add per-layer deepstack projector indices (Granite4 Vision format)""" + self.add_array(Keys.LLM.DEEPSTACK_MAPPING.format(arch=self.arch), list(layers)) + def add_rope_dimension_count(self, count: int) -> None: self.add_uint32(Keys.Rope.DIMENSION_COUNT.format(arch=self.arch), count) @@ -1184,6 +1189,15 @@ class GGUFWriter: def add_vision_preproc_image_size(self, value: int) -> None: self.add_uint32(Keys.ClipVision.PREPROC_IMAGE_SIZE, value) + def add_vision_projector_query_side(self, value: int) -> None: + self.add_uint32(Keys.ClipVision.Projector.QUERY_SIDE, value) + + def add_vision_projector_window_side(self, value: int) -> None: + self.add_uint32(Keys.ClipVision.Projector.WINDOW_SIDE, value) + + def add_vision_spatial_offsets(self, layers: Sequence[int]) -> None: + self.add_array(Keys.ClipVision.Projector.SPATIAL_OFFSETS, layers) + def add_vision_image_mean(self, values: Sequence[float]) -> None: self.add_array(Keys.ClipVision.IMAGE_MEAN, values) @@ -1240,6 +1254,12 @@ class GGUFWriter: def add_vision_window_size(self, value: int) -> None: self.add_uint32(Keys.ClipVision.WINDOW_SIZE, value) + def add_vision_feature_layers(self, layers: Sequence[int]) -> None: + self.add_array(Keys.ClipVision.FEATURE_LAYERS, layers) + + def add_vision_image_grid_pinpoints(self, layers: Sequence[Sequence[int]]) -> None: + self.add_array(Keys.ClipVision.IMAGE_GRID_PINPOINTS, layers) + def add_vision_sam_layers_count(self, value: int) -> None: self.add_uint32(Keys.ClipVision.SAM.BLOCK_COUNT, value) diff --git a/gguf-py/gguf/tensor_mapping.py b/gguf-py/gguf/tensor_mapping.py index 82f26e7b3..3e63b2165 100644 --- a/gguf-py/gguf/tensor_mapping.py +++ b/gguf-py/gguf/tensor_mapping.py @@ -1408,6 +1408,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_EMBD_PATCH: ( + "model.vision_tower.vision_model.embeddings.patch_embedding", # Granite4Vision "vision_tower.vision_model.embeddings.patch_embedding", "model.vision_tower.embeddings.patch_embedding", # minicpmv4_6 "model.vision_tower.embeddings.patch_embeddings.projection", # Intern-S1 @@ -1439,6 +1440,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_EMBD_POS: ( + "model.vision_tower.vision_model.embeddings.position_embedding", # Granite4Vision "vision_tower.vision_model.embeddings.position_embedding", "model.vision_tower.embeddings.position_embedding", # minicpmv4_6 "model.vision_tower.embeddings.position_embeddings", # Intern-S1 @@ -1456,8 +1458,9 @@ class TensorNameMap: "model.vision_embedder.pos_embedding", # gemma4 unified ), + # TODO: I think these should all be moved to mapping_cfg? MODEL_TENSOR.V_ENC_EMBD_IMGNL: ( - "model.image_newline", # Deepseek-OCR + "model.image_newline", # Deepseek-OCR, Granite4Vision "vit.perceive.image_newline", # HunyuanVL ), @@ -1477,6 +1480,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_ATTN_Q: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.self_attn.q_proj", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.self_attn.q_proj", "model.vision_tower.encoder.layers.{bid}.self_attn.q_proj", # minicpmv4_6 "model.vision_tower.encoder.layer.{bid}.attention.q_proj", # Intern-S1 @@ -1502,6 +1506,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_ATTN_K: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.self_attn.k_proj", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.self_attn.k_proj", "model.vision_tower.encoder.layers.{bid}.self_attn.k_proj", # minicpmv4_6 "model.vision_tower.encoder.layer.{bid}.attention.k_proj", # Intern-S1 @@ -1527,6 +1532,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_ATTN_V: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.self_attn.v_proj", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.self_attn.v_proj", "model.vision_tower.encoder.layers.{bid}.self_attn.v_proj", # minicpmv4_6 "model.vision_tower.encoder.layer.{bid}.attention.v_proj", # Intern-S1 @@ -1545,6 +1551,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_INPUT_NORM: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.layer_norm1", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.layer_norm1", "model.vision_tower.encoder.layers.{bid}.layer_norm1", # minicpmv4_6 "vision_tower.vision_model.encoder.layers.{bid}.norm1", # InternVL @@ -1567,6 +1574,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_ATTN_O: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.self_attn.out_proj", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.self_attn.out_proj", "model.vision_tower.encoder.layers.{bid}.self_attn.out_proj", # minicpmv4_6 "vision_tower.vision_model.encoder.layers.{bid}.attn.proj", # InternVL @@ -1595,6 +1603,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_POST_ATTN_NORM: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.layer_norm2", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.layer_norm2", "model.vision_tower.encoder.layers.{bid}.layer_norm2", # minicpmv4_6 "vision_tower.vision_model.encoder.layers.{bid}.norm2", # InternVL @@ -1618,6 +1627,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_FFN_UP: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.mlp.fc1", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.mlp.fc1", "model.vision_tower.encoder.layers.{bid}.mlp.fc1", # minicpmv4_6 "model.vision_tower.encoder.layer.{bid}.mlp.fc1", # Intern-S1 @@ -1649,6 +1659,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_FFN_DOWN: ( + "model.vision_tower.vision_model.encoder.layers.{bid}.mlp.fc2", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.mlp.fc2", "model.vision_tower.encoder.layers.{bid}.mlp.fc2", # minicpmv4_6 "model.vision_tower.encoder.layer.{bid}.mlp.fc2", # Intern-S1 @@ -1706,6 +1717,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_POST_NORM: ( + "model.vision_tower.vision_model.post_layernorm", # Granite4Vision "vision_tower.vision_model.post_layernorm", "model.vision_tower.post_layernorm", # minicpmv4_6 "model.vision_model.post_layernorm", # SmolVLM @@ -1952,6 +1964,82 @@ class TensorNameMap: "model.vision_tower.std_scale", # gemma4 ), + # For these tensors, bid => projector ID + MODEL_TENSOR.V_MULTI_PROJ_IMG_POS: ( + "model.layerwise_projectors.{bid}.image_positions", # Granite4 Vision + "model.spatial_projectors.{bid}.image_positions", # Granite4 Vision + ), + MODEL_TENSOR.V_MULTI_PROJ_QUERY: ( + "model.layerwise_projectors.{bid}.query", # Granite4 Vision + "model.spatial_projectors.{bid}.query", # Granite4 Vision + ), + MODEL_TENSOR.V_MULTI_PROJ_LINEAR: ( + "model.layerwise_projectors.{bid}.out_linear", # Granite4 Vision + "model.spatial_projectors.{bid}.out_linear", # Granite4 Vision + ), + MODEL_TENSOR.V_MULTI_PROJ_NORM: ( + "model.layerwise_projectors.{bid}.norm", # Granite4 Vision + "model.spatial_projectors.{bid}.norm", # Granite4 Vision + ), + MODEL_TENSOR.V_MULTI_PROJ_POST_NORM: ( + "model.layerwise_projectors.{bid}.qformer.layernorm", # Granite4 Vision + "model.spatial_projectors.{bid}.qformer.layernorm", # Granite4 Vision + ), + + # For these tensors, bid => proj-id + MODEL_TENSOR.V_QF_SELF_ATTN_Q: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.attention.attention.query", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.attention.attention.query", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_SELF_ATTN_K: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.attention.attention.key", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.attention.attention.key", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_SELF_ATTN_V: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.attention.attention.value", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.attention.attention.value", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_SELF_ATTN_O: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.attention.output.dense", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.attention.output.dense", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_SELF_ATTN_NORM: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.attention.output.LayerNorm", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.attention.output.LayerNorm", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_CROSS_ATTN_Q: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.crossattention.attention.query", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.crossattention.attention.query", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_CROSS_ATTN_K: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.crossattention.attention.key", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.crossattention.attention.key", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_CROSS_ATTN_V: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.crossattention.attention.value", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.crossattention.attention.value", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_CROSS_ATTN_O: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.crossattention.output.dense", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.crossattention.output.dense", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_CROSS_ATTN_NORM: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.crossattention.output.LayerNorm", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.crossattention.output.LayerNorm", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_FFN_UP: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.intermediate_query.dense", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.intermediate_query.dense", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_FFN_DOWN: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.output_query.dense", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.output_query.dense", # Granite4 Vision + ), + MODEL_TENSOR.V_QF_FFN_NORM: ( + "model.layerwise_projectors.qformer.encoder.layer.{bid}.output_query.LayerNorm", # Granite4 Vision + "model.spatial_projectors.qformer.encoder.layer.{bid}.output_query.LayerNorm", # Granite4 Vision + ), + # audio (mtmd) MODEL_TENSOR.A_ENC_EMBD_POS: ( diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index fea898dea..52963f8f1 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -196,6 +196,7 @@ static const std::map LLM_KV_NAMES = { { LLM_KV_MOE_LATENT_SIZE, "%s.moe_latent_size" }, { LLM_KV_NEXTN_PREDICT_LAYERS, "%s.nextn_predict_layers" }, { LLM_KV_NUM_DEEPSTACK_LAYERS, "%s.n_deepstack_layers" }, + { LLM_KV_DEEPSTACK_MAPPING, "%s.deepstack_mapping" }, { LLM_KV_HIDDEN_ACT, "%s.hidden_activation" }, { LLM_KV_POOLING_TYPE, "%s.pooling_type" }, { LLM_KV_LOGIT_SCALE, "%s.logit_scale" }, diff --git a/src/llama-arch.h b/src/llama-arch.h index f364f6b0b..dc9bca9bf 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -200,6 +200,7 @@ enum llm_kv { LLM_KV_MOE_LATENT_SIZE, LLM_KV_NEXTN_PREDICT_LAYERS, LLM_KV_NUM_DEEPSTACK_LAYERS, + LLM_KV_DEEPSTACK_MAPPING, LLM_KV_HIDDEN_ACT, LLM_KV_POOLING_TYPE, LLM_KV_LOGIT_SCALE, diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index 172edf24c..3b8125cde 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -1859,7 +1859,12 @@ ggml_tensor * llm_graph_context::build_inp_embd(ggml_tensor * tok_embd) const { res->t_inp_embd = cur; // For Granite architecture - if (hparams.f_embedding_scale != 0.0f) { + // NOTE: Only apply scale to token inputs. Raw embeddings are assumed to be + // multimodal inputs that should not be scaled. + if (ubatch.token && hparams.f_embedding_scale != 0.0f) { + if (!ggml_is_contiguous(cur)) { + cur = ggml_cont(ctx0, cur); + } cur = ggml_scale(ctx0, cur, hparams.f_embedding_scale); } diff --git a/src/llama-hparams.h b/src/llama-hparams.h index fde6183e8..87db4a0dd 100644 --- a/src/llama-hparams.h +++ b/src/llama-hparams.h @@ -219,8 +219,18 @@ struct llama_hparams { uint32_t indexer_top_k = 0; // qwen3vl deepstack + // When parsed from GGUF, this implies the first N layers consume the first + // N deepstack embeddings. Use deepstack_mapping_arr if you need a more + // complex mapping. If using deepstack_mapping_arr, also make sure to set + // n_deepstack_layers to the number of unique deepstack layers so that + // n_embd_imp is accurate (see granite.cpp). uint32_t n_deepstack_layers = 0; + // deepstack layer array (Granite4 Vision) + // -1 => no deepstack + // >=0 => input embedding index for deepstack injection + std::array deepstack_mapping_arr; + // gemma4 per-layer embedding uint32_t n_embd_per_layer = 0; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index ba08a19ac..0d1cf3cc3 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -393,6 +393,7 @@ namespace GGUFMeta { } template bool llama_model_loader::get_arr>(enum llm_kv kid, std::vector & result, bool required); + template bool llama_model_loader::get_arr>(enum llm_kv kid, std::array & result, bool required); template bool llama_model_loader::get_key(const std::string & key, T & result, bool required) { diff --git a/src/llama-model-saver.cpp b/src/llama-model-saver.cpp index b05228780..67d4a9df0 100644 --- a/src/llama-model-saver.cpp +++ b/src/llama-model-saver.cpp @@ -229,6 +229,7 @@ void llama_model_saver::add_kv_from_model() { add_kv(LLM_KV_MOE_EVERY_N_LAYERS, hparams.moe_every_n_layers); add_kv(LLM_KV_NEXTN_PREDICT_LAYERS, hparams.n_layer_nextn); add_kv(LLM_KV_NUM_DEEPSTACK_LAYERS, hparams.n_deepstack_layers); + add_kv(LLM_KV_DEEPSTACK_MAPPING, hparams.deepstack_mapping_arr); add_kv(LLM_KV_POOLING_TYPE, uint32_t(hparams.pooling_type)); add_kv(LLM_KV_LOGIT_SCALE, hparams.f_logit_scale); add_kv(LLM_KV_DECODER_START_TOKEN_ID, hparams.dec_start_token_id); diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 784deb70a..6808ad044 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1100,6 +1100,9 @@ void llama_model_base::load_hparams(llama_model_loader & ml) { ml.get_key_or_arr(LLM_KV_FEED_FORWARD_LENGTH, hparams.n_ff_arr, hparams.n_layer(), false); ml.get_key_or_arr(LLM_KV_ATTENTION_HEAD_COUNT, hparams.n_head_arr, hparams.n_layer(), false); + // Populate deepstack_mapping_arr - initialized to -1 (no deepstack) + std::fill(hparams.deepstack_mapping_arr.begin(), hparams.deepstack_mapping_arr.end(), -1); + // n_head_kv is optional, default to n_head hparams.n_head_kv_arr = hparams.n_head_arr; @@ -1678,10 +1681,10 @@ uint64_t llama_model::n_elements() const { void llama_model::print_info() const { const std::string rope_scaling_type = llama_rope_scaling_type_name(hparams.rope_scaling_type_train); - auto print_f = [](const std::function & f, uint32_t n) { + auto print_f = [](const std::function & f, uint32_t n) { bool is_var = false; - std::vector v; + std::vector v; for (uint32_t i = 0; i < n; ++i) { v.push_back(f(i)); if (v[i] != v[0]) { @@ -1755,6 +1758,14 @@ void llama_model::print_info() const { LLAMA_LOG_INFO("%s: n_ctx_orig_yarn = %u\n", __func__, hparams.n_ctx_orig_yarn); LLAMA_LOG_INFO("%s: rope_yarn_log_mul = %.4f\n", __func__, hparams.rope_yarn_log_mul); LLAMA_LOG_INFO("%s: rope_finetuned = %s\n", __func__, hparams.rope_finetuned ? "yes" : "unknown"); + if (arch == LLM_ARCH_GRANITE && + std::any_of(hparams.deepstack_mapping_arr.begin(), + hparams.deepstack_mapping_arr.end(), + [](const auto & entry) { return entry >= 0; })) { + LLAMA_LOG_INFO("%s: deepstack_mapping_arr = %s\n", __func__, + print_f([&](uint32_t il) { return hparams.deepstack_mapping_arr[il]; }, + hparams.n_layer).c_str()); + } // MRoPE (Multi-axis Rotary Position Embedding) sections if (const auto & s = hparams.rope_sections; s[0] || s[1] || s[2] || s[3]) { LLAMA_LOG_INFO("%s: mrope sections = [%d, %d, %d, %d]\n", __func__, s[0], s[1], s[2], s[3]); diff --git a/src/models/granite.cpp b/src/models/granite.cpp index 7aff942da..4a75c5ff3 100644 --- a/src/models/granite.cpp +++ b/src/models/granite.cpp @@ -1,5 +1,7 @@ #include "models.h" +#include + void llama_model_granite::load_arch_hparams(llama_model_loader & ml) { ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps); ml.get_key(LLM_KV_LOGIT_SCALE, hparams.f_logit_scale); @@ -7,6 +9,27 @@ void llama_model_granite::load_arch_hparams(llama_model_loader & ml) { ml.get_key(LLM_KV_EMBEDDING_SCALE, hparams.f_embedding_scale, false); ml.get_key(LLM_KV_ATTENTION_SCALE, hparams.f_attention_scale, false); + // Granite4 Vision uses array deepstack_mapping + ml.get_arr(LLM_KV_DEEPSTACK_MAPPING, hparams.deepstack_mapping_arr, false); + + // Count the unique deepstack input indices + std::unordered_set unique_deepstack_idxs; + for (const auto val : hparams.deepstack_mapping_arr) { + if (val >= 0) { + unique_deepstack_idxs.insert(val); + } + } + hparams.n_deepstack_layers = unique_deepstack_idxs.size(); + + // Ensure all values are valid (avoid overflow attacks) + for (const auto val : unique_deepstack_idxs) { + if (val > hparams.n_deepstack_layers) { + std::stringstream ss; + ss << "Invalid deepstack index: " << val << " > " << hparams.n_deepstack_layers; + throw std::runtime_error(ss.str()); + } + } + // Granite uses rope_finetuned as a switch for rope, so default to true bool rope_finetuned = true; ml.get_key(LLM_KV_ROPE_SCALING_FINETUNED, rope_finetuned, false); @@ -112,6 +135,20 @@ llama_model_granite::graph::graph( ggml_tensor * inp_out_ids = build_inp_out_ids(); for (int il = 0; il < n_layer; ++il) { + + // Granite Vision 4.1 deepstack: inject the projector stream that + // targets decoder layer `il` before the decoder runs. + // NOTE: skip the first deepstack layer since that's inpL + const auto & deepstack_emb_idx = hparams.deepstack_mapping_arr[il]; + if (il > 0 && deepstack_emb_idx >= 0) { + ggml_tensor * ds = ggml_view_2d(ctx0, + res->t_inp_embd, n_embd, n_tokens, + res->t_inp_embd->nb[1], + deepstack_emb_idx * n_embd * sizeof(float)); + inpL = ggml_add(ctx0, inpL, ds); + cb(inpL, "deepstack_in", il); + } + ggml_tensor * inpSA = inpL; // norm diff --git a/tools/mtmd/CMakeLists.txt b/tools/mtmd/CMakeLists.txt index 93f005652..20c531786 100644 --- a/tools/mtmd/CMakeLists.txt +++ b/tools/mtmd/CMakeLists.txt @@ -25,6 +25,7 @@ add_library(mtmd models/gemma4uv.cpp models/glm4v.cpp models/granite-speech.cpp + models/granite4-vision.cpp models/hunyuanvl.cpp models/internvl.cpp models/kimivl.cpp diff --git a/tools/mtmd/clip-impl.h b/tools/mtmd/clip-impl.h index c055cfb75..393e085f7 100644 --- a/tools/mtmd/clip-impl.h +++ b/tools/mtmd/clip-impl.h @@ -35,20 +35,22 @@ #define KEY_LAYER_NORM_EPS "clip.%s.attention.layer_norm_epsilon" // vision-specific -#define KEY_VISION_PROJ_TYPE "clip.vision.projector_type" // for models with mixed modalities -#define KEY_IMAGE_SIZE "clip.vision.image_size" -#define KEY_IMAGE_MIN_PIXELS "clip.vision.image_min_pixels" -#define KEY_IMAGE_MAX_PIXELS "clip.vision.image_max_pixels" -#define KEY_PREPROC_MIN_TILES "clip.vision.preproc_min_tiles" -#define KEY_PREPROC_MAX_TILES "clip.vision.preproc_max_tiles" -#define KEY_PREPROC_IMAGE_SIZE "clip.vision.preproc_image_size" -#define KEY_PATCH_SIZE "clip.vision.patch_size" -#define KEY_IMAGE_MEAN "clip.vision.image_mean" -#define KEY_IMAGE_STD "clip.vision.image_std" -#define KEY_FEATURE_LAYER "clip.vision.feature_layer" -#define KEY_PROJ_SCALE_FACTOR "clip.vision.projector.scale_factor" -#define KEY_SPATIAL_MERGE_SIZE "clip.vision.spatial_merge_size" -#define KEY_IS_DEEPSTACK_LAYERS "clip.vision.is_deepstack_layers" +#define KEY_VISION_PROJ_TYPE "clip.vision.projector_type" // for models with mixed modalities +#define KEY_IMAGE_SIZE "clip.vision.image_size" +#define KEY_IMAGE_MIN_PIXELS "clip.vision.image_min_pixels" +#define KEY_IMAGE_MAX_PIXELS "clip.vision.image_max_pixels" +#define KEY_PREPROC_MIN_TILES "clip.vision.preproc_min_tiles" +#define KEY_PREPROC_MAX_TILES "clip.vision.preproc_max_tiles" +#define KEY_PREPROC_IMAGE_SIZE "clip.vision.preproc_image_size" +#define KEY_PATCH_SIZE "clip.vision.patch_size" +#define KEY_IMAGE_MEAN "clip.vision.image_mean" +#define KEY_IMAGE_STD "clip.vision.image_std" +#define KEY_FEATURE_LAYER "clip.vision.feature_layer" +#define KEY_PROJ_SCALE_FACTOR "clip.vision.projector.scale_factor" +#define KEY_PROJ_SAMPLE_QUERY_SIDE "clip.vision.projector.query_side" +#define KEY_PROJ_SAMPLE_WINDOW_SIDE "clip.vision.projector.window_side" +#define KEY_PROJ_SPATIAL_OFFSETS "clip.vision.projector.spatial_offsets" +#define KEY_SPATIAL_MERGE_SIZE "clip.vision.spatial_merge_size" #define KEY_MM_PATCH_MERGE_TYPE "clip.vision.mm_patch_merge_type" #define KEY_IMAGE_GRID_PINPOINTS "clip.vision.image_grid_pinpoints" @@ -72,7 +74,6 @@ #define KEY_A_PROJ_DOWNSAMPLE_RATE "clip.audio.projector.downsample_rate" #define KEY_A_PROJ_HEAD_COUNT "clip.audio.projector.head_count" - // // tensor name constants // @@ -210,22 +211,28 @@ #define TN_CTC_OUT_MID "a.enc_ctc_out_mid.%s" #define TN_ATTN_REL_POS_EMB "%s.blk.%d.attn_rel_pos_emb" // qformer projector -#define TN_QF_PROJ_QUERY "a.proj_query" -#define TN_QF_PROJ_NORM "a.proj_norm.%s" -#define TN_QF_PROJ_LINEAR "a.proj_linear.%s" -#define TN_QF_SELF_ATTN_Q "a.proj_blk.%d.self_attn_q.%s" -#define TN_QF_SELF_ATTN_K "a.proj_blk.%d.self_attn_k.%s" -#define TN_QF_SELF_ATTN_V "a.proj_blk.%d.self_attn_v.%s" -#define TN_QF_SELF_ATTN_O "a.proj_blk.%d.self_attn_out.%s" -#define TN_QF_SELF_ATTN_N "a.proj_blk.%d.self_attn_norm.%s" -#define TN_QF_CROSS_ATTN_Q "a.proj_blk.%d.cross_attn_q.%s" -#define TN_QF_CROSS_ATTN_K "a.proj_blk.%d.cross_attn_k.%s" -#define TN_QF_CROSS_ATTN_V "a.proj_blk.%d.cross_attn_v.%s" -#define TN_QF_CROSS_ATTN_O "a.proj_blk.%d.cross_attn_out.%s" -#define TN_QF_CROSS_ATTN_N "a.proj_blk.%d.cross_attn_norm.%s" -#define TN_QF_FFN_UP "a.proj_blk.%d.ffn_up.%s" -#define TN_QF_FFN_DOWN "a.proj_blk.%d.ffn_down.%s" -#define TN_QF_FFN_NORM "a.proj_blk.%d.ffn_norm.%s" +#define TN_QF_PROJ_QUERY "%s.proj_query" +#define TN_QF_PROJ_NORM "%s.proj_norm.%s" +#define TN_QF_PROJ_LINEAR "%s.proj_linear.%s" +#define TN_QF_SELF_ATTN_Q "%s.proj_blk.%d.self_attn_q.%s" +#define TN_QF_SELF_ATTN_K "%s.proj_blk.%d.self_attn_k.%s" +#define TN_QF_SELF_ATTN_V "%s.proj_blk.%d.self_attn_v.%s" +#define TN_QF_SELF_ATTN_O "%s.proj_blk.%d.self_attn_out.%s" +#define TN_QF_SELF_ATTN_N "%s.proj_blk.%d.self_attn_norm.%s" +#define TN_QF_CROSS_ATTN_Q "%s.proj_blk.%d.cross_attn_q.%s" +#define TN_QF_CROSS_ATTN_K "%s.proj_blk.%d.cross_attn_k.%s" +#define TN_QF_CROSS_ATTN_V "%s.proj_blk.%d.cross_attn_v.%s" +#define TN_QF_CROSS_ATTN_O "%s.proj_blk.%d.cross_attn_out.%s" +#define TN_QF_CROSS_ATTN_N "%s.proj_blk.%d.cross_attn_norm.%s" +#define TN_QF_FFN_UP "%s.proj_blk.%d.ffn_up.%s" +#define TN_QF_FFN_DOWN "%s.proj_blk.%d.ffn_down.%s" +#define TN_QF_FFN_NORM "%s.proj_blk.%d.ffn_norm.%s" +// multi-projector qformer (bid => projector ID) +#define TN_MULTI_PROJ_IMG_POS "v.proj_blk.%d.img_pos" +#define TN_MULTI_PROJ_QUERY "%s.proj_blk.%d.query" +#define TN_MULTI_PROJ_LINEAR "%s.proj_blk.%d.linear.%s" +#define TN_MULTI_PROJ_NORM "%s.proj_blk.%d.norm.%s" +#define TN_MULTI_PROJ_POST_NORM "%s.proj_blk.%d.post_norm.%s" // gemma4 audio conformer #define TN_A_MM_INP_PROJ "mm.a.input_projection.%s" @@ -354,6 +361,7 @@ enum projector_type { PROJECTOR_TYPE_MINICPMV4_6, PROJECTOR_TYPE_GRANITE_SPEECH, PROJECTOR_TYPE_MIMOVL, + PROJECTOR_TYPE_GRANITE4_VISION, PROJECTOR_TYPE_UNKNOWN, }; @@ -407,6 +415,7 @@ static std::map PROJECTOR_TYPE_NAMES = { { PROJECTOR_TYPE_MINICPMV4_6, "minicpmv4_6"}, { PROJECTOR_TYPE_GRANITE_SPEECH, "granite_speech"}, { PROJECTOR_TYPE_MIMOVL, "mimovl"}, + { PROJECTOR_TYPE_GRANITE4_VISION, "granite4_vision"}, }; static projector_type clip_projector_type_from_string(const std::string & str) { @@ -438,6 +447,8 @@ struct clip_image_f32 { // marks the global view in e.g., DeepSeek-OCR Models bool add_viewsep = false; + // whether a learned newline token should be appended after the image (eg Granite4 Vision) + bool add_newline = false; }; // diff --git a/tools/mtmd/clip-model.h b/tools/mtmd/clip-model.h index 238f805a9..48796b630 100644 --- a/tools/mtmd/clip-model.h +++ b/tools/mtmd/clip-model.h @@ -4,6 +4,7 @@ #include "clip.h" #include "clip-impl.h" +#include #include #include #include @@ -90,7 +91,7 @@ struct clip_hparams { float eps = 1e-6; float rope_theta = 0.0; - std::unordered_set vision_feature_layer; + std::vector vision_feature_layer; int32_t attn_window_size = 0; int32_t n_wa_pattern = 0; std::unordered_set wa_layer_indexes; // explicit layer indexes that use full attention (for irregular patterns like YoutuVL) @@ -101,6 +102,11 @@ struct clip_hparams { int32_t sam_n_head = 0; int32_t sam_n_embd = 0; + // Granite4 Vision + std::vector proj_spatial_offsets; + int32_t downsample_query_side; + int32_t downsample_window_side; + // audio int32_t n_mel_bins = 0; // whisper preprocessor int32_t proj_stack_factor = 0; // ultravox @@ -158,6 +164,10 @@ struct clip_hparams { return false; } + + bool is_vision_feature_layer(int32_t layer) const { + return std::find(vision_feature_layer.begin(), vision_feature_layer.end(), layer) != vision_feature_layer.end(); + } }; struct clip_layer { @@ -325,6 +335,20 @@ struct yasa2_stage { std::vector blocks; }; +// QFormer projector block for models with 1 (or more) QFormer projectors +// Granite Speech, Granite4 Vision +struct qf_block { + ggml_tensor * qf_proj_query = nullptr; + ggml_tensor * qf_proj_norm_w = nullptr; + ggml_tensor * qf_proj_norm_b = nullptr; + ggml_tensor * qf_proj_linear_w = nullptr; + ggml_tensor * qf_proj_linear_b = nullptr; + ggml_tensor * qf_proj_post_norm_w = nullptr; + ggml_tensor * qf_proj_post_norm_b = nullptr; + ggml_tensor * qf_proj_img_pos = nullptr; // Vision only + std::vector qf_proj_layers; +}; + struct clip_model { clip_modality modality = CLIP_MODALITY_VISION; projector_type proj_type = PROJECTOR_TYPE_MLP; @@ -589,13 +613,8 @@ struct clip_model { ggml_tensor * ctc_out_b = nullptr; ggml_tensor * ctc_out_mid_w = nullptr; ggml_tensor * ctc_out_mid_b = nullptr; - // qformer projector - ggml_tensor * qf_proj_query = nullptr; - ggml_tensor * qf_proj_norm_w = nullptr; - ggml_tensor * qf_proj_norm_b = nullptr; - ggml_tensor * qf_proj_linear_w = nullptr; - ggml_tensor * qf_proj_linear_b = nullptr; - std::vector qf_proj_layers; + // qformer projector(s) + std::vector qf_proj_blocks; bool audio_has_avgpool() const { return proj_type == PROJECTOR_TYPE_QWEN2A diff --git a/tools/mtmd/clip.cpp b/tools/mtmd/clip.cpp index 80136ed86..c12c910a1 100644 --- a/tools/mtmd/clip.cpp +++ b/tools/mtmd/clip.cpp @@ -997,6 +997,10 @@ static ggml_cgraph * clip_image_build_graph(clip_ctx * ctx, const clip_image_f32 { builder = std::make_unique(ctx, img); } break; + case PROJECTOR_TYPE_GRANITE4_VISION: + { + builder = std::make_unique(ctx, img); + } break; default: GGML_ABORT("missing cgraph builder"); } @@ -1234,12 +1238,7 @@ struct clip_model_loader { // to form the final visual features. // NOTE: gguf conversions should standardize the values of the vision feature layer to // be non-negative, since we use -1 to mark values as unset here. - std::vector vision_feature_layer; - get_arr_int(KEY_FEATURE_LAYER, vision_feature_layer, false); - // convert std::vector to std::unordered_set - for (auto & layer : vision_feature_layer) { - hparams.vision_feature_layer.insert(layer); - } + get_arr_int(KEY_FEATURE_LAYER, hparams.vision_feature_layer, false); // model-specific params switch (model.proj_type) { @@ -1627,6 +1626,23 @@ struct clip_model_loader { hparams.image_pad_color = {127, 127, 127}; hparams.image_resize_algo = RESIZE_ALGO_BILINEAR; } break; + case PROJECTOR_TYPE_GRANITE4_VISION: + { + // SigLIP tower. + hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW; + hparams.image_resize_pad = PAD_CEIL; + + get_arr_int(KEY_FEATURE_LAYER, hparams.vision_feature_layer); + get_arr_int(KEY_PROJ_SPATIAL_OFFSETS, hparams.proj_spatial_offsets); + if (hparams.vision_feature_layer.size() != hparams.proj_spatial_offsets.size()) { + throw std::runtime_error(string_format("%s: vision_feature_layer.size() %d != proj_spatial_offsets.size() %d", + hparams.vision_feature_layer.size(), hparams.proj_spatial_offsets.size())); + } + + get_u32(KEY_PROJ_SAMPLE_QUERY_SIDE, hparams.downsample_query_side); + get_u32(KEY_PROJ_SAMPLE_WINDOW_SIDE, hparams.downsample_window_side); + hparams.warmup_image_size = hparams.image_size; + } break; default: throw std::runtime_error(string_format("%s: unknown vision projector type %s\n", __func__, proj_type.c_str())); } @@ -2628,47 +2644,106 @@ struct clip_model_loader { layer.conv_pw2_b = get_tensor(string_format(TN_CONV_PW2, prefix, il, "bias")); } - model.qf_proj_query = get_tensor(TN_QF_PROJ_QUERY); - model.qf_proj_norm_w = get_tensor(string_format(TN_QF_PROJ_NORM, "weight")); - model.qf_proj_norm_b = get_tensor(string_format(TN_QF_PROJ_NORM, "bias")); - model.qf_proj_linear_w = get_tensor(string_format(TN_QF_PROJ_LINEAR, "weight")); - model.qf_proj_linear_b = get_tensor(string_format(TN_QF_PROJ_LINEAR, "bias")); + model.qf_proj_blocks.resize(1); + auto & qf = model.qf_proj_blocks[0]; + qf.qf_proj_query = get_tensor(string_format(TN_QF_PROJ_QUERY, prefix)); + qf.qf_proj_norm_w = get_tensor(string_format(TN_QF_PROJ_NORM, prefix, "weight")); + qf.qf_proj_norm_b = get_tensor(string_format(TN_QF_PROJ_NORM, prefix, "bias")); + qf.qf_proj_linear_w = get_tensor(string_format(TN_QF_PROJ_LINEAR, prefix, "weight")); + qf.qf_proj_linear_b = get_tensor(string_format(TN_QF_PROJ_LINEAR, prefix, "bias")); const int n_proj_layers = 2; - model.qf_proj_layers.resize(n_proj_layers); + qf.qf_proj_layers.resize(n_proj_layers); for (int il = 0; il < n_proj_layers; ++il) { - auto & pl = model.qf_proj_layers[il]; + auto & pl = qf.qf_proj_layers[il]; - pl.q_w = get_tensor(string_format(TN_QF_SELF_ATTN_Q, il, "weight")); - pl.q_b = get_tensor(string_format(TN_QF_SELF_ATTN_Q, il, "bias")); - pl.k_w = get_tensor(string_format(TN_QF_SELF_ATTN_K, il, "weight")); - pl.k_b = get_tensor(string_format(TN_QF_SELF_ATTN_K, il, "bias")); - pl.v_w = get_tensor(string_format(TN_QF_SELF_ATTN_V, il, "weight")); - pl.v_b = get_tensor(string_format(TN_QF_SELF_ATTN_V, il, "bias")); - pl.o_w = get_tensor(string_format(TN_QF_SELF_ATTN_O, il, "weight")); - pl.o_b = get_tensor(string_format(TN_QF_SELF_ATTN_O, il, "bias")); - pl.ln_1_w = get_tensor(string_format(TN_QF_SELF_ATTN_N, il, "weight")); - pl.ln_1_b = get_tensor(string_format(TN_QF_SELF_ATTN_N, il, "bias")); + pl.q_w = get_tensor(string_format(TN_QF_SELF_ATTN_Q, prefix, il, "weight")); + pl.q_b = get_tensor(string_format(TN_QF_SELF_ATTN_Q, prefix, il, "bias")); + pl.k_w = get_tensor(string_format(TN_QF_SELF_ATTN_K, prefix, il, "weight")); + pl.k_b = get_tensor(string_format(TN_QF_SELF_ATTN_K, prefix, il, "bias")); + pl.v_w = get_tensor(string_format(TN_QF_SELF_ATTN_V, prefix, il, "weight")); + pl.v_b = get_tensor(string_format(TN_QF_SELF_ATTN_V, prefix, il, "bias")); + pl.o_w = get_tensor(string_format(TN_QF_SELF_ATTN_O, prefix, il, "weight")); + pl.o_b = get_tensor(string_format(TN_QF_SELF_ATTN_O, prefix, il, "bias")); + pl.ln_1_w = get_tensor(string_format(TN_QF_SELF_ATTN_N, prefix, il, "weight")); + pl.ln_1_b = get_tensor(string_format(TN_QF_SELF_ATTN_N, prefix, il, "bias")); - pl.cross_attn_q_w = get_tensor(string_format(TN_QF_CROSS_ATTN_Q, il, "weight")); - pl.cross_attn_q_b = get_tensor(string_format(TN_QF_CROSS_ATTN_Q, il, "bias")); - pl.cross_attn_k_w = get_tensor(string_format(TN_QF_CROSS_ATTN_K, il, "weight")); - pl.cross_attn_k_b = get_tensor(string_format(TN_QF_CROSS_ATTN_K, il, "bias")); - pl.cross_attn_v_w = get_tensor(string_format(TN_QF_CROSS_ATTN_V, il, "weight")); - pl.cross_attn_v_b = get_tensor(string_format(TN_QF_CROSS_ATTN_V, il, "bias")); - pl.cross_attn_o_w = get_tensor(string_format(TN_QF_CROSS_ATTN_O, il, "weight")); - pl.cross_attn_o_b = get_tensor(string_format(TN_QF_CROSS_ATTN_O, il, "bias")); - pl.cross_attn_norm_w = get_tensor(string_format(TN_QF_CROSS_ATTN_N, il, "weight")); - pl.cross_attn_norm_b = get_tensor(string_format(TN_QF_CROSS_ATTN_N, il, "bias")); + pl.cross_attn_q_w = get_tensor(string_format(TN_QF_CROSS_ATTN_Q, prefix, il, "weight")); + pl.cross_attn_q_b = get_tensor(string_format(TN_QF_CROSS_ATTN_Q, prefix, il, "bias")); + pl.cross_attn_k_w = get_tensor(string_format(TN_QF_CROSS_ATTN_K, prefix, il, "weight")); + pl.cross_attn_k_b = get_tensor(string_format(TN_QF_CROSS_ATTN_K, prefix, il, "bias")); + pl.cross_attn_v_w = get_tensor(string_format(TN_QF_CROSS_ATTN_V, prefix, il, "weight")); + pl.cross_attn_v_b = get_tensor(string_format(TN_QF_CROSS_ATTN_V, prefix, il, "bias")); + pl.cross_attn_o_w = get_tensor(string_format(TN_QF_CROSS_ATTN_O, prefix, il, "weight")); + pl.cross_attn_o_b = get_tensor(string_format(TN_QF_CROSS_ATTN_O, prefix, il, "bias")); + pl.cross_attn_norm_w = get_tensor(string_format(TN_QF_CROSS_ATTN_N, prefix, il, "weight")); + pl.cross_attn_norm_b = get_tensor(string_format(TN_QF_CROSS_ATTN_N, prefix, il, "bias")); - pl.ff_up_w = get_tensor(string_format(TN_QF_FFN_UP, il, "weight")); - pl.ff_up_b = get_tensor(string_format(TN_QF_FFN_UP, il, "bias")); - pl.ff_down_w = get_tensor(string_format(TN_QF_FFN_DOWN, il, "weight")); - pl.ff_down_b = get_tensor(string_format(TN_QF_FFN_DOWN, il, "bias")); - pl.ln_2_w = get_tensor(string_format(TN_QF_FFN_NORM, il, "weight")); - pl.ln_2_b = get_tensor(string_format(TN_QF_FFN_NORM, il, "bias")); + pl.ff_up_w = get_tensor(string_format(TN_QF_FFN_UP, prefix, il, "weight")); + pl.ff_up_b = get_tensor(string_format(TN_QF_FFN_UP, prefix, il, "bias")); + pl.ff_down_w = get_tensor(string_format(TN_QF_FFN_DOWN, prefix, il, "weight")); + pl.ff_down_b = get_tensor(string_format(TN_QF_FFN_DOWN, prefix, il, "bias")); + pl.ln_2_w = get_tensor(string_format(TN_QF_FFN_NORM, prefix, il, "weight")); + pl.ln_2_b = get_tensor(string_format(TN_QF_FFN_NORM, prefix, il, "bias")); } } break; + case PROJECTOR_TYPE_GRANITE4_VISION: + { + // image_newline lives at the top-level. + model.image_newline = get_tensor(TN_IMAGE_NEWLINE); + + // Load separate layerwise and spatial projector tensors + const auto projector_count = hparams.vision_feature_layer.size(); + model.qf_proj_blocks.resize(projector_count); + for (size_t bid = 0; bid < projector_count; ++bid) { + auto & b = model.qf_proj_blocks[bid]; + + // non-layerwise tensors + b.qf_proj_img_pos = get_tensor(string_format(TN_MULTI_PROJ_IMG_POS, bid)); + b.qf_proj_query = get_tensor(string_format(TN_MULTI_PROJ_QUERY, prefix, bid)); + b.qf_proj_linear_w = get_tensor(string_format(TN_MULTI_PROJ_LINEAR, prefix, bid, "weight")); + b.qf_proj_linear_b = get_tensor(string_format(TN_MULTI_PROJ_LINEAR, prefix, bid, "bias")); + b.qf_proj_norm_w = get_tensor(string_format(TN_MULTI_PROJ_NORM, prefix, bid, "weight")); + b.qf_proj_norm_b = get_tensor(string_format(TN_MULTI_PROJ_NORM, prefix, bid, "bias")); + b.qf_proj_post_norm_w = get_tensor(string_format(TN_MULTI_PROJ_POST_NORM, prefix, bid, "weight")); + b.qf_proj_post_norm_b = get_tensor(string_format(TN_MULTI_PROJ_POST_NORM, prefix, bid, "bias")); + + // laywerwise tensors + // NOTE: If any model uses multi-layer qformers, this will need to change + b.qf_proj_layers.resize(1); + auto & pl = b.qf_proj_layers[0]; + + pl.q_w = get_tensor(string_format(TN_QF_SELF_ATTN_Q, prefix, bid, "weight")); + pl.q_b = get_tensor(string_format(TN_QF_SELF_ATTN_Q, prefix, bid, "bias")); + pl.k_w = get_tensor(string_format(TN_QF_SELF_ATTN_K, prefix, bid, "weight")); + pl.k_b = get_tensor(string_format(TN_QF_SELF_ATTN_K, prefix, bid, "bias")); + pl.v_w = get_tensor(string_format(TN_QF_SELF_ATTN_V, prefix, bid, "weight")); + pl.v_b = get_tensor(string_format(TN_QF_SELF_ATTN_V, prefix, bid, "bias")); + pl.o_w = get_tensor(string_format(TN_QF_SELF_ATTN_O, prefix, bid, "weight")); + pl.o_b = get_tensor(string_format(TN_QF_SELF_ATTN_O, prefix, bid, "bias")); + pl.ln_1_w = get_tensor(string_format(TN_QF_SELF_ATTN_N, prefix, bid, "weight")); + pl.ln_1_b = get_tensor(string_format(TN_QF_SELF_ATTN_N, prefix, bid, "bias")); + + pl.cross_attn_q_w = get_tensor(string_format(TN_QF_CROSS_ATTN_Q, prefix, bid, "weight")); + pl.cross_attn_q_b = get_tensor(string_format(TN_QF_CROSS_ATTN_Q, prefix, bid, "bias")); + pl.cross_attn_k_w = get_tensor(string_format(TN_QF_CROSS_ATTN_K, prefix, bid, "weight")); + pl.cross_attn_k_b = get_tensor(string_format(TN_QF_CROSS_ATTN_K, prefix, bid, "bias")); + pl.cross_attn_v_w = get_tensor(string_format(TN_QF_CROSS_ATTN_V, prefix, bid, "weight")); + pl.cross_attn_v_b = get_tensor(string_format(TN_QF_CROSS_ATTN_V, prefix, bid, "bias")); + pl.cross_attn_o_w = get_tensor(string_format(TN_QF_CROSS_ATTN_O, prefix, bid, "weight")); + pl.cross_attn_o_b = get_tensor(string_format(TN_QF_CROSS_ATTN_O, prefix, bid, "bias")); + pl.cross_attn_norm_w = get_tensor(string_format(TN_QF_CROSS_ATTN_N, prefix, bid, "weight")); + pl.cross_attn_norm_b = get_tensor(string_format(TN_QF_CROSS_ATTN_N, prefix, bid, "bias")); + + pl.ff_up_w = get_tensor(string_format(TN_QF_FFN_UP, prefix, bid, "weight")); + pl.ff_up_b = get_tensor(string_format(TN_QF_FFN_UP, prefix, bid, "bias")); + pl.ff_down_w = get_tensor(string_format(TN_QF_FFN_DOWN, prefix, bid, "weight")); + pl.ff_down_b = get_tensor(string_format(TN_QF_FFN_DOWN, prefix, bid, "bias")); + pl.ln_2_w = get_tensor(string_format(TN_QF_FFN_NORM, prefix, bid, "weight")); + pl.ln_2_b = get_tensor(string_format(TN_QF_FFN_NORM, prefix, bid, "bias")); + } + + } break; default: GGML_ASSERT(false && "unknown projector type"); } @@ -3085,10 +3160,6 @@ void clip_build_img_from_pixels(const unsigned char * rgb_pixels, int nx, int ny memcpy(img->buf.data(), rgb_pixels, img->buf.size()); } -ggml_tensor * clip_get_newline_tensor(const struct clip_ctx * ctx) { - return ctx->model.image_newline; -} - void clip_free(clip_ctx * ctx) { if (ctx == nullptr) { return; @@ -3397,6 +3468,23 @@ int clip_n_output_tokens(const struct clip_ctx * ctx, struct clip_image_f32 * im const int ds = ctx->model.hparams.audio_proj_downsample_rate; n_patches = ((img->nx + ws - 1) / ws) * (ws / ds); } break; + case PROJECTOR_TYPE_GRANITE4_VISION: + { + // Per-tile output token count: each projector block outputs + // query_side^2 tokens per window × n^2 windows. + // For 384×384 input: n = 24/8 = 3, query_side = 4 → 144. + const int window_side = ctx->model.hparams.downsample_window_side; + const int query_side = ctx->model.hparams.downsample_query_side; + const int side = img->nx / params.patch_size; + const int n = side / window_side; + n_patches = (query_side * n) * (query_side * n); + if (img->add_newline) { + // For single-tile case: append 1 newline row. + // For multi-tile rowwise: handled by caller, but here we + // report the per-tile count including one trailing newline. + n_patches += 1; + } + } break; default: GGML_ABORT("unsupported projector type"); } @@ -4229,6 +4317,82 @@ bool clip_image_batch_encode(clip_ctx * ctx, const int n_threads, const clip_ima set_input_f32("attn_mask", mask); } } break; + case PROJECTOR_TYPE_GRANITE4_VISION: + { + // Granite Vision 4.1 uses precomputed permutation index + // tensors to express the _win / _unwin / spatial sampling + // reshapes as ggml_get_rows gathers. The names are set + // by g4v_gather() in models/granite4-vision.cpp. + const int patch_size = model.hparams.patch_size; + const int image_side = imgs.entries.front()->nx / patch_size; + const int window_side = hparams.downsample_window_side; + const int query_side = hparams.downsample_query_side; + const int n = image_side / window_side; + const int new_side = n * query_side; + + // Builds the raster→window permutation indices for a + // (side, side) grid split into (n × n) windows of (win × win) + // tokens each. dst[w * win*win + p] = source raster index. + auto make_win_idx = [](int side, int win) { + const int nn = side / win; + std::vector idx(static_cast(side) * side); + for (int wy = 0; wy < nn; ++wy) { + for (int wx = 0; wx < nn; ++wx) { + for (int iy = 0; iy < win; ++iy) { + for (int ix = 0; ix < win; ++ix) { + const int w = wy * nn + wx; + const int p = iy * win + ix; + const int y = wy * win + iy; + const int x = wx * win + ix; + idx[static_cast(w) * (win*win) + p] = y * side + x; + } + } + } + } + return idx; + }; + + auto make_unwin_idx = [&](int side, int win) { + const std::vector fwd = make_win_idx(side, win); + std::vector inv(fwd.size()); + for (size_t i = 0; i < fwd.size(); ++i) { + inv[fwd[i]] = static_cast(i); + } + return inv; + }; + + auto make_spatial_idx = [](int side, int offset) { + const int off_y = (offset >> 1) & 1; + const int off_x = offset & 1; + const int new_s = side / 2; + std::vector idx(static_cast(new_s) * new_s); + for (int y = 0; y < new_s; ++y) { + for (int x = 0; x < new_s; ++x) { + idx[y * new_s + x] = (y * 2 + off_y) * side + (x * 2 + off_x); + } + } + return idx; + }; + + auto upload = [&](const std::string & name, const std::vector & idx) { + ggml_tensor * t = ggml_graph_get_tensor(gf, name.c_str()); + GGML_ASSERT(t); + ggml_backend_tensor_set(t, idx.data(), 0, idx.size() * sizeof(int32_t)); + }; + + // Stage 1b only uses block 0's permutations; future stages + // will upload all blocks. + for (size_t bid = 0; bid < hparams.vision_feature_layer.size(); ++bid) { + const std::string prefix = "g4v_blk" + std::to_string(bid) + "_"; + upload(prefix + "win_idx", make_win_idx(image_side, window_side)); + upload(prefix + "qwin_idx", make_win_idx(new_side, query_side)); + upload(prefix + "unwin_idx", make_unwin_idx(new_side, query_side)); + const auto spatial_offset = hparams.proj_spatial_offsets[bid]; + if (spatial_offset >= 0) { + upload(prefix + "spatial_idx", make_spatial_idx(image_side,spatial_offset)); + } + } + } break; default: GGML_ABORT("Unknown projector type"); } @@ -4384,7 +4548,9 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) { case PROJECTOR_TYPE_LFM2A: return ctx->model.position_embeddings->ne[0]; case PROJECTOR_TYPE_GRANITE_SPEECH: - return ctx->model.qf_proj_linear_w->ne[1]; + return ctx->model.qf_proj_blocks[0].qf_proj_linear_w->ne[1]; + case PROJECTOR_TYPE_GRANITE4_VISION: + return ctx->model.qf_proj_blocks.size() * ctx->model.hparams.projection_dim; case PROJECTOR_TYPE_GLM4V: return ctx->model.mm_ffn_down_w->ne[1]; default: diff --git a/tools/mtmd/clip.h b/tools/mtmd/clip.h index 9b807ffa7..a62c9d618 100644 --- a/tools/mtmd/clip.h +++ b/tools/mtmd/clip.h @@ -100,8 +100,6 @@ struct clip_image_f32 * clip_image_f32_get_img(const struct clip_image_f32_batch */ void clip_build_img_from_pixels(const unsigned char * rgb_pixels, int nx, int ny, struct clip_image_u8 * img); -struct ggml_tensor * clip_get_newline_tensor(const struct clip_ctx * ctx); - bool clip_image_encode (struct clip_ctx * ctx, int n_threads, struct clip_image_f32 * img, float * vec); bool clip_image_batch_encode(struct clip_ctx * ctx, int n_threads, const struct clip_image_f32_batch * imgs, float * vec); diff --git a/tools/mtmd/models/granite-speech.cpp b/tools/mtmd/models/granite-speech.cpp index c7e3794a4..5e66f75d0 100644 --- a/tools/mtmd/models/granite-speech.cpp +++ b/tools/mtmd/models/granite-speech.cpp @@ -199,8 +199,8 @@ ggml_cgraph * clip_graph_granite_speech::build() { ggml_tensor * enc_windows = ggml_reshape_3d(ctx0, cur, n_embd, window_size, nblocks_proj); - ggml_tensor * queries = build_norm(model.qf_proj_query, - model.qf_proj_norm_w, model.qf_proj_norm_b, + ggml_tensor * queries = build_norm(model.qf_proj_blocks[0].qf_proj_query, + model.qf_proj_blocks[0].qf_proj_norm_w, model.qf_proj_blocks[0].qf_proj_norm_b, NORM_TYPE_NORMAL, proj_eps, -1); { ggml_tensor * q_3d = ggml_reshape_3d(ctx0, queries, n_embd, num_queries, 1); @@ -209,8 +209,8 @@ ggml_cgraph * clip_graph_granite_speech::build() { queries = ggml_repeat(ctx0, q_3d, q_shape); } - for (int il = 0; il < (int)model.qf_proj_layers.size(); il++) { - const auto & pl = model.qf_proj_layers[il]; + for (int il = 0; il < (int)model.qf_proj_blocks[0].qf_proj_layers.size(); il++) { + const auto & pl = model.qf_proj_blocks[0].qf_proj_layers[il]; // self-attention { @@ -265,7 +265,7 @@ ggml_cgraph * clip_graph_granite_speech::build() { } cur = ggml_reshape_2d(ctx0, queries, n_embd, num_queries * nblocks_proj); - cur = ggml_add(ctx0, build_mm(model.qf_proj_linear_w, cur), model.qf_proj_linear_b); + cur = ggml_add(ctx0, build_mm(model.qf_proj_blocks[0].qf_proj_linear_w, cur), model.qf_proj_blocks[0].qf_proj_linear_b); cb(cur, "projector_out", -1); } diff --git a/tools/mtmd/models/granite4-vision.cpp b/tools/mtmd/models/granite4-vision.cpp new file mode 100644 index 000000000..9adb6f0fd --- /dev/null +++ b/tools/mtmd/models/granite4-vision.cpp @@ -0,0 +1,339 @@ +#include "models.h" +#include "../clip-impl.h" +#include "../clip-model.h" + +#include +#include +#include +#include +#include + +/* + * Granite Vision 4.1 clip graph + * + * Stage 1a: SigLIP vision tower (N layers, post-norm) + * Stage 1b: WindowQFormer blocks (deepstack + spatial) + * Stage 1c: Concatenate and pack outputs + * Stage 1d: Append newline tokens if add_newline is set + */ + +// --------------------------------------------------------------------------- +// Member method implementations +// --------------------------------------------------------------------------- + +ggml_tensor * clip_graph_granite4_vision::gather( + ggml_tensor * src, + const std::string & name, + int idx_len) { + ggml_tensor * idx = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, idx_len); + ggml_set_name(idx, name.c_str()); + ggml_set_input(idx); + return ggml_get_rows(ctx0, src, idx); +} + +ggml_tensor * clip_graph_granite4_vision::interp_down( + ggml_tensor * src, + int side, + int new_side) { + const int n_embd = src->ne[0]; + ggml_tensor * t = ggml_reshape_4d(ctx0, src, n_embd, side, side, 1); + t = ggml_cont(ctx0, ggml_permute(ctx0, t, 2, 0, 1, 3)); + const int kernel = side / new_side; + t = ggml_pool_2d(ctx0, t, GGML_OP_POOL_AVG, kernel, kernel, kernel, kernel, 0, 0); + t = ggml_cont(ctx0, ggml_permute(ctx0, t, 1, 2, 0, 3)); + return ggml_reshape_2d(ctx0, t, n_embd, new_side * new_side); +} + +// --------------------------------------------------------------------------- +// build_block - WindowQFormer block implementation +// --------------------------------------------------------------------------- + +ggml_tensor * clip_graph_granite4_vision::build_block( + const qf_block & blk, + ggml_tensor * h, + int bid, + int spatial_offset, + int image_side, + int window_side, + int query_side, + float qformer_eps) { + + const int n_embd = h->ne[0]; + GGML_ASSERT(h->ne[1] == image_side * image_side); + const int n = image_side / window_side; + const int new_side = n * query_side; + const int n_windows = n * n; + const int enc_len = window_side * window_side; + const int query_len = query_side * query_side; + + auto cbx = [&](ggml_tensor * & t, const char * step) { + const std::string name = "g4v_blk" + std::to_string(bid) + "_" + step; + ggml_set_name(t, name.c_str()); + }; + + // 1. Top-level LN + cbx(h, "inp"); + ggml_tensor * x = build_norm(h, blk.qf_proj_norm_w, blk.qf_proj_norm_b, NORM_TYPE_NORMAL, eps, bid); + cbx(x, "norm"); + + // 2. enc = _win(x, image_side, window_side) + ggml_tensor * enc; + { + ggml_tensor * enc_flat = gather(x, + "g4v_blk" + std::to_string(bid) + "_win_idx", + image_side * image_side); + enc = ggml_reshape_3d(ctx0, enc_flat, n_embd, enc_len, n_windows); + } + cbx(enc, "enc"); + + // 3. downsampled = downsampler(x) + ggml_tensor * d; + (void) spatial_offset; + if (spatial_offset >= 0) { + d = gather(x, + "g4v_blk" + std::to_string(bid) + "_spatial_idx", + new_side * new_side); + } else { + d = interp_down(x, image_side, new_side); + } + cbx(d, "downsampled"); + + // 4. query_embeds = query + _win(d, new_side, query_side) + ggml_tensor * q_in; + { + ggml_tensor * dw_flat = gather(d, + "g4v_blk" + std::to_string(bid) + "_qwin_idx", + new_side * new_side); + ggml_tensor * dw = ggml_reshape_3d(ctx0, dw_flat, n_embd, query_len, n_windows); + q_in = ggml_add(ctx0, dw, blk.qf_proj_query); + } + cbx(q_in, "query_embeds"); + + // 5. encoder_embeds = enc + image_positions → (C, enc_len, n_windows) + ggml_tensor * e_in = ggml_add(ctx0, enc, blk.qf_proj_img_pos); + cbx(e_in, "encoder_embeds"); + + // 6. Qformer forward. + ggml_tensor * q = build_norm(q_in, blk.qf_proj_post_norm_w, blk.qf_proj_post_norm_b, NORM_TYPE_NORMAL, qformer_eps, bid); + + // Helper for linear projections with window batching + auto linear = [&](ggml_tensor * x, ggml_tensor * w, ggml_tensor * b) -> ggml_tensor * { + ggml_tensor * t = ggml_reshape_2d(ctx0, x, x->ne[0], x->ne[1] * x->ne[2]); + t = build_mm(w, t); + if (b) t = ggml_add(ctx0, t, b); + return t; + }; + + // Get the single QFormer layer + GGML_ASSERT(blk.qf_proj_layers.size() == 1); + const auto & pl = blk.qf_proj_layers[0]; + + // 6a. Self-attention + ggml_tensor * sa_out; + { + const int d_h = 64; + const int n_head = n_embd / d_h; + const int nq = q->ne[1]; + const float scale = 1.0f / std::sqrt((float) d_h); + + ggml_tensor * Q = linear(q, pl.q_w, pl.q_b); + ggml_tensor * K = linear(q, pl.k_w, pl.k_b); + ggml_tensor * V = linear(q, pl.v_w, pl.v_b); + + Q = ggml_reshape_4d(ctx0, Q, d_h, n_head, nq, n_windows); + K = ggml_reshape_4d(ctx0, K, d_h, n_head, nq, n_windows); + V = ggml_reshape_4d(ctx0, V, d_h, n_head, nq, n_windows); + + sa_out = build_attn(pl.o_w, pl.o_b, Q, K, V, nullptr, scale, bid); + sa_out = ggml_reshape_3d(ctx0, sa_out, n_embd, nq, n_windows); + + sa_out = ggml_add(ctx0, sa_out, q); + sa_out = build_norm(sa_out, pl.ln_1_w, pl.ln_1_b, + NORM_TYPE_NORMAL, qformer_eps, bid); + } + cbx(sa_out, "sa_out"); + + // 6b. Cross-attention + ggml_tensor * ca_out; + { + const int d_h = 64; + const int n_head = n_embd / d_h; + const int nq = sa_out->ne[1]; + const int nkv = e_in->ne[1]; + const float scale = 1.0f / std::sqrt((float) d_h); + + ggml_tensor * Q = linear(sa_out, pl.cross_attn_q_w, pl.cross_attn_q_b); + ggml_tensor * K = linear(e_in, pl.cross_attn_k_w, pl.cross_attn_k_b); + ggml_tensor * V = linear(e_in, pl.cross_attn_v_w, pl.cross_attn_v_b); + + Q = ggml_reshape_4d(ctx0, Q, d_h, n_head, nq, n_windows); + K = ggml_reshape_4d(ctx0, K, d_h, n_head, nkv, n_windows); + V = ggml_reshape_4d(ctx0, V, d_h, n_head, nkv, n_windows); + + ca_out = build_attn(pl.cross_attn_o_w, pl.cross_attn_o_b, + Q, K, V, nullptr, scale, bid); + ca_out = ggml_reshape_3d(ctx0, ca_out, n_embd, nq, n_windows); + + ca_out = ggml_add(ctx0, ca_out, sa_out); + ca_out = build_norm(ca_out, pl.cross_attn_norm_w, pl.cross_attn_norm_b, + NORM_TYPE_NORMAL, qformer_eps, bid); + } + cbx(ca_out, "ca_out"); + + // 6c. FFN + ggml_tensor * ffn; + { + ggml_tensor * t = ggml_reshape_2d(ctx0, ca_out, n_embd, query_len * n_windows); + t = build_mm(pl.ff_up_w, t); + if (pl.ff_up_b) t = ggml_add(ctx0, t, pl.ff_up_b); + t = ggml_gelu_erf(ctx0, t); + t = build_mm(pl.ff_down_w, t); + if (pl.ff_down_b) t = ggml_add(ctx0, t, pl.ff_down_b); + t = ggml_reshape_3d(ctx0, t, n_embd, query_len, n_windows); + ffn = ggml_add(ctx0, t, ca_out); + ffn = build_norm(ffn, pl.ln_2_w, pl.ln_2_b, NORM_TYPE_NORMAL, qformer_eps, bid); + } + cbx(ffn, "qformer_out"); + + // 7. _unwin back to raster + ggml_tensor * unwinned; + { + ggml_tensor * flat = ggml_reshape_2d(ctx0, ffn, n_embd, query_len * n_windows); + unwinned = gather(flat, + "g4v_blk" + std::to_string(bid) + "_unwin_idx", + new_side * new_side); + } + cbx(unwinned, "unwin"); + + // 8. out_linear + ggml_tensor * out = build_mm(blk.qf_proj_linear_w, unwinned); + if (blk.qf_proj_linear_b) out = ggml_add(ctx0, out, blk.qf_proj_linear_b); + cbx(out, "out"); + + return out; +} + +// --------------------------------------------------------------------------- +// build() - top-level graph +// --------------------------------------------------------------------------- + +// Build the K-tiled, base-scaled newline row tensor. +// Shape: (n_mmproj_embd, 1) +ggml_tensor * clip_graph_granite4_vision::build_newline_row(ggml_context * ctx0) { + const int K = (int) model.qf_proj_blocks.size(); + GGML_ASSERT(K > 0); + GGML_ASSERT(n_mmproj_embd % K == 0); + const int projection_dim = n_mmproj_embd / K; + GGML_ASSERT(model.image_newline != nullptr); + GGML_ASSERT(ggml_nelements(model.image_newline) == projection_dim); + + // Build newline_row[k*projection_dim + d] = nl[d] * (k == 0 ? base : 1.0) + ggml_tensor * nl = model.image_newline; // (projection_dim,) + ggml_tensor * nl_first_2d = ggml_reshape_2d(ctx0, nl, projection_dim, 1); + ggml_tensor * nl_row_2d; + if (K == 1) { + nl_row_2d = nl_first_2d; + } else { + ggml_tensor * nl_2d = ggml_reshape_2d(ctx0, nl, projection_dim, 1); + ggml_tensor * rest_template = ggml_new_tensor_2d( + ctx0, GGML_TYPE_F32, projection_dim, K - 1); + ggml_tensor * nl_rest = ggml_repeat(ctx0, nl_2d, rest_template); + nl_row_2d = ggml_concat(ctx0, nl_first_2d, nl_rest, 1); // (projection_dim, K) + } + nl_row_2d = ggml_cont(ctx0, nl_row_2d); + return ggml_reshape_2d(ctx0, nl_row_2d, n_mmproj_embd, 1); +} + +// Append a single newline row at the end of the tile output. +ggml_tensor * clip_graph_granite4_vision::append_rowwise_newlines(ggml_context * ctx0, ggml_tensor * tile_output) { + // For the single-tile case, append one newline row at the end. + // For the multi-tile rowwise case, this will be called per-tile + // (though currently only the single-tile path uses it). + ggml_tensor * nl_row = build_newline_row(ctx0); + return ggml_concat(ctx0, tile_output, nl_row, 1); +} + +ggml_cgraph * clip_graph_granite4_vision::build() { + GGML_ASSERT(model.patch_embeddings_0 != nullptr); + GGML_ASSERT(model.position_embeddings != nullptr); + GGML_ASSERT(model.class_embedding == nullptr); + GGML_ASSERT(!model.qf_proj_blocks.empty()); + + // --- Stage 1a: SigLIP encoder producing intermediate hidden states --- + ggml_tensor * inp = build_inp(); + inp = ggml_add(ctx0, inp, model.position_embeddings); + cb(inp, "pos_embed", -1); + + ggml_tensor * inpL = inp; + std::vector layer_outs(n_layer, nullptr); + + for (int il = 0; il < n_layer; ++il) { + const auto & layer = model.layers[il]; + ggml_tensor * cur = inpL; + + cur = build_norm(cur, layer.ln_1_w, layer.ln_1_b, NORM_TYPE_NORMAL, eps, il); + + // Self-attention + ggml_tensor * Qcur = build_mm(layer.q_w, cur); + if (layer.q_b) Qcur = ggml_add(ctx0, Qcur, layer.q_b); + ggml_tensor * Kcur = build_mm(layer.k_w, cur); + if (layer.k_b) Kcur = ggml_add(ctx0, Kcur, layer.k_b); + ggml_tensor * Vcur = build_mm(layer.v_w, cur); + if (layer.v_b) Vcur = ggml_add(ctx0, Vcur, layer.v_b); + + Qcur = ggml_reshape_3d(ctx0, Qcur, d_head, n_head, n_patches); + Kcur = ggml_reshape_3d(ctx0, Kcur, d_head, n_head, n_patches); + Vcur = ggml_reshape_3d(ctx0, Vcur, d_head, n_head, n_patches); + + cur = build_attn(layer.o_w, layer.o_b, + Qcur, Kcur, Vcur, nullptr, kq_scale, il); + + cur = ggml_add(ctx0, cur, inpL); + inpL = cur; + + cur = build_norm(cur, layer.ln_2_w, layer.ln_2_b, NORM_TYPE_NORMAL, eps, il); + cur = build_ffn(cur, + layer.ff_up_w, layer.ff_up_b, + layer.ff_gate_w, layer.ff_gate_b, + layer.ff_down_w, layer.ff_down_b, + hparams.ffn_op, il); + cur = ggml_add(ctx0, inpL, cur); + cb(cur, "layer_out", il); + layer_outs[il] = cur; + inpL = cur; + } + + // --- Stage 1b/1c: WindowQFormer blocks --- + const int projector_count = hparams.vision_feature_layer.size(); + const float qformer_eps = 1e-12f; + + ggml_tensor * mmproj = nullptr; + for (int bid = 0; bid < projector_count; ++bid) { + const auto & blk = model.qf_proj_blocks[bid]; + + int vlayer = hparams.vision_feature_layer[bid]; + GGML_ASSERT(vlayer >= 0 && vlayer < n_layer); + ggml_tensor * h = layer_outs[vlayer]; + + ggml_tensor * stream = build_block( + blk, h, bid, + hparams.proj_spatial_offsets[bid], + n_patches_x, + hparams.downsample_window_side, + hparams.downsample_query_side, + qformer_eps); + cb(stream, (std::string("proj_") + std::to_string(bid) + std::string("_v_out")).c_str(), vlayer); + mmproj = mmproj ? ggml_concat(ctx0, mmproj, stream, 0) : stream; + } + + // --- Stage 1d: Append newline tokens if add_newline is set --- + if (add_newline) { + mmproj = append_rowwise_newlines(ctx0, mmproj); + ggml_set_name(mmproj, "g4v_mmproj_out_nl"); + } else { + ggml_set_name(mmproj, "g4v_mmproj_out"); + } + ggml_build_forward_expand(gf, mmproj); + + return gf; +} diff --git a/tools/mtmd/models/llava.cpp b/tools/mtmd/models/llava.cpp index 4af17ccfe..5aa3d2f0f 100644 --- a/tools/mtmd/models/llava.cpp +++ b/tools/mtmd/models/llava.cpp @@ -51,7 +51,6 @@ ggml_cgraph * clip_graph_llava::build() { } std::vector embedding_stack; - const auto & vision_feature_layer = hparams.vision_feature_layer; // loop over layers for (int il = 0; il < max_feature_layer; il++) { @@ -60,7 +59,7 @@ ggml_cgraph * clip_graph_llava::build() { // If this is an embedding feature layer, save the output. // NOTE: 0 index here refers to the input to the encoder. - if (vision_feature_layer.find(il) != vision_feature_layer.end()) { + if (hparams.is_vision_feature_layer(il)) { embedding_stack.push_back(cur); } @@ -135,7 +134,7 @@ ggml_cgraph * clip_graph_llava::build() { // process vision feature layers (used by granite) { // final layer is a vision feature layer - if (vision_feature_layer.find(max_feature_layer) != vision_feature_layer.end()) { + if (hparams.is_vision_feature_layer(max_feature_layer)) { embedding_stack.push_back(inpL); } diff --git a/tools/mtmd/models/models.h b/tools/mtmd/models/models.h index b882f800d..d1865103b 100644 --- a/tools/mtmd/models/models.h +++ b/tools/mtmd/models/models.h @@ -211,3 +211,26 @@ struct clip_graph_exaone4_5 : clip_graph { clip_graph_exaone4_5(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {} ggml_cgraph * build() override; }; + +struct clip_graph_granite4_vision : clip_graph { + clip_graph_granite4_vision(clip_ctx * ctx, const clip_image_f32 & img) + : clip_graph(ctx, img), + add_newline(img.add_newline) {} + + ggml_cgraph * build() override; + +private: + // The graph is per-tile since only batch-size 1 is supported in clip. As + // such, this value is set at construct time based on the tile that will be + // encoded, then used during build to determine how to handle newlines. + const bool add_newline; + + ggml_tensor * gather(ggml_tensor * src, const std::string & name, int idx_len); + ggml_tensor * interp_down(ggml_tensor * src, int side, int new_side); + ggml_tensor * build_block(const qf_block & blk, ggml_tensor * h, int bid, + int spatial_offset, int image_side, int window_side, + int query_side, float qformer_eps); + + ggml_tensor * build_newline_row(ggml_context * ctx0); + ggml_tensor * append_rowwise_newlines(ggml_context * ctx0, ggml_tensor * tile_output); +}; diff --git a/tools/mtmd/mtmd.cpp b/tools/mtmd/mtmd.cpp index 0b5caa6cb..260f30756 100644 --- a/tools/mtmd/mtmd.cpp +++ b/tools/mtmd/mtmd.cpp @@ -513,6 +513,12 @@ struct mtmd_context { img_end = ""; image_preproc = std::make_unique(ctx_v); } break; + case PROJECTOR_TYPE_GRANITE4_VISION: + { + img_beg = ""; + img_end = ""; + image_preproc = std::make_unique(ctx_v); + } break; default: throw std::runtime_error(string_format("%s: unexpected vision projector type %d\n", __func__, proj)); } @@ -808,6 +814,21 @@ struct mtmd_tokenizer { return 2; } + // Annotate llava-next style tiles so clip_n_output_tokens accounts + // for per-tile newline injection. + if (ctx->proj_type_v() == PROJECTOR_TYPE_GRANITE4_VISION) { + if (batch_f32.entries.size() == 1) { + // Single-tile (overview only): append one newline row. + batch_f32.entries[0]->add_newline = true; + } else { + // Multi-tile: overview gets no newline, grid tiles get one. + batch_f32.entries[0]->add_newline = false; + for (size_t i = 1; i < batch_f32.entries.size(); ++i) { + batch_f32.entries[i]->add_newline = true; + } + } + } + // handle llava-uhd style preprocessing const bool has_tiling_grid = batch_f32.grid_x > 0 && batch_f32.grid_y > 0; if ( @@ -872,9 +893,10 @@ struct mtmd_tokenizer { } } else { + size_t n_tokens = 0; - for (const auto & entry : batch_f32.entries) { - n_tokens += clip_n_output_tokens(ctx->ctx_v, entry.get()); + for (const auto & e : batch_f32.entries) { + n_tokens += clip_n_output_tokens(ctx->ctx_v, e.get()); } mtmd_image_tokens_ptr image_tokens(new mtmd_image_tokens); @@ -1111,7 +1133,8 @@ int32_t mtmd_encode(mtmd_context * ctx, const mtmd_image_tokens * image_tokens) || proj_type == PROJECTOR_TYPE_MINICPMV || proj_type == PROJECTOR_TYPE_GLM_EDGE || proj_type == PROJECTOR_TYPE_INTERNVL - || proj_type == PROJECTOR_TYPE_DEEPSEEKOCR2) { + || proj_type == PROJECTOR_TYPE_DEEPSEEKOCR2 + || proj_type == PROJECTOR_TYPE_GRANITE4_VISION) { // TODO @ngxson : llava does not support batched encoding ; this should be fixed inside clip_image_batch_encode() const auto & entries = image_tokens->batch_f32.entries; // entries may have different token counts From c4a278d68efa17811006f2123a84081dac03fac7 Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Fri, 5 Jun 2026 18:12:27 +0200 Subject: [PATCH 2/8] model: fix build failed (#24193) --- src/llama-model.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 6808ad044..137d3501e 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1764,7 +1764,7 @@ void llama_model::print_info() const { [](const auto & entry) { return entry >= 0; })) { LLAMA_LOG_INFO("%s: deepstack_mapping_arr = %s\n", __func__, print_f([&](uint32_t il) { return hparams.deepstack_mapping_arr[il]; }, - hparams.n_layer).c_str()); + hparams.n_layer()).c_str()); } // MRoPE (Multi-axis Rotary Position Embedding) sections if (const auto & s = hparams.rope_sections; s[0] || s[1] || s[2] || s[3]) { From e82beaa60d3bef6b6234656389239facf5ba2a3a Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Fri, 5 Jun 2026 19:44:40 +0200 Subject: [PATCH 3/8] vulkan: add fwht support for Intel with shmem reduction (#23964) * vulkan: add fwht support for Intel with shmem reduction * don't use N as workgroup size * disable subgroup shuffle on MoltenVK AMD * disable fwht shader on Intel Windows due to driver bug --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 13 ++++ ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp | 78 +++++++++++++++---- .../vulkan-shaders/vulkan-shaders-gen.cpp | 1 + 3 files changed, 76 insertions(+), 16 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index e7d04634b..df410368a 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -5084,6 +5084,14 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { } ++idx; } + } else if (device->driver_id != vk::DriverId::eIntelProprietaryWindows) { + // Disabled on Intel Windows due to a driver bug: https://github.com/ggml-org/llama.cpp/pull/23964#issuecomment-4598226147 + int idx = 0; + for (uint32_t n : {64, 128, 256, 512}) { + const uint32_t block_size = std::min(device->subgroup_size, n); + ggml_vk_create_pipeline(device, device->pipeline_fwht_f32[idx], "fwht_shmem_f32", fwht_shmem_f32_len, fwht_shmem_f32_data, "main", 2, sizeof(vk_op_fwht_push_constants), {1, 1, 1}, { block_size, n }, 1); + ++idx; + } } const uint32_t cumsum_elem_per_thread = (device->vendor_id == VK_VENDOR_ID_AMD || device->vendor_id == VK_VENDOR_ID_INTEL) ? 2 : 4; @@ -5630,6 +5638,11 @@ static vk_device ggml_vk_get_device(size_t idx) { #endif device->subgroup_shuffle = (vk11_props.subgroupSupportedStages & vk::ShaderStageFlagBits::eCompute) && (vk11_props.subgroupSupportedOperations & vk::SubgroupFeatureFlagBits::eShuffle); +#ifdef __APPLE__ + if (device->vendor_id == VK_VENDOR_ID_AMD) { + device->subgroup_shuffle = false; + } +#endif device->subgroup_clustered = (vk11_props.subgroupSupportedStages & vk::ShaderStageFlagBits::eCompute) && (vk11_props.subgroupSupportedOperations & vk::SubgroupFeatureFlagBits::eClustered); diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp b/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp index 72059d4af..a2069964a 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/fwht.comp @@ -1,14 +1,16 @@ #version 450 #extension GL_EXT_control_flow_attributes : require +#ifndef FWHT_SHMEM #extension GL_KHR_shader_subgroup_basic : enable #extension GL_KHR_shader_subgroup_shuffle : enable +#endif + +layout(constant_id = 0) const uint BLOCK_SIZE = 32; +layout(constant_id = 1) const uint N = 128; layout(local_size_x_id = 0, local_size_y = 4, local_size_z = 1) in; -layout(constant_id = 0) const uint WARP_SIZE = 32; -layout(constant_id = 1) const uint N = 128; - layout(push_constant) uniform parameter { uint n_rows; @@ -20,35 +22,72 @@ layout(push_constant) uniform parameter layout(binding = 0, std430) readonly buffer A { float data_a[]; }; layout(binding = 1, std430) writeonly buffer D { float data_d[]; }; -const uint EL_W = N / WARP_SIZE; +const uint EL_W = N / BLOCK_SIZE; + +#ifdef FWHT_SHMEM +shared float shmem[4 * N]; +#endif void main() { - const uint lane = gl_SubgroupInvocationID; - for (uint row = gl_WorkGroupID.x * gl_WorkGroupSize.y + gl_SubgroupID; - row < n_rows; - row += gl_NumWorkGroups.x * gl_WorkGroupSize.y) { +#ifdef FWHT_SHMEM + const uint tid = gl_LocalInvocationID.x; + const uint shmem_base = gl_LocalInvocationID.y * N; + const uint row_id = gl_LocalInvocationID.y; +#else + const uint tid = gl_SubgroupInvocationID; + const uint row_id = gl_SubgroupID; +#endif + + for (uint base_row = gl_WorkGroupID.x * gl_WorkGroupSize.y; + base_row < n_rows; + base_row += gl_NumWorkGroups.x * gl_WorkGroupSize.y) { + const uint row = base_row + row_id; const uint row_offset = row * N; +#ifndef FWHT_SHMEM + if (row >= n_rows) { + continue; + } +#endif + float reg[EL_W]; [[unroll]] for (uint i = 0; i < EL_W; ++i) { - reg[i] = data_a[src_offset + row_offset + i * WARP_SIZE + lane] * scale; + reg[i] = row < n_rows ? data_a[src_offset + row_offset + i * BLOCK_SIZE + tid] * scale : 0.0; } +#ifdef FWHT_SHMEM [[unroll]] - for (uint h = 1; h < WARP_SIZE; h <<= 1) { + for (uint h = 1; h < BLOCK_SIZE; h <<= 1) { + [[unroll]] + for (uint i = 0; i < EL_W; ++i) { + shmem[shmem_base + i * BLOCK_SIZE + tid] = reg[i]; + } + barrier(); + [[unroll]] + for (uint j = 0; j < EL_W; ++j) { + const float val = reg[j]; + const float other = shmem[shmem_base + j * BLOCK_SIZE + (tid ^ h)]; + reg[j] = (tid & h) == 0 ? val + other : other - val; + } + barrier(); + } +#else + [[unroll]] + for (uint h = 1; h < BLOCK_SIZE; h <<= 1) { [[unroll]] for (uint j = 0; j < EL_W; ++j) { const float val = reg[j]; const float val2 = subgroupShuffleXor(val, h); - reg[j] = (lane & h) == 0 ? val + val2 : val2 - val; + reg[j] = (tid & h) == 0 ? val + val2 : val2 - val; } } +#endif [[unroll]] - for (uint h = WARP_SIZE; h < N; h <<= 1) { - const uint step = h / WARP_SIZE; + for (uint h = BLOCK_SIZE; h < N; h <<= 1) { + const uint step = h / BLOCK_SIZE; [[unroll]] for (uint j = 0; j < EL_W; j += 2 * step) { [[unroll]] @@ -61,9 +100,16 @@ void main() { } } - [[unroll]] - for (uint i = 0; i < EL_W; ++i) { - data_d[dst_offset + row_offset + i * WARP_SIZE + lane] = reg[i]; +#ifdef FWHT_SHMEM + if (row < n_rows) { +#endif + [[unroll]] + for (uint i = 0; i < EL_W; ++i) { + data_d[dst_offset + row_offset + i * BLOCK_SIZE + tid] = reg[i]; + } +#ifdef FWHT_SHMEM } + barrier(); +#endif } } diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index de7dbec2c..d65cd12b2 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -957,6 +957,7 @@ void process_shaders() { string_to_spv("argmax_f32", "argmax.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "int"}})); string_to_spv("sum_rows_f32", "sum_rows.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}})); string_to_spv("fwht_f32", "fwht.comp", {}); + string_to_spv("fwht_shmem_f32", "fwht.comp", {{"FWHT_SHMEM", "1"}}); string_to_spv("count_equal_i32", "count_equal.comp", merge_maps(base_dict, {{"A_TYPE", "int"}, {"B_TYPE", "int"}, {"D_TYPE", "int"}})); string_to_spv("cumsum_f32", "cumsum.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}})); string_to_spv("cumsum_multipass1_f32", "cumsum_multipass1.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}})); From da87e9b612592ccd2a2e783eff631346459b9de0 Mon Sep 17 00:00:00 2001 From: Tarek Dakhran Date: Fri, 5 Jun 2026 21:31:56 +0200 Subject: [PATCH 4/8] common/chat : unify and fix LFM2/LFM2.5 tool parser (#24178) --- common/chat-peg-parser.cpp | 60 +++++++++++++--- common/chat-peg-parser.h | 7 +- common/chat.cpp | 141 +++++++------------------------------ tests/test-chat.cpp | 68 +++++++++++++++--- 4 files changed, 143 insertions(+), 133 deletions(-) diff --git a/common/chat-peg-parser.cpp b/common/chat-peg-parser.cpp index 12e747d1c..9bc5ac98b 100644 --- a/common/chat-peg-parser.cpp +++ b/common/chat-peg-parser.cpp @@ -87,6 +87,8 @@ static std::string normalize_quotes_to_json(const std::string & input) { bool in_single_quoted = false; bool in_double_quoted = false; + auto is_word_char = [](char ch) { return std::isalnum(static_cast(ch)) || ch == '_'; }; + for (size_t i = 0; i < input.size(); ++i) { char c = input[i]; @@ -151,6 +153,29 @@ static std::string normalize_quotes_to_json(const std::string & input) { in_single_quoted = true; result += '"'; } + } else if (!in_single_quoted && !in_double_quoted && (c == 'T' || c == 'F' || c == 'N') && + (i == 0 || !is_word_char(input[i - 1]))) { + // Python literals -> JSON; prefix match keeps streamed partials monotonic. + static constexpr std::pair literals[] = { + { "True", "true" }, { "False", "false" }, { "None", "null" }, + }; + size_t n = 0; + while (i + n < input.size() && is_word_char(input[i + n])) { + ++n; + } + std::string_view token(input.data() + i, n); + bool matched = false; + for (const auto & [py, js] : literals) { + if (py.substr(0, n) == token) { + result += js.substr(0, n); + i += n - 1; + matched = true; + break; + } + } + if (!matched) { + result += c; + } } else { result += c; } @@ -353,12 +378,8 @@ void common_chat_peg_mapper::map(const common_peg_ast_node & node) { } value_to_add += escape_json_string_inner(value_content); } else if (!value_content.empty()) { - // For potential containers, normalize Python-style single quotes to JSON double quotes - bool is_potential_container = value_content[0] == '[' || value_content[0] == '{'; - if (is_potential_container) { - value_content = normalize_container_value(value_content); - } - value_to_add += value_content; + // Pythonic scalars/containers -> JSON. + value_to_add += normalize_container_value(value_content); } args_target() += value_to_add; @@ -466,11 +487,34 @@ common_peg_parser common_chat_peg_builder::standard_constructed_tools( return force_tool_calls ? section : optional(section); } +// Like python_value(), but the leaf also accepts JSON-cased true/false/null, used by LFM2/LFM2.5 +common_peg_parser common_chat_peg_builder::python_or_json_value() { + return rule("python-or-json-value", [this]() { + auto ws = space(); + auto value = python_or_json_value(); + + auto member = sequence({ python_string(), ws, literal(":"), ws, value }); + auto members = sequence({ member, zero_or_more(sequence({ ws, literal(","), ws, member })) }); + auto dict = rule("python-or-json-dict", [&]() { + return sequence({ literal("{"), ws, choice({ literal("}"), sequence({ members, ws, literal("}") }) }), ws }); + }); + + auto elements = sequence({ value, zero_or_more(sequence({ literal(","), ws, value })) }); + auto array = rule("python-or-json-array", [&]() { + return sequence({ literal("["), ws, choice({ literal("]"), sequence({ elements, ws, literal("]") }) }), ws }); + }); + + return choice({ dict, array, python_string(), python_number(), + python_bool(), python_null(), json_bool(), json_null() }); + }); +} + // Python-style tool calls: name(arg1="value1", arg2=123) // Used only by LFM2 for now, so we don't merge it into autoparser common_peg_parser common_chat_peg_builder::python_style_tool_calls( const ordered_json & tools, - bool parallel_tool_calls) { + bool parallel_tool_calls, + bool allow_json_literals) { if (!tools.is_array() || tools.empty()) { return eps(); } @@ -504,7 +548,7 @@ common_peg_parser common_chat_peg_builder::python_style_tool_calls( if (is_string_type) { arg_value_parser = string_value_parser; } else { - arg_value_parser = tool_arg_value(python_value()); + arg_value_parser = tool_arg_value(allow_json_literals ? python_or_json_value() : python_value()); } // Full argument: name="value" or name=value diff --git a/common/chat-peg-parser.h b/common/chat-peg-parser.h index be92f17d9..a4643fbea 100644 --- a/common/chat-peg-parser.h +++ b/common/chat-peg-parser.h @@ -132,9 +132,13 @@ class common_chat_peg_builder : public common_peg_parser_builder { // Helper for Python-style function call format: name(arg1="value1", arg2=123) // Used by LFM2 and similar templates common_peg_parser python_style_tool_calls(const nlohmann::ordered_json & tools, - bool parallel_tool_calls); + bool parallel_tool_calls, + bool allow_json_literals); private: + // Python values plus JSON true/false/null. + common_peg_parser python_or_json_value(); + // Implementation helpers for standard_json_tools — one per JSON tool call layout mode common_peg_parser build_json_tools_function_is_key(const nlohmann::ordered_json & tools, const std::string & args_key, @@ -195,4 +199,3 @@ struct tagged_peg_parser { tagged_peg_parser build_tagged_peg_parser( const std::function & fn); - diff --git a/common/chat.cpp b/common/chat.cpp index ef151691c..b8f248dab 100644 --- a/common/chat.cpp +++ b/common/chat.cpp @@ -1608,42 +1608,40 @@ static common_chat_params common_chat_params_init_kimi_k2(const common_chat_temp return data; } -// LFM2 format: uses <|tool_list_start|>[...]<|tool_list_end|> in system prompt -// and <|tool_call_start|>[name(arg="val")]<|tool_call_end|> for tool calls. -// - Reasoning: {reasoning} (optional) -// - Content: text before a tool call (optional) -// - Tool calls: Python-style, e.g. [function_name(arg1="value1", arg2="value2")] -// Tool calls can appear multiple times (parallel tool calls supported) -static common_chat_params common_chat_params_init_lfm2(const common_chat_template & tmpl, - const autoparser::generation_params & inputs) { +// LFM2/LFM2.5 parser. Tool calls are almost Python-style and parallel-capable +// (except dotted names and JSON literals true/false/null). +// Always wrapped in <|tool_call_start|>[name(args)]<|tool_call_end|> with optional reasoning. +// tool_list_tokens preserves LFM2 system tool-list markers. +static common_chat_params common_chat_params_init_lfm2(const common_chat_template & tmpl, + const autoparser::generation_params & inputs, + bool tool_list_tokens) { common_chat_params data; + const std::string TOOL_CALL_START = "<|tool_call_start|>"; + const std::string TOOL_CALL_END = "<|tool_call_end|>"; + const std::string TOOL_LIST_START = "<|tool_list_start|>"; + const std::string TOOL_LIST_END = "<|tool_list_end|>"; + const std::string THINK_START = ""; + const std::string THINK_END = ""; + const std::string GEN_PROMPT = "<|im_start|>assistant\n"; + data.prompt = common_chat_template_direct_apply_impl(tmpl, inputs); data.generation_prompt = common_chat_template_generation_prompt_impl(tmpl, inputs); data.format = COMMON_CHAT_FORMAT_PEG_NATIVE; data.supports_thinking = true; - data.preserved_tokens = { - "<|tool_list_start|>", - "<|tool_list_end|>", - "<|tool_call_start|>", - "<|tool_call_end|>", - "", - "", - }; + data.preserved_tokens = { TOOL_CALL_START, TOOL_CALL_END, THINK_START, THINK_END }; + if (tool_list_tokens) { + data.preserved_tokens.push_back(TOOL_LIST_START); + data.preserved_tokens.push_back(TOOL_LIST_END); + } + + data.thinking_start_tag = THINK_START; + data.thinking_end_tag = THINK_END; auto has_tools = inputs.tools.is_array() && !inputs.tools.empty(); auto extract_reasoning = inputs.reasoning_format != COMMON_REASONING_FORMAT_NONE; auto include_grammar = has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE; - const std::string TOOL_CALL_START = "<|tool_call_start|>"; - const std::string TOOL_CALL_END = "<|tool_call_end|>"; - const std::string THINK_START = ""; - const std::string THINK_END = ""; - const std::string GEN_PROMPT = "<|im_start|>assistant\n"; - - data.thinking_start_tag = THINK_START; - data.thinking_end_tag = THINK_END; - if (inputs.has_continuation()) { const auto & msg = inputs.continue_msg; @@ -1670,7 +1668,7 @@ static common_chat_params common_chat_params_init_lfm2(const common_chat_templat auto tool_calls = p.rule("tool-calls", p.trigger_rule("tool-call", p.literal(TOOL_CALL_START) + - p.python_style_tool_calls(inputs.tools, inputs.parallel_tool_calls) + + p.python_style_tool_calls(inputs.tools, inputs.parallel_tool_calls, /* allow_json_literals = */ true) + p.literal(TOOL_CALL_END) ) ); @@ -1697,93 +1695,6 @@ static common_chat_params common_chat_params_init_lfm2(const common_chat_templat { COMMON_GRAMMAR_TRIGGER_TYPE_WORD, TOOL_CALL_START } }; } - return data; -} - -// LFM2.5 format: uses plain "List of tools: [...]" in system prompt, no wrapper tokens. -// Tool calls are bare [name(arg="val")], though model may optionally emit <|tool_call_start|>. -// - Reasoning: {reasoning} (optional) -// - Content: text before a tool call (optional) -// - Tool calls: Python-style, e.g. [function_name(arg1="value1", arg2="value2")] -// Tool calls can appear multiple times (parallel tool calls supported) -static common_chat_params common_chat_params_init_lfm2_5(const common_chat_template & tmpl, - const autoparser::generation_params & inputs) { - common_chat_params data; - - data.prompt = common_chat_template_direct_apply_impl(tmpl, inputs); - data.generation_prompt = common_chat_template_generation_prompt_impl(tmpl, inputs); - data.format = COMMON_CHAT_FORMAT_PEG_NATIVE; - data.supports_thinking = true; - data.preserved_tokens = { - "<|tool_call_start|>", - "<|tool_call_end|>", - "", - "", - }; - - auto has_tools = inputs.tools.is_array() && !inputs.tools.empty(); - auto extract_reasoning = inputs.reasoning_format != COMMON_REASONING_FORMAT_NONE; - auto include_grammar = has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE; - - const std::string THINK_START = ""; - const std::string THINK_END = ""; - const std::string GEN_PROMPT = "<|im_start|>assistant\n"; - - data.thinking_start_tag = THINK_START; - data.thinking_end_tag = THINK_END; - - if (inputs.has_continuation()) { - const auto & msg = inputs.continue_msg; - - data.generation_prompt = GEN_PROMPT + THINK_START + msg.reasoning_content; - if (inputs.continue_final_message == COMMON_CHAT_CONTINUATION_CONTENT) { - data.generation_prompt += THINK_END + msg.render_content(); - } - - data.prompt += data.generation_prompt; - } - - auto parser = build_chat_peg_parser([&](common_chat_peg_builder & p) { - auto generation_prompt = p.literal(GEN_PROMPT); - auto end = p.end(); - - auto reasoning = p.eps(); - if (extract_reasoning && inputs.enable_thinking) { - reasoning = p.optional(THINK_START + p.reasoning(p.until(THINK_END)) + THINK_END); - } - - if (!has_tools || inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_NONE) { - return generation_prompt + reasoning + p.content(p.rest()) + end; - } - - auto tool_calls = p.rule("tool-calls", - p.trigger_rule("tool-call", - p.python_style_tool_calls(inputs.tools, inputs.parallel_tool_calls) - ) - ); - - auto content = p.content(p.until_one_of({"<|tool_call_start|>", "["})); - auto maybe_start = p.optional(p.literal("<|tool_call_start|>")); - return generation_prompt + reasoning + content + maybe_start + tool_calls + end; - }); - - data.parser = parser.save(); - - if (include_grammar) { - data.grammar_lazy = inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_AUTO; - data.grammar = build_grammar([&](const common_grammar_builder & builder) { - foreach_function(inputs.tools, [&](const json & tool) { - const auto & function = tool.at("function"); - auto schema = function.at("parameters"); - builder.resolve_refs(schema); - }); - parser.build_grammar(builder, data.grammar_lazy); - }); - foreach_function(inputs.tools, [&](const json & tool) { - const std::string name = tool.at("function").at("name"); - data.grammar_triggers.push_back({ COMMON_GRAMMAR_TRIGGER_TYPE_WORD, "[" + name + "(" }); - }); - } return data; } @@ -2298,14 +2209,14 @@ std::optional common_chat_try_specialized_template( if (is_lfm2_template(src)) { LOG_DBG("Using specialized template: LFM2\n"); - return common_chat_params_init_lfm2(tmpl, params); + return common_chat_params_init_lfm2(tmpl, params, /* tool_list_tokens = */ true); } // LFM2.5 format detection: template uses plain "List of tools: [...]" with no special tokens if (src.find("List of tools: [") != std::string::npos && src.find("<|tool_list_start|>") == std::string::npos) { LOG_DBG("Using specialized template: LFM2.5\n"); - return common_chat_params_init_lfm2_5(tmpl, params); + return common_chat_params_init_lfm2(tmpl, params, /* tool_list_tokens = */ false); } // GigaChatV3 format detection diff --git a/tests/test-chat.cpp b/tests/test-chat.cpp index 30ea2c072..3107045b4 100644 --- a/tests/test-chat.cpp +++ b/tests/test-chat.cpp @@ -684,6 +684,20 @@ static common_chat_tool config_tool{ })", }; +static common_chat_tool calendar_create_event_tool{ + /* .name = */ "Calendar.create_event", + /* .description = */ "Create a calendar event", + /* .parameters = */ R"({ + "type": "object", + "properties": { + "title": { "type": "string" }, + "participants": { "type": "array", "items": { "type": "string" } }, + "metadata": { "type": "object" } + }, + "required": ["title", "participants", "metadata"] + })", +}; + static common_chat_tool imaginary_number_tool{ /* .name = */ "imaginary_number", /* .description = */ "Imaginary number converter", @@ -4130,7 +4144,7 @@ static void test_template_output_peg_parsers(bool detailed_debug) { .run(); } - // LFM2.5 tests - uses plain "List of tools: [...]" and bare [name(args)] without wrapper tokens + // LFM2.5 tests - format <|tool_call_start|>[name(args)]<|tool_call_end|> { auto tst = peg_tester("models/templates/LFM2.5-Instruct.jinja", detailed_debug); @@ -4138,19 +4152,57 @@ static void test_template_output_peg_parsers(bool detailed_debug) { tst.test("Hello, world!\nWhat's up?").expect(message_assist).run(); // Single tool call without reasoning - tst.test("[special_function(arg1=1)]") + tst.test("<|tool_call_start|>[special_function(arg1=1)]<|tool_call_end|>") .tools({ special_function_tool }) .expect(message_assist_call) .run(); // Tool call with string argument - tst.test("[get_time(city=\"XYZCITY\")]") + tst.test("<|tool_call_start|>[get_time(city=\"XYZCITY\")]<|tool_call_end|>") .tools({ get_time_tool }) .expect(message_with_tool_calls("get_time", "{\"city\":\"XYZCITY\"}")) .run(); + // Python literals become JSON. + tst.test("<|tool_call_start|>[toggle(enabled=True)]<|tool_call_end|>") + .tools({ toggle_tool }) + .expect(message_with_tool_calls("toggle", R"({"enabled": true})")) + .run(); + + tst.test("<|tool_call_start|>[set_nullable(value=None)]<|tool_call_end|>") + .tools({ nullable_tool }) + .expect(message_with_tool_calls("set_nullable", R"({"value": null})")) + .run(); + + // Nested Python literal. + tst.test("<|tool_call_start|>[set_config(config={\"enabled\": True, \"count\": 3})]<|tool_call_end|>") + .tools({ config_tool }) + .expect(message_with_tool_calls("set_config", R"({"config": {"enabled": true, "count": 3}})")) + .run(); + + // JSON literals are accepted too. + tst.test("<|tool_call_start|>[set_config(config={\"enabled\": true, \"note\": null})]<|tool_call_end|>") + .tools({ config_tool }) + .expect(message_with_tool_calls("set_config", R"({"config": {"enabled": true, "note": null}})")) + .run(); + + // Dotted function name with structured args. + tst.test("<|tool_call_start|>[Calendar.create_event(title=\"demo\", participants=[\"Alice\", \"Bob\"], " + "metadata={\"priority\": \"high\", \"reminder\": true})]<|tool_call_end|>") + .tools({ calendar_create_event_tool }) + .expect(message_with_tool_calls( + "Calendar.create_event", + R"({"title": "demo", "participants": ["Alice", "Bob"], "metadata": {"priority": "high", "reminder": true}})")) + .run(); + + // Markdown links stay content. + tst.test("Use this format: [link text](url). Example: [Wikipedia](https://www.wikipedia.org).") + .tools({ get_time_tool }) + .expect(simple_assist_msg("Use this format: [link text](url). Example: [Wikipedia](https://www.wikipedia.org).")) + .run(); + // Tool call with reasoning (enable_thinking=true) - tst.test("I'm\nthinking[special_function(arg1=1)]") + tst.test("I'm\nthinking<|tool_call_start|>[special_function(arg1=1)]<|tool_call_end|>") .enable_thinking(true) .reasoning_format(COMMON_REASONING_FORMAT_AUTO) .tools({ special_function_tool }) @@ -4158,7 +4210,7 @@ static void test_template_output_peg_parsers(bool detailed_debug) { .run(); // Multiple tool calls (parallel) - tst.test("[special_function(arg1=1), special_function_with_opt(arg1=1, arg2=2)]") + tst.test("<|tool_call_start|>[special_function(arg1=1), special_function_with_opt(arg1=1, arg2=2)]<|tool_call_end|>") .parallel_tool_calls(true) .tools({ special_function_tool, special_function_tool_with_optional_param @@ -4170,7 +4222,7 @@ static void test_template_output_peg_parsers(bool detailed_debug) { .run(); // Tool call with content before tool call - tst.test("Let me check the time.[get_time(city=\"Paris\")]") + tst.test("Let me check the time.<|tool_call_start|>[get_time(city=\"Paris\")]<|tool_call_end|>") .tools({ get_time_tool }) .expect(message_with_reasoning_content_and_multiple_tool_calls( "", "Let me check the time.", { { "get_time", "{\"city\":\"Paris\"}" } } @@ -4178,14 +4230,14 @@ static void test_template_output_peg_parsers(bool detailed_debug) { .run(); // Partial tool call (streaming) - tst.test("[special_function(arg1=") + tst.test("<|tool_call_start|>[special_function(arg1=") .tools({ special_function_tool }) .is_partial(true) .expect(simple_assist_msg("", "", "special_function", "{\"arg1\": ")) .run(); // Tool call with empty arguments - tst.test("[empty_args()]") + tst.test("<|tool_call_start|>[empty_args()]<|tool_call_end|>") .tools({ empty_args_tool }) .expect(simple_assist_msg("", "", "empty_args", "{}")) .run(); From 308f61c31f083251ce8150f10b9ef97679b500b5 Mon Sep 17 00:00:00 2001 From: lhez Date: Fri, 5 Jun 2026 13:45:25 -0700 Subject: [PATCH 5/8] opencl: improve get_rows, cpy, concat and q6_k flat gemv (#24160) * opencl: allow multiple workgroups for large rows * opencl: improve small cpy * opencl: packed concat for small input * opencl: tweak flat q6_K gemv, increase N_DST and remap threads --- ggml/src/ggml-opencl/ggml-opencl.cpp | 71 +++++++++--- ggml/src/ggml-opencl/kernels/concat.cl | 67 +++++++++++ ggml/src/ggml-opencl/kernels/cpy.cl | 59 ++++++++++ ggml/src/ggml-opencl/kernels/get_rows.cl | 24 ++-- .../kernels/mul_mv_q6_k_f32_flat.cl | 106 ++++++++---------- 5 files changed, 244 insertions(+), 83 deletions(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index c411e4aea..2a41215fd 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -558,7 +558,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_set_rows_f32_i64, kernel_set_rows_f32_i32, kernel_set_rows_f16_i64, kernel_set_rows_f16_i32; cl_kernel kernel_rope_norm_f32, kernel_rope_norm_f16, kernel_rope_neox_f32, kernel_rope_neox_f16; cl_kernel kernel_rope_multi_f32, kernel_rope_multi_f16, kernel_rope_vision_f32, kernel_rope_vision_f16; - cl_kernel kernel_cpy_f16_f16, kernel_cpy_f16_f32, kernel_cpy_f32_f16, kernel_cpy_f32_f32, kernel_cpy_i32_i32; + cl_kernel kernel_cpy_f16_f16, kernel_cpy_f16_f32, kernel_cpy_f32_f16, kernel_cpy_f32_f32, kernel_cpy_f32_f32_pack, kernel_cpy_i32_i32; cl_kernel kernel_mul_mat_f32_f32; cl_kernel kernel_mul_mat_f16_f16; cl_kernel kernel_mul_mat_f16_f32_1row; @@ -639,7 +639,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_softplus_f16, kernel_softplus_f16_4, kernel_softplus_f16_nc; cl_kernel kernel_upscale; cl_kernel kernel_upscale_bilinear; - cl_kernel kernel_concat_f32; + cl_kernel kernel_concat_f32, kernel_concat_f32_pack; cl_kernel kernel_conv_2d_f16; cl_kernel kernel_conv_2d_f32; cl_kernel kernel_conv_2d_f16_f32; @@ -1121,6 +1121,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { CL_CHECK((backend_ctx->kernel_cpy_f16_f32 = clCreateKernel(prog, "kernel_cpy_f16_f32", &err), err)); CL_CHECK((backend_ctx->kernel_cpy_f32_f16 = clCreateKernel(prog, "kernel_cpy_f32_f16", &err), err)); CL_CHECK((backend_ctx->kernel_cpy_f32_f32 = clCreateKernel(prog, "kernel_cpy_f32_f32", &err), err)); + CL_CHECK((backend_ctx->kernel_cpy_f32_f32_pack = clCreateKernel(prog, "kernel_cpy_f32_f32_pack", &err), err)); CL_CHECK((backend_ctx->kernel_cpy_i32_i32 = clCreateKernel(prog, "kernel_cpy_i32_i32", &err), err)); GGML_LOG_CONT("."); } @@ -2615,6 +2616,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { cl_program prog = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); CL_CHECK((backend_ctx->kernel_concat_f32 = clCreateKernel(prog, "kernel_concat_f32", &err), err)); + CL_CHECK((backend_ctx->kernel_concat_f32_pack = clCreateKernel(prog, "kernel_concat_f32_pack", &err), err)); CL_CHECK(clReleaseProgram(prog)); GGML_LOG_CONT("."); } @@ -8552,7 +8554,14 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c nth *= 2; } - size_t global_work_size[] = {(size_t)ne10*nth, (size_t)ne11, (size_t)ne12}; + int nchunks = 1; + if (src0->type == GGML_TYPE_F32) { + const int chunk_target = nth * 4; + nchunks = (ne00 + chunk_target - 1) / chunk_target; + nchunks = MAX(1, MIN(nchunks, 64)); + } + + size_t global_work_size[] = {(size_t)ne10*nth*nchunks, (size_t)ne11, (size_t)ne12}; size_t local_work_size[] = {(size_t)nth, 1, 1}; backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); @@ -11128,7 +11137,9 @@ static void ggml_cl_concat(ggml_backend_t backend, const ggml_tensor * src0, con int nth = MIN(64, ne0); - cl_kernel kernel = backend_ctx->kernel_concat_f32; + const bool concat_pack = (dim == 0 && ne0 < 32); + cl_kernel kernel = concat_pack ? backend_ctx->kernel_concat_f32_pack + : backend_ctx->kernel_concat_f32; CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device)); CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0)); @@ -11155,10 +11166,28 @@ static void ggml_cl_concat(ggml_backend_t backend, const ggml_tensor * src0, con CL_CHECK(clSetKernelArg(kernel, 22, sizeof(cl_ulong), &nb3)); CL_CHECK(clSetKernelArg(kernel, 23, sizeof(cl_int), &dim)); - size_t global_work_size[] = {(size_t)ne1*nth, (size_t)ne2, (size_t)ne3}; - size_t local_work_size[] = {(size_t)nth, 1, 1}; + if (concat_pack) { + // packed kernel needs the dst dims to unflatten its 1-D row index. + CL_CHECK(clSetKernelArg(kernel, 24, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 25, sizeof(int), &ne2)); + CL_CHECK(clSetKernelArg(kernel, 26, sizeof(int), &ne3)); - backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + const int maxwg = (int)backend_ctx->get_kernel_workgroup_size(kernel); + const int base = MIN(64, maxwg); + const int tpr = MIN(ne0, base); // threads per row + const int rpw = MAX(1, base / tpr); // rows per workgroup + const int lsz = tpr * rpw; + const int nrows = ne1*ne2*ne3; + const int nwg = (nrows + rpw - 1) / rpw; + size_t global_work_size[] = {(size_t)nwg*lsz, 1, 1}; + size_t local_work_size[] = {(size_t)lsz, 1, 1}; + backend_ctx->enqueue_ndrange_kernel(kernel, 1, global_work_size, local_work_size, dst); + } else { + size_t global_work_size[] = {(size_t)ne1*nth, (size_t)ne2, (size_t)ne3}; + size_t local_work_size[] = {(size_t)nth, 1, 1}; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + } } static void ggml_cl_timestep_embedding(ggml_backend_t backend, const ggml_tensor * src0, ggml_tensor * dst) { @@ -14536,7 +14565,7 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co } else if (backend_ctx->gpu_family == ADRENO) { nth0 = 64; nth1 = 2; - ndst = 4; + ndst = 16; } else { GGML_ASSERT(false && "TODO: Unknown GPU"); } @@ -16633,7 +16662,8 @@ static void ggml_cl_cpy(ggml_backend_t backend, const ggml_tensor * src0, const kernel = backend_ctx->kernel_cpy_f32_f16; break; case GGML_TYPE_F32: - kernel = backend_ctx->kernel_cpy_f32_f32; + kernel = ne00 < 32 ? backend_ctx->kernel_cpy_f32_f32_pack + : backend_ctx->kernel_cpy_f32_f32; break; default: GGML_ASSERT(false && "not implemented"); @@ -16685,12 +16715,27 @@ static void ggml_cl_cpy(ggml_backend_t backend, const ggml_tensor * src0, const CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_ulong), &nb12)); CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_ulong), &nb13)); - const int nth = MIN(64, ne00); + if (kernel == backend_ctx->kernel_cpy_f32_f32_pack) { + const int maxwg = (int)backend_ctx->get_kernel_workgroup_size(kernel); + const int base = MIN(64, maxwg); + const int tpr = MIN(ne00, base); // threads per row + const int rpw = MAX(1, base / tpr); // rows per workgroup + const int lsz = tpr * rpw; // <= base <= maxwg + const int nrows = ne01*ne02*ne03; + const int nwg = (nrows + rpw - 1) / rpw; - size_t global_work_size[] = {(size_t)ne01*nth, (size_t)ne02, (size_t)ne03}; - size_t local_work_size[] = {(size_t)nth, 1, 1}; + size_t global_work_size[] = {(size_t)nwg*lsz, 1, 1}; + size_t local_work_size[] = {(size_t)lsz, 1, 1}; - backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, src1); + backend_ctx->enqueue_ndrange_kernel(kernel, 1, global_work_size, local_work_size, src1); + } else { + const int nth = MIN(64, ne00); + + size_t global_work_size[] = {(size_t)ne01*nth, (size_t)ne02, (size_t)ne03}; + size_t local_work_size[] = {(size_t)nth, 1, 1}; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, src1); + } } static void ggml_cl_dup(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { diff --git a/ggml/src/ggml-opencl/kernels/concat.cl b/ggml/src/ggml-opencl/kernels/concat.cl index 0c1b3d785..2fbd7851d 100644 --- a/ggml/src/ggml-opencl/kernels/concat.cl +++ b/ggml/src/ggml-opencl/kernels/concat.cl @@ -49,3 +49,70 @@ kernel void kernel_concat_f32( *y = *x; } } + +kernel void kernel_concat_f32_pack( + global const char * src0, + ulong offset0, + global const char * src1, + ulong offset1, + global char * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne03, + ulong nb00, + ulong nb01, + ulong nb02, + ulong nb03, + ulong nb10, + ulong nb11, + ulong nb12, + ulong nb13, + int ne0, + ulong nb0, + ulong nb1, + ulong nb2, + ulong nb3, + int dim, + int ne1, + int ne2, + int ne3 +) { + src0 = src0 + offset0; + src1 = src1 + offset1; + dst = dst + offsetd; + + int lsz = get_local_size(0); + int tpr = min(ne0, lsz); // threads per row + int rpw = lsz / tpr; // rows per workgroup + int lid = get_local_id(0); + int row = get_group_id(0)*rpw + lid / tpr; + int lane = lid - (lid / tpr) * tpr; + + int nrows = ne1*ne2*ne3; + if (row >= nrows) { + return; + } + + int i1 = row % ne1; + int t = row / ne1; + int i2 = t % ne2; + int i3 = t / ne2; + + int o[4] = {0, 0, 0, 0}; + o[dim] = dim == 0 ? ne00 : (dim == 1 ? ne01 : (dim == 2 ? ne02 : ne03)); + + for (int i0 = lane; i0 < ne0; i0 += tpr) { + global const float * x; + if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) { + x = (global const float *)(src0 + (i3 )*nb03 + (i2 )*nb02 + (i1 )*nb01 + (i0 )*nb00); + } else { + x = (global const float *)(src1 + (i3 - o[3])*nb13 + (i2 - o[2])*nb12 + (i1 - o[1])*nb11 + (i0 - o[0])*nb10); + } + + global float * y = (global float *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); + + *y = *x; + } +} diff --git a/ggml/src/ggml-opencl/kernels/cpy.cl b/ggml/src/ggml-opencl/kernels/cpy.cl index 820aa538a..adbd2e766 100644 --- a/ggml/src/ggml-opencl/kernels/cpy.cl +++ b/ggml/src/ggml-opencl/kernels/cpy.cl @@ -183,6 +183,65 @@ kernel void kernel_cpy_f32_f32( } } +kernel void kernel_cpy_f32_f32_pack( + global float * src0, + ulong offset0, + global float * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne03, + ulong nb00, + ulong nb01, + ulong nb02, + ulong nb03, + int ne0, + int ne1, + int ne2, + int ne3, + ulong nb0, + ulong nb1, + ulong nb2, + ulong nb3 +) { + src0 = (global float*)((global char*)src0 + offset0); + dst = (global float*)((global char*)dst + offsetd); + + int lsz = get_local_size(0); + int tpr = min(ne00, lsz); // threads per row + int rpw = lsz / tpr; // rows per workgroup + int lid = get_local_id(0); + int row = get_group_id(0)*rpw + lid / tpr; + int lane = lid - (lid / tpr) * tpr; + + int nrows = ne01*ne02*ne03; + if (row >= nrows) { + return; + } + + int i01 = row % ne01; + int t = row / ne01; + int i02 = t % ne02; + int i03 = t / ne02; + + // linear index of the first element of this row, unflattened over dst dims + long n = (long)row * ne00; + int i3 = (int)(n / ((long)ne2*ne1*ne0)); + long rm = n - (long)i3*ne2*ne1*ne0; + int i2 = (int)(rm / ((long)ne1*ne0)); + rm -= (long)i2*ne1*ne0; + int i1 = (int)(rm / ne0); + int i0 = (int)(rm - (long)i1*ne0); + + global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); + + for (int i00 = lane; i00 < ne00; i00 += tpr) { + global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00); + dst_data[i00] = src[0]; + } +} + kernel void kernel_cpy_i32_i32( global int * src0, ulong offset0, diff --git a/ggml/src/ggml-opencl/kernels/get_rows.cl b/ggml/src/ggml-opencl/kernels/get_rows.cl index c2962edc9..9ae4fff09 100644 --- a/ggml/src/ggml-opencl/kernels/get_rows.cl +++ b/ggml/src/ggml-opencl/kernels/get_rows.cl @@ -82,21 +82,27 @@ kernel void kernel_get_rows_f32( src1 = (global int*)((global char*)src1 + offset1); dst = (global float*)((global char*)dst + offsetd); - int i10 = get_group_id(0); - int i11 = get_group_id(1); - int i12 = get_group_id(2); + int nchunks = get_num_groups(0) / ne10; + int g = get_group_id(0); + int i10 = g / nchunks; + int chunk = g - i10 * nchunks; + int i11 = get_group_id(1); + int i12 = get_group_id(2); int r = ((global int *) ((global char *) src1 + i12*nb12 + i11*nb11 + i10*nb10))[0]; int i02 = i11; int i03 = i12; - for (int ind = get_local_id(0); ind < ne00; ind += get_local_size(0)) { - if (ind >= ne00) { - return; - } - ((global float *) ((global char *) dst + i12*nb3 + i11*nb2 + i10*nb1))[ind] = - ((global float *) ((global char *) src0 + r*nb01 + i02*nb02 + i03*nb03))[ind]; + global float * dst_row = (global float *) ((global char *) dst + i12*nb3 + i11*nb2 + i10*nb1); + global float * src_row = (global float *) ((global char *) src0 + r*nb01 + i02*nb02 + i03*nb03); + + int span = (ne00 + nchunks - 1) / nchunks; + int start = chunk * span; + int end = min(start + span, ne00); + + for (int ind = start + get_local_id(0); ind < end; ind += get_local_size(0)) { + dst_row[ind] = src_row[ind]; } } diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q6_k_f32_flat.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q6_k_f32_flat.cl index 86fe09c6d..57b90c05a 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mv_q6_k_f32_flat.cl +++ b/ggml/src/ggml-opencl/kernels/mul_mv_q6_k_f32_flat.cl @@ -33,13 +33,15 @@ inline float block_q_6_K_dot_y_flat( global uchar * blk_qh, global char * blk_scales, global half * blk_d, - global float * yy, int ib, int ip, int is, - int l0 + int l0, + float4 y0, + float4 y1, + float4 y2, + float4 y3 ) { - int y_offset = 128*ip + l0; int q_offset_l = 64*ip + l0; int q_offset_h = 32*ip + l0; @@ -48,36 +50,28 @@ inline float block_q_6_K_dot_y_flat( global uchar * qh = blk_qh + ib*64 + q_offset_h; global char * sc = blk_scales + ib*16 + is; - global float * y = yy + ib * QK_K + y_offset; - float dall = blk_d[ib]; - float sumf = 0; - float4 sums = {0.f, 0.f, 0.f, 0.f}; + // Vectorized loads: 3 uchar4 weight loads instead of 12 scalar byte reads. + // q_offset_l/h are 4-aligned, so these are aligned vector loads. + uchar4 q1v = vload4(0, q1); + uchar4 q2v = vload4(0, q2); + uchar4 qhv = vload4(0, qh); - sums.s0 += y[0+ 0] * ((float)((q1[0] & 0xF) | ((qh[0] & Q6_K_MASK1) << 4)) - 32.f); - sums.s1 += y[0+32] * ((float)((q2[0] & 0xF) | ((qh[0] & Q6_K_MASK2) << 2)) - 32.f); - sums.s2 += y[0+64] * ((float)((q1[0] >> 4) | ((qh[0] & Q6_K_MASK3) << 0)) - 32.f); - sums.s3 += y[0+96] * ((float)((q2[0] >> 4) | ((qh[0] & Q6_K_MASK4) >> 2)) - 32.f); + int4 q1i = convert_int4(q1v); + int4 q2i = convert_int4(q2v); + int4 qhi = convert_int4(qhv); - sums.s0 += y[1+ 0] * ((float)((q1[1] & 0xF) | ((qh[1] & Q6_K_MASK1) << 4)) - 32.f); - sums.s1 += y[1+32] * ((float)((q2[1] & 0xF) | ((qh[1] & Q6_K_MASK2) << 2)) - 32.f); - sums.s2 += y[1+64] * ((float)((q1[1] >> 4) | ((qh[1] & Q6_K_MASK3) << 0)) - 32.f); - sums.s3 += y[1+96] * ((float)((q2[1] >> 4) | ((qh[1] & Q6_K_MASK4) >> 2)) - 32.f); + // Reconstruct the four 6-bit weight groups (low/high nibble of ql OR'd with the + // matching 2-bit plane of qh), same arithmetic as the scalar version, then dot() + // against the cached activation lanes. + float4 w0 = convert_float4((q1i & 0xF) | ((qhi & Q6_K_MASK1) << 4)) - 32.f; + float4 w1 = convert_float4((q2i & 0xF) | ((qhi & Q6_K_MASK2) << 2)) - 32.f; + float4 w2 = convert_float4((q1i >> 4) | ((qhi & Q6_K_MASK3) )) - 32.f; + float4 w3 = convert_float4((q2i >> 4) | ((qhi & Q6_K_MASK4) >> 2)) - 32.f; - sums.s0 += y[2+ 0] * ((float)((q1[2] & 0xF) | ((qh[2] & Q6_K_MASK1) << 4)) - 32.f); - sums.s1 += y[2+32] * ((float)((q2[2] & 0xF) | ((qh[2] & Q6_K_MASK2) << 2)) - 32.f); - sums.s2 += y[2+64] * ((float)((q1[2] >> 4) | ((qh[2] & Q6_K_MASK3) << 0)) - 32.f); - sums.s3 += y[2+96] * ((float)((q2[2] >> 4) | ((qh[2] & Q6_K_MASK4) >> 2)) - 32.f); - - sums.s0 += y[3+ 0] * ((float)((q1[3] & 0xF) | ((qh[3] & Q6_K_MASK1) << 4)) - 32.f); - sums.s1 += y[3+32] * ((float)((q2[3] & 0xF) | ((qh[3] & Q6_K_MASK2) << 2)) - 32.f); - sums.s2 += y[3+64] * ((float)((q1[3] >> 4) | ((qh[3] & Q6_K_MASK3) << 0)) - 32.f); - sums.s3 += y[3+96] * ((float)((q2[3] >> 4) | ((qh[3] & Q6_K_MASK4) >> 2)) - 32.f); - - sumf += dall * (sums.s0 * sc[0] + sums.s1 * sc[2] + sums.s2 * sc[4] + sums.s3 * sc[6]); - - return sumf; + return dall * (dot(y0, w0) * sc[0] + dot(y1, w1) * sc[2] + + dot(y2, w2) * sc[4] + dot(y3, w3) * sc[6]); } #undef N_DST @@ -89,7 +83,7 @@ inline float block_q_6_K_dot_y_flat( #define N_SIMDGROUP 2 #define N_SIMDWIDTH 16 #elif defined (ADRENO_GPU) -#define N_DST 4 +#define N_DST 16 #define N_SIMDGROUP 2 #define N_SIMDWIDTH 64 #endif @@ -146,49 +140,39 @@ kernel void kernel_mul_mv_q6_K_f32_flat( global half * blk_d = (global half *) src0_d + offset_src0_d; global float * yy = (global float *) src1 + r1*ne10 + im*ne00*ne1; - int tid = get_sub_group_local_id()/BLOCK_STRIDE; // first block_stride groups have tid=0 - int ix = get_sub_group_local_id()%BLOCK_STRIDE; // first block is 0..block_stride-1 + int tid = get_sub_group_local_id()%(N_SIMDWIDTH/BLOCK_STRIDE); // within-super-block part, 0..15 + int ix = get_sub_group_local_id()/(N_SIMDWIDTH/BLOCK_STRIDE); // super-block selector, 0..BLOCK_STRIDE-1 int ip = tid/8; // first or second half of (super) block (0 or 1) int il = tid%8; // each half has 8 parts, one per scale int n = 4; // 4 scales at a time (and 4 sums) int l0 = n*il; // offset into half-block, 0..28 int is = 8*ip + l0/16; // 0, 1, 8, 9 - float4 sumf = 0; + float sumf[N_DST]; + for (int row = 0; row < N_DST; row++) { + sumf[row] = 0.f; + } for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) { - if (first_row + 0 < ne01) { - sumf.s0 += block_q_6_K_dot_y_flat(blk_ql + 0*nb*128, blk_qh + 0*nb*64, blk_scales + 0*nb*16, blk_d + 0*nb, yy, ib, ip, is, l0); - } - if (first_row + 1 < ne01) { - sumf.s1 += block_q_6_K_dot_y_flat(blk_ql + 1*nb*128, blk_qh + 1*nb*64, blk_scales + 1*nb*16, blk_d + 1*nb, yy, ib, ip, is, l0); - } - if (first_row + 2 < ne01) { - sumf.s2 += block_q_6_K_dot_y_flat(blk_ql + 2*nb*128, blk_qh + 2*nb*64, blk_scales + 2*nb*16, blk_d + 2*nb, yy, ib, ip, is, l0); - } - if (first_row + 3 < ne01) { - sumf.s3 += block_q_6_K_dot_y_flat(blk_ql + 3*nb*128, blk_qh + 3*nb*64, blk_scales + 3*nb*16, blk_d + 3*nb, yy, ib, ip, is, l0); + global float * y = yy + ib * QK_K + 128*ip + l0; + float4 y0 = vload4(0, y + 0); + float4 y1 = vload4(0, y + 32); + float4 y2 = vload4(0, y + 64); + float4 y3 = vload4(0, y + 96); + + for (int row = 0; row < N_DST; row++) { + if (first_row + row < ne01) { + sumf[row] += block_q_6_K_dot_y_flat( + blk_ql + row*nb*128, blk_qh + row*nb*64, blk_scales + row*nb*16, blk_d + row*nb, + ib, ip, is, l0, y0, y1, y2, y3); + } } } - float4 tot = (float4)( - sub_group_reduce_add(sumf.s0), - sub_group_reduce_add(sumf.s1), - sub_group_reduce_add(sumf.s2), - sub_group_reduce_add(sumf.s3) - ); - if (get_sub_group_local_id() == 0) { - if (first_row + 0 < ne01) { - dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0; - } - if (first_row + 1 < ne01) { - dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1; - } - if (first_row + 2 < ne01) { - dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2; - } - if (first_row + 3 < ne01) { - dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3; + for (int row = 0; row < N_DST; row++) { + float tot = sub_group_reduce_add(sumf[row]); + if (get_sub_group_local_id() == 0 && first_row + row < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + row] = tot; } } } From 603300b008e01934c242c541259449ad95e91a43 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?= Date: Sat, 6 Jun 2026 06:06:47 +0200 Subject: [PATCH 6/8] context : fix off-by-one comparisons to n_gpu_layers (#24208) --- src/llama-context.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/llama-context.cpp b/src/llama-context.cpp index eff1d8f89..d0c314199 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -341,7 +341,7 @@ llama_context::llama_context( // enabling pipeline parallelism in the scheduler increases memory usage, so it is only done when necessary bool pipeline_parallel = model.n_devices() > 1 && - model.n_gpu_layers() > model.hparams.n_layer() && + model.n_gpu_layers() > model.hparams.n_layer_all && model.split_mode() == LLAMA_SPLIT_MODE_LAYER && cparams.offload_kqv && !model.has_tensor_overrides(); @@ -2351,7 +2351,7 @@ llm_graph_cb llama_context::graph_get_cb() const { // norm may be automatically assigned to the backend of the previous layer, increasing data transfer between backends // FIXME: fix in ggml_backend_sched - const bool full_offload = model.n_gpu_layers() > model.hparams.n_layer(); + const bool full_offload = model.n_gpu_layers() > model.hparams.n_layer_all; if (ubatch.n_tokens < 32 || full_offload) { if (il != -1 && strcmp(name, "norm") == 0) { const auto & dev_layer = model.dev_layer(il); From 5343f4502ab5273d7cef85012af020cad0182376 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?= Date: Sat, 6 Jun 2026 06:07:20 +0200 Subject: [PATCH 7/8] model : rename local n_layer_all variable (#24209) --- src/llama-model.cpp | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 137d3501e..0d23a605e 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1205,7 +1205,7 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { const auto & use_mlock = params.use_mlock; const auto & tensor_split = params.tensor_split; - const int n_layer = hparams.n_layer_all; + const int n_layer_all = hparams.n_layer_all; const int n_gpu_layers = this->n_gpu_layers(); const bool use_mmap_buffer = true; @@ -1262,10 +1262,10 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { splits[i] /= split_sum; } - const int i_gpu_start = std::max(n_layer + 1 - n_gpu_layers, 0); - const int act_gpu_layers = devices.empty() ? 0 : std::min(n_gpu_layers, n_layer + 1); + const int i_gpu_start = std::max(n_layer_all + 1 - n_gpu_layers, 0); + const int act_gpu_layers = devices.empty() ? 0 : std::min(n_gpu_layers, n_layer_all + 1); auto get_layer_buft_list = [&](int il) -> llama_model::impl::layer_dev { - const bool is_swa = il < n_layer && hparams.is_swa(il); + const bool is_swa = il < n_layer_all && hparams.is_swa(il); if (il < i_gpu_start || (il - i_gpu_start) >= act_gpu_layers) { LLAMA_LOG_DEBUG("load_tensors: layer %3d assigned to device %s, is_swa = %d\n", il, ggml_backend_dev_name(cpu_dev), is_swa); return {cpu_dev, &pimpl->cpu_buft_list}; @@ -1281,13 +1281,13 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { pimpl->dev_input = { cpu_dev, &pimpl->cpu_buft_list }; // assign the repeating layers to the devices according to the splits - pimpl->dev_layer.resize(n_layer); - for (int il = 0; il < n_layer; ++il) { + pimpl->dev_layer.resize(n_layer_all); + for (int il = 0; il < n_layer_all; ++il) { pimpl->dev_layer[il] = get_layer_buft_list(il); } // assign the output layer - pimpl->dev_output = get_layer_buft_list(n_layer); + pimpl->dev_output = get_layer_buft_list(n_layer_all); const auto TENSOR_NOT_REQUIRED = llama_model_loader::TENSOR_NOT_REQUIRED; @@ -1303,14 +1303,14 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { throw std::runtime_error("model has expert layers but no expert layers are used"); } - layers.resize(n_layer); + layers.resize(n_layer_all); // call the per-model loading function load_arch_tensors(ml); // generic pass: load optional per-tensor/per-expert ".scale" tensors (e.g. NVFP4 scale2) // this avoids having to add scale loading to every architecture - for (int i = 0; i < n_layer; ++i) { + for (int i = 0; i < n_layer_all; ++i) { auto & layer = layers[i]; // attention weight scales (per-tensor, shape {1}) @@ -1568,7 +1568,7 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { } if (llama_supports_gpu_offload()) { - const int n_gpu = std::min(n_gpu_layers, n_layer); + const int n_gpu = std::min(n_gpu_layers, n_layer_all); int n_repeating = n_gpu; if (n_repeating > 0) { @@ -1577,8 +1577,8 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { } LLAMA_LOG_INFO("%s: offloading %d repeating layers to GPU\n", __func__, n_repeating); - const int max_backend_supported_layers = n_layer + 1; - const int max_offloadable_layers = n_layer + 1; + const int max_backend_supported_layers = n_layer_all + 1; + const int max_offloadable_layers = n_layer_all + 1; LLAMA_LOG_INFO("%s: offloaded %d/%d layers to GPU\n", __func__, std::min(n_gpu_layers, max_offloadable_layers), max_backend_supported_layers); } From 5a69c974392020e514c3b2b2910bb92f847cb4c9 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Sat, 6 Jun 2026 09:11:35 +0200 Subject: [PATCH 8/8] vulkan: check coopmat2 features before reporting support (#24186) --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index df410368a..fc9bc8fe3 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -6349,6 +6349,15 @@ static void ggml_vk_print_gpu_info(size_t idx) { } #endif +#if defined(VK_NV_cooperative_matrix2) + VkPhysicalDeviceCooperativeMatrix2FeaturesNV coopmat2_features {}; + coopmat2_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_MATRIX_2_FEATURES_NV; + if (coopmat2_support) { + last_struct->pNext = (VkBaseOutStructure *)&coopmat2_features; + last_struct = (VkBaseOutStructure *)&coopmat2_features; + } +#endif + VkPhysicalDeviceCooperativeMatrixDecodeVectorFeaturesNV coopmat2_decode_vector_features {}; coopmat2_decode_vector_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_MATRIX_DECODE_VECTOR_FEATURES_NV; if (coopmat2_decode_vector_support) { @@ -6380,6 +6389,19 @@ static void ggml_vk_print_gpu_info(size_t idx) { #endif && ggml_vk_khr_cooperative_matrix_support(props2.properties, driver_props, device_architecture); +#if defined(VK_NV_cooperative_matrix2) && defined(GGML_VULKAN_COOPMAT2_GLSLC_SUPPORT) + coopmat2_support = coopmat2_support && + coopmat2_features.cooperativeMatrixWorkgroupScope && + coopmat2_features.cooperativeMatrixFlexibleDimensions && + coopmat2_features.cooperativeMatrixReductions && + coopmat2_features.cooperativeMatrixConversions && + coopmat2_features.cooperativeMatrixPerElementOperations && + coopmat2_features.cooperativeMatrixTensorAddressing && + coopmat2_features.cooperativeMatrixBlockLoads; +#else + coopmat2_support = false; +#endif + coopmat2_decode_vector_support = coopmat2_decode_vector_support && coopmat2_decode_vector_features.cooperativeMatrixDecodeVector; #if !defined(GGML_VULKAN_COOPMAT2_DECODE_VECTOR_GLSLC_SUPPORT) coopmat2_decode_vector_support = false;