model, mtmd: drama-3 tts support (#31613 )

* model, mtmd: drama-3 tts support * clean up * mtmd: drama3 helper stops at the codebook tower limit --------- Co-authored-by: Pascal <admin@serveurperso.com>
ui: show text and page thumbnails for PDF attachments, stop endless conversion retries (#31468 )
2026-11-04 09:57:44 -06:00 · 2026-11-04 15:15:09 +01:00 · 2026-11-04 15:46:51 +02:00 · 2026-11-04 15:46:06 +02:00 · 2026-11-04 15:45:29 +02:00 · 2026-11-04 14:42:39 +01:00
69 changed files with 4091 additions and 428 deletions
@@ -64,7 +64,7 @@ RUN mkdir -p /app/full \
    && cp requirements.txt /app/full \
    && cp .devops/tools.sh /app/full/tools.sh

-FROM docker.io/intel/oneapi-toolkit:$ONEAPI_VERSION AS base
+FROM docker.io/intel/deep-learning-essentials:$ONEAPI_VERSION AS base

 ARG BUILD_DATE=N/A
 ARG APP_VERSION=N/A
@@ -79,6 +79,8 @@ LABEL org.opencontainers.image.created=$BUILD_DATE \
      org.opencontainers.image.url=$IMAGE_URL \
      org.opencontainers.image.source=$IMAGE_SOURCE

+COPY --from=build /opt/intel/oneapi/dnnl /opt/intel/oneapi/dnnl
+
 #Following versions are for multiple GPUs, since 26.x has known issue:
 #   https://github.com/ggml-org/llama.cpp/issues/21747,
 #   https://github.com/intel/compute-runtime/issues/921.
@@ -1,4 +1,4 @@
-# date: Wed Oct 28 12:38:26 EET 2026
+# date: Wed Nov  4 12:03:31 EET 2026
 # this file is auto-generated by scripts/gen-authors.sh

 Нияз Гарифзянов <112617865+garrnizon@users.noreply.github.com>
@@ -48,6 +48,7 @@ Abheek Gulati <abheekg@hotmail.com>
 abhijain1204fujitsu <139222713+abhijain1204fujitsu@users.noreply.github.com>
 Abhijit Ramesh <abhijitramesh2k@gmail.com>
 abhijitb11 <113058133+abhijitb11@users.noreply.github.com>
+Abhijith Jajam <129675815+jajamabhijith@users.noreply.github.com>
 Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com>
 Abhinay Krishna <abhinaykrishna60@gmail.com>
 Abhiram <78226909+geckguy@users.noreply.github.com>
@@ -203,14 +204,17 @@ Ankit Khandelwal <ankk98@gmail.com>
 Ankur Verma <31362771+ankurvdev@users.noreply.github.com>
 anon998 <131767832+anon998@users.noreply.github.com>
 Anri Lombard <anri.m.lombard@gmail.com>
+Anthony Maltsev <77708730+anthonymaltsev@users.noreply.github.com>
 Anthony Umfer <aumfer@gmail.com>
 Anthony Van de Gejuchte <anthonyvdgent@gmail.com>
 Anthony Wang <anthonywang903@gmail.com>
 antichristHater <142441588+antichristHater@users.noreply.github.com>
 Antoine Viallon <antoine@lesviallon.fr>
+Anton I <inflyer@hotmail.com>
 Anton Mitkov <anton_b_mitkov@abv.bg>
 Anton Mitkov <anton.mitkov@codeplay.com>
 Antonis Makropoulos <benuix@gmail.com>
+anubhavfujitsu <Anubhav.Srivastava@fujitsu.com>
 Anudit Nagar <nagaranudit@gmail.com>
 Anuj Attri <anujattri01@gmail.com>
 anujj <ajalota@nvidia.com>
@@ -293,6 +297,7 @@ Bernat Vadell <hounter.caza@gmail.com>
 Bernhard M. Wiedemann <githubbmwprimary@lsmod.de>
 Bert Wagner <github@bertwagner.com>
 Bertay Eren <39909689+bertaye@users.noreply.github.com>
+Bhargav Krish <bkrishnamurt@nvidia.com>
 Bhavik Sharda <10757940+BLSharda@users.noreply.github.com>
 bhubbb <79117352+bhubbb@users.noreply.github.com>
 Bill Sideris <bill88t@feline.gr>
@@ -533,6 +538,7 @@ Don Mahurin <dmahurin@users.noreply.github.com>
 Dong Won Kim <63934649+ddwkim@users.noreply.github.com>
 Donghyeon Jeong <54725479+djeong20@users.noreply.github.com>
 Dongliang Wei <121270393+wdl339@users.noreply.github.com>
+Dongsheng Luo <luods1995@outlook.com>
 Doomsdayrs <38189170+Doomsdayrs@users.noreply.github.com>
 DooWoong Lee (David) <manics99@naver.com>
 DorianRudolph <dorianrudo97@googlemail.com>
@@ -612,6 +618,7 @@ Faez Shakil <faez.shakil@gmail.com>
 fairydreaming <166155368+fairydreaming@users.noreply.github.com>
 Faisal Zaghloul <faisal.zaghloul@gmail.com>
 Faisal Zaghloul <quic_fzaghlou@quicinc.com>
+FamousMonsterr <famousmonster@ya.ru>
 Fan Shupei <dymarkfan@outlook.com>
 FantasyGmm <16450052+FantasyGmm@users.noreply.github.com>
 fanyang <fanyang89@outlook.com>
@@ -654,6 +661,7 @@ Frederik Vogel <Schaltfehler@users.noreply.github.com>
 Fredrik Hultin <noname@nurd.se>
 fredzillman <fzillman@gmail.com>
 frob <rick+github@frob.com.au>
+frogscooper <frogscooper@gmail.com>
 Frost-54 <91168148+Frost-54@users.noreply.github.com>
 Frosty40 <newjordan@gmail.com>
 Funtowicz Morgan <mfuntowicz@users.noreply.github.com>
@@ -694,6 +702,7 @@ Gilad S. <7817232+giladgd@users.noreply.github.com>
 Giovanni Rivera <rivera.giovanni271@gmail.com>
 github-actions[bot] <github-actions[bot]@users.noreply.github.com>
 GittyBurstein <g0534163997@gmail.com>
+Giulianno <lernexbusiness@gmail.com>
 Giuseppe Scrivano <giuseppe@scrivano.org>
 Giuseppe Scrivano <gscrivan@redhat.com>
 GiviMAD <GiviMAD@users.noreply.github.com>
@@ -762,6 +771,7 @@ Henry147147 <44851451+Henry147147@users.noreply.github.com>
 Herman Semenoff <GermanAizek@yandex.ru>
 Herman Semenov <GermanAizek@yandex.ru>
 Hesen Peng <hesen.peng@gmail.com>
+hesiru <siruhe666@gmail.com>
 hey-gm <heitorgm@outlook.com>
 HighDoping <highdoping@gmail.com>
 HimariO <dsfhe49854@gmail.com>
@@ -919,6 +929,7 @@ Jesse Ikonen <jesse.ikonen@gmail.com>
 Jesse Jojo Johnson <williamsaintgeorge@gmail.com>
 Jesse LaRose <jesse@taey.ai>
 Jesse Posner <jesse.posner@gmail.com>
+Jesse Symons <155509533+Torchit1@users.noreply.github.com>
 Jesus Gulfo <jesdga95@gmail.com>
 Jesus Talavera <145992175+jesus-talavera-ibm@users.noreply.github.com>
 Jetson Tan <tanzongyouyi@outlook.com>
@@ -1071,6 +1082,7 @@ Kevin Liu <4396kevinliu@gmail.com>
 Kevin Lo <kevlo@kevlo.org>
 Kevin Pouget <kpouget@redhat.com>
 Kevin Wang <kevmo314@gmail.com>
+Keyu Zhu <inuebisu@outlook.com>
 Khashayar Ghafouri <43180261+khashayarghafouri@users.noreply.github.com>
 khimaros <me@khimaros.com>
 Kilian Hu <90606809+kilian-hu@users.noreply.github.com>
@@ -1078,6 +1090,7 @@ Kilian Krampf <kilian@krampf.de>
 kiltyj <kiltyj@gmail.com>
 Kim S. <polydecay@users.noreply.github.com>
 kimminsu <80271594+kimminsu38oo@users.noreply.github.com>
+Kirill A. Korinsky <kirill@korins.ky>
 KITAITI Makoto <KitaitiMakoto@gmail.com>
 kiwi <122582483+kiwi142857@users.noreply.github.com>
 klosax <131523366+klosax@users.noreply.github.com>
@@ -1136,6 +1149,7 @@ leejet <leejet714@gmail.com>
 Leng Yue <lengyue@lengyue.me>
 Lennart Austenfeld <53152202+l-austenfeld@users.noreply.github.com>
 leo-pony <nengjunma@outlook.com>
+Leon <118652405+LeON-Nie-code@users.noreply.github.com>
 Leon Knauer <git@leonknauer.com>
 Leonard Mosescu <tlemo@users.noreply.github.com>
 leonardHONG <2695316095@qq.com>
@@ -1146,8 +1160,10 @@ lexasub <lexakopp2212@gmail.com>
 lgai-exaone <exaonemodels@lgresearch.ai>
 lhez <lih@qti.qualcomm.com>
 lhez <quic_lih@quicinc.com>
+Li Chuyang <1418587903@qq.com>
 Li Pengzhan <151381994+Lpzhan931@users.noreply.github.com>
 Li Tan <tanliboy@gmail.com>
+li-lizhe <82501526+li-lizhe@users.noreply.github.com>
 liangliangchang <lichang@amd.com>
 liminfei-amd <91481003+liminfei-amd@users.noreply.github.com>
 limitedAtonement <limitedAtonement@users.noreply.github.com>
@@ -1166,6 +1182,7 @@ LoganDark <github@logandark.mozmail.com>
 Loïc Carrère <loic.carrere@gmail.com>
 lon <114724657+longregen@users.noreply.github.com>
 loonerin <132926317+loonerin@users.noreply.github.com>
+Lorenzo Bona <lorenzobona@tuta.com>
 losewayy <1134285612@qq.com>
 LostRuins <39025047+LostRuins@users.noreply.github.com>
 LostRuins Concedo <39025047+LostRuins@users.noreply.github.com>
@@ -1253,6 +1270,7 @@ Masaya, Kato <62578291+msy-kato@users.noreply.github.com>
 mashdragon <122402293+mashdragon@users.noreply.github.com>
 Mason Milburn <masonmilby@gmail.com>
 MasterYi1024 <39848311+MasterYi1024@users.noreply.github.com>
+Maternion <98753158+maternion@users.noreply.github.com>
 Mateusz Charytoniuk <mateusz.charytoniuk@protonmail.com>
 Matheus C. França <matheus-catarino@hotmail.com>
 Matheus Gabriel Alves Silva <matheusgasource@gmail.com>
@@ -1350,6 +1368,7 @@ Miwa / Ensan <63481257+ensan-hcl@users.noreply.github.com>
 miyan <1138989048@qq.com>
 mj-shifu <77107165+mj-shifu@users.noreply.github.com>
 mkoker <132301062+mkoker@users.noreply.github.com>
+mmalpe-nv <mmalpe@nvidia.com>
 mmyjona <jonathan.gonse@gmail.com>
 mnehete32 <33429707+mnehete32@users.noreply.github.com>
 Mohamed Elashri <git@elashri.com>
@@ -1368,6 +1387,7 @@ mtmcp <141645996+mtmcp@users.noreply.github.com>
 muggle-stack <promuggle@qq.com>
 Muhammad Saad <saadydev@gmail.com>
 Muhammad Salem <salem.ebo@gmail.com>
+MuoDoo <36472129+MuoDoo@users.noreply.github.com>
 Murilo Santana <mvrilo@gmail.com>
 Musab Gultekin <musabgultekin@users.noreply.github.com>
 musoles <135031143+musoles@users.noreply.github.com>
@@ -1416,10 +1436,12 @@ Nicolas Budyn <16596345+nbud@users.noreply.github.com>
 Nicolas Mowen <nickmowen213@gmail.com>
 Nicolás Pérez <nicolas_perez@brown.edu>
 Nicolò Scipione <nicolo.scipione@codeplay.com>
+Nidhi <nidhisharma16635@gmail.com>
 Nigel Bosch <pnigelb@gmail.com>
 Nigel Jones <nigel.l.jones+git@gmail.com>
 Nik Bogatyrev <50734723+notforu777@users.noreply.github.com>
 Nikhil Jain <nikhil.jain0987@gmail.com>
+Nikhil Sharma <150424968+nikhilfujitsu@users.noreply.github.com>
 Nikita Gordeev <iamfobey@gmail.com>
 Nikita Sarychev <42014488+sARY77@users.noreply.github.com>
 Niklas Korz <niklas@niklaskorz.de>
@@ -1592,6 +1614,7 @@ Raman Shinde <raman.shinde15@gmail.com>
 Rand Xie <randxiexyy29@gmail.com>
 Randall Fitzgerald <randall@dasaku.net>
 Random Fly <renfei8@live.cn>
+randomized <286843573@qq.com>
 rankaiyx <rankaiyx@foxmail.com>
 rankaiyx <rankaiyx@rankaiyx.com>
 RapidMark <32768622+RapidMark@users.noreply.github.com>
@@ -2009,6 +2032,7 @@ Tungsten842 <quantmint@protonmail.com>
 Tushar <ditsuke@protonmail.com>
 tv1wnd <55383215+tv1wnd@users.noreply.github.com>
 tyronecai <tyronecai@163.com>
+Uber <20715993+YagUber@users.noreply.github.com>
 ubergarm <leimgrub@gmail.com>
 ubik2 <ubik2@users.noreply.github.com>
 UEXTM.com <84163508+uextm@users.noreply.github.com>
@@ -2242,4 +2266,5 @@ Zijun Yu <zijun.yu@intel.com>
 zql <37731799+zqlcode@users.noreply.github.com>
 zrm <trustiosity.zrm@gmail.com>
 Zsapi <martin1.zsapka@gmail.com>
+zyang-dev <267119621+zyang-dev@users.noreply.github.com>
 zzzzwc <tiddar@foxmail.com>
@@ -4,7 +4,7 @@ include(CheckIncludeFileCXX)

 ### llama.cpp version
 set(LLAMA_VERSION_MAJOR 0)
-set(LLAMA_VERSION_MINOR 8)
+set(LLAMA_VERSION_MINOR 9)
 set(LLAMA_VERSION_PATCH 0)
 set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")

@@ -710,7 +710,7 @@ static void common_params_apply_system_config(common_params & params, llama_exam
    std::vector<std::filesystem::path> paths;

 #if defined(_WIN32)
-    const std::filesystem::path program_data = common_get_path_from_env("PROGRAMDATA");
+    const std::filesystem::path program_data = common_get_env("PROGRAMDATA");
    if (!program_data.empty()) {
        paths.push_back(program_data / "llama.cpp" / "config.ini");
    }
@@ -983,20 +983,9 @@ void common_set_env(const std::string & name, const std::string & value) {
 #endif
 }

-std::filesystem::path common_get_path_from_env(const std::string & name) {
-#if defined(_WIN32)
-    const std::wstring wname = utf8_to_wstring(name);
-    const wchar_t * wvalue = _wgetenv(wname.c_str());
-    return wvalue ? std::filesystem::path(wvalue) : std::filesystem::path();
-#else
-    const char * value = std::getenv(name.c_str());
-    return value ? std::filesystem::path(value) : std::filesystem::path();
-#endif
-}
-
 #if !defined(_WIN32)
 static std::filesystem::path get_home_directory() {
-    std::filesystem::path home = common_get_path_from_env("HOME");
+    std::filesystem::path home = common_get_env("HOME");
    if (!home.empty()) {
        return home;
    }
@@ -1009,19 +998,19 @@ static std::filesystem::path get_home_directory() {
 #endif

 std::filesystem::path fs_get_cache_directory() {
-    std::filesystem::path cache_directory = common_get_path_from_env("LLAMA_CACHE");
+    std::filesystem::path cache_directory = common_get_env("LLAMA_CACHE");
    if (!cache_directory.empty()) {
        return cache_directory;
    }
 #if defined(_WIN32)
-    cache_directory = common_get_path_from_env("LOCALAPPDATA");
+    cache_directory = common_get_env("LOCALAPPDATA");
    if (cache_directory.empty()) {
        throw std::runtime_error("Failed to find %LOCALAPPDATA% directory");
    }
 #elif defined(__APPLE__)
    cache_directory = get_home_directory() / "Library/Caches";
 #else
-    cache_directory = common_get_path_from_env("XDG_CACHE_HOME");
+    cache_directory = common_get_env("XDG_CACHE_HOME");
    if (cache_directory.empty()) {
        cache_directory = get_home_directory() / ".cache";
    }
@@ -1032,12 +1021,12 @@ std::filesystem::path fs_get_cache_directory() {
 std::filesystem::path fs_get_config_directory() {
    std::filesystem::path config_directory;
 #if defined(_WIN32)
-    config_directory = common_get_path_from_env("APPDATA");
+    config_directory = common_get_env("APPDATA");
    if (config_directory.empty()) {
        throw std::runtime_error("Failed to find %APPDATA% directory");
    }
 #else
-    config_directory = common_get_path_from_env("XDG_CONFIG_HOME");
+    config_directory = common_get_env("XDG_CONFIG_HOME");
    if (config_directory.empty()) {
        config_directory = get_home_directory() / ".config";
    }
@@ -919,9 +919,6 @@ std::string  wstring_to_utf8(const std::wstring & str);
 std::string common_get_env(const std::string & name);
 void        common_set_env(const std::string & name, const std::string & value);

-// reads a path from the environment, an unset variable gives an empty path
-std::filesystem::path common_get_path_from_env(const std::string & name);
-
 //
 // Filesystem utils
 //
@@ -748,7 +748,8 @@ static hf_cache::hf_file find_best_model(const hf_cache::hf_files & files,
    }

    for (const auto & t : tags) {
-        std::regex pattern(t + "[.-]", std::regex::icase);
+        // match the tag as a full token, so ":F16" does not match a BF16 file
+        std::regex pattern("(^|[-./])" + t + "[.-]", std::regex::icase);
        for (const auto & f : files) {
            if (gguf_filename_is_model(f.path) &&
                std::regex_search(f.path, pattern)) {
@@ -45,7 +45,7 @@ static fs::path get_cache_directory() {
            {HOME_DIR,                fs::path(".cache") / "huggingface" / "hub"}
        };
        for (const auto & entry : entries) {
-            if (fs::path base = common_get_path_from_env(entry.var); !base.empty()) {
+            if (fs::path base = common_get_env(entry.var); !base.empty()) {
                return entry.path.empty() ? base : base / entry.path;
            }
        }
@@ -95,6 +95,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
    "FalconForCausalLM": "falcon",
    "FalconH1ForCausalLM": "falcon_h1",
    "FalconMambaForCausalLM": "mamba",
+    "FishOmniModel": "drama3",
    "GPT2LMHeadModel": "gpt2",
    "GPTBigCodeForCausalLM": "starcoder",
    "GPTNeoXForCausalLM": "gptneox",
@@ -346,6 +347,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
    "DotsOCRForCausalLM": "dotsocr",
    "EmbeddingGemma2Model": "gemma",
    "Exaone4_5_ForConditionalGeneration": "exaone",
+    "FishOmniModel": "drama3",
    "Gemma3ForConditionalGeneration": "gemma",
    "Gemma3nForConditionalGeneration": "gemma",
    "Gemma4ForConditionalGeneration": "gemma",
@@ -494,7 +494,7 @@ class ModelBase:
                        s = self.model_tensors[name]
                        self.model_tensors[weight_name] = lambda w=w, s=s: dequant_bitnet(w(), s())
                        tensors_to_remove.append(name)
-            elif quant_method == "fp8":
+            elif quant_method in ("fp8", "bsgemm-moe"):  # bsgemm-moe: drama-3, same block layout
                block_size = quant_config.get("weight_block_size")
                for name in self.model_tensors.keys():
                    if name.endswith("_scale_inv"):
@@ -0,0 +1,251 @@
+from __future__ import annotations
+
+from typing import Any, Callable, Iterable, TYPE_CHECKING
+
+import torch
+
+if TYPE_CHECKING:
+    from torch import Tensor
+
+from .base import ModelBase, MmprojModel, gguf
+from .qwen import Qwen3_5MoeTextModel
+
+# Fish Audio Drama 3 (FishOmniModel)
+# text model: Qwen3.5-MoE backbone, samples the main token of each frame
+# mmproj: codebook tower and Fish DAC decoder (codec.pth)
+
+_TOWER_SHARED = "audio_decoder_mtp_shared.layers."
+_TOWER_RESIDUAL = "audio_decoder_mtp.layers."
+
+
+@ModelBase.register("FishOmniModel")
+@ModelBase.example("fishaudio/drama-3")
+class Drama3TextModel(Qwen3_5MoeTextModel):
+    model_arch = gguf.MODEL_ARCH.QWEN35MOE
+    no_mtp = True  # the checkpoint has no MTP head
+
+    @classmethod
+    def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
+        if item[0].startswith("audio_decoder"):
+            return None
+        return super().filter_tensors(item)
+
+
+@ModelBase.register("FishOmniModel")
+@ModelBase.example("fishaudio/drama-3")
+class Drama3MmprojModel(MmprojModel):
+    has_vision_encoder = False
+    has_audio_encoder = False
+
+    _LAYER_MAP = {
+        "enorm.weight":                                 (gguf.MODEL_TENSOR.A_GEN_CODE_ENORM, True),
+        "hnorm.weight":                                 (gguf.MODEL_TENSOR.A_GEN_CODE_HNORM, True),
+        "eh_proj.weight":                               (gguf.MODEL_TENSOR.A_GEN_CODE_EH_PROJ, False),
+        "semantic_embedding.weight":                    (gguf.MODEL_TENSOR.A_GEN_CODE_BLK_EMBD, False),
+        "codebook_embedding.weight":                    (gguf.MODEL_TENSOR.A_GEN_CODE_BLK_EMBD, False),
+        "final_layernorm.weight":                       (gguf.MODEL_TENSOR.A_GEN_CODE_BLK_OUT_NORM, True),
+        "transformer_layer.attention_norm.weight":      (gguf.MODEL_TENSOR.A_GEN_CODE_ATTN_NORM, True),
+        "transformer_layer.ffn_norm.weight":            (gguf.MODEL_TENSOR.A_GEN_CODE_FFN_NORM, True),
+        "transformer_layer.attention.q_norm.weight":    (gguf.MODEL_TENSOR.A_GEN_CODE_ATTN_Q_NORM, True),
+        "transformer_layer.attention.k_norm.weight":    (gguf.MODEL_TENSOR.A_GEN_CODE_ATTN_K_NORM, True),
+        "transformer_layer.attention.wo.weight":        (gguf.MODEL_TENSOR.A_GEN_CODE_ATTN_OUT, False),
+        "transformer_layer.feed_forward.w1.weight":     (gguf.MODEL_TENSOR.A_GEN_CODE_FFN_GATE, False),
+        "transformer_layer.feed_forward.w3.weight":     (gguf.MODEL_TENSOR.A_GEN_CODE_FFN_UP, False),
+        "transformer_layer.feed_forward.w2.weight":     (gguf.MODEL_TENSOR.A_GEN_CODE_FFN_DOWN, False),
+    }
+
+    _heads: dict[int, Tensor]
+
+    def __init__(self, *args, **kwargs):
+        super().__init__(*args, **kwargs)
+        self._heads = {}
+
+    def get_audio_config(self) -> dict[str, Any] | None:
+        cfg = dict(self.global_config["audio_decoder_config"])
+        cfg["num_hidden_layers"] = cfg["shared_enrichment_layers"] + cfg["num_codebooks"] - 1
+        return cfg
+
+    def set_gguf_parameters(self):
+        cfg = self.hparams
+        self.gguf_writer.add_file_type(self.ftype)
+        self.gguf_writer.add_clip_has_gen_audio_encoder(True)
+        self.gguf_writer.add_clip_gen_audio_projector_type(gguf.VisionProjectorType.DRAMA3_GEN)
+        self.gguf_writer.add_gen_audio_projection_dim(self.n_embd_text)
+        self.gguf_writer.add_gen_audio_embedding_length(cfg["dim"])
+        self.gguf_writer.add_gen_audio_feed_forward_length(cfg["intermediate_size"])
+        self.gguf_writer.add_gen_audio_block_count(cfg["num_hidden_layers"])
+        self.gguf_writer.add_gen_audio_head_count(cfg["n_head"])
+        self.gguf_writer.add_gen_audio_head_count_kv(cfg["n_local_heads"])
+        self.gguf_writer.add_gen_audio_attention_layernorm_eps(cfg["norm_eps"])
+        # other tower and codec hparams are hardcoded in clip.cpp
+
+    @classmethod
+    def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
+        name = item[0]
+        if name.startswith("audio_decoder") or name == "model.language_model.embed_tokens.weight":
+            return item
+        return None
+
+    def tensor_force_quant(self, name, new_name, bid, n_dims):
+        # conv1d/conv1d_dw kernels must be F16, ggml_conv_1d(_dw) has no BF16 path
+        if new_name.endswith(".weight") and (
+            new_name in ("a.gen.wav.dac.entry.weight", "a.gen.wav.dac.post_conv.weight")
+            or (".up.blk." in new_name and new_name.endswith(".dwconv.weight"))
+            or (".dac.blk." in new_name and (new_name.endswith(".conv1.weight") or new_name.endswith(".conv2.weight")))
+        ):
+            return gguf.GGMLQuantizationType.F16
+        # keep ConvTranspose1d kernels and folded codebook tables in F32
+        if new_name.endswith(".conv.weight") and (".up.blk." in new_name or ".dac.blk." in new_name):
+            return gguf.GGMLQuantizationType.F32
+        if new_name.startswith("a.gen.wav.quant."):
+            return gguf.GGMLQuantizationType.F32
+        return super().tensor_force_quant(name, new_name, bid, n_dims)
+
+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        T = gguf.MODEL_TENSOR
+        cfg = self.hparams
+
+        if name.startswith("a.gen.wav."):
+            yield (name, data_torch)
+            return
+
+        if name == "model.language_model.embed_tokens.weight":
+            # semantic band of the backbone embeddings
+            begin = self.global_config["semantic_start_token_id"]
+            yield (self.format_tensor_name(T.A_GEN_CODE_OUT_EMBD), data_torch[begin:begin + cfg["vocab_size"]])
+            return
+
+        if name == "audio_decoder.codebook_embeddings.weight":
+            n_cb = cfg["num_codebooks"]
+            yield (self.format_tensor_name(T.A_GEN_CODE_EMBD), data_torch.reshape(n_cb, -1, data_torch.shape[-1]))
+            return
+
+        if name.startswith("audio_decoder_mtp.heads."):
+            n_heads = cfg["num_codebooks"] - 1
+            self._heads[int(name.split(".")[2])] = data_torch
+            if len(self._heads) < n_heads:
+                return
+            yield (self.format_tensor_name(T.A_GEN_CODE_HEAD), torch.stack([self._heads.pop(i) for i in range(n_heads)]))
+            return
+
+        for prefix, offset in ((_TOWER_SHARED, 0), (_TOWER_RESIDUAL, cfg["shared_enrichment_layers"])):
+            if not name.startswith(prefix):
+                continue
+            idx, key = name[len(prefix):].split(".", 1)
+            il = int(idx) + offset
+            if key == "transformer_layer.attention.wqkv.weight":
+                # rows are [query | gate | key | value]
+                n_q = cfg["n_head"] * cfg["head_dim"]
+                n_kv = cfg["n_local_heads"] * cfg["head_dim"]
+                q, gate, k, v = data_torch.split([n_q, n_q, n_kv, n_kv], dim=0)
+                yield (self.format_tensor_name(T.A_GEN_CODE_ATTN_Q, il), q)
+                yield (self.format_tensor_name(T.A_GEN_CODE_ATTN_GATE, il), gate)
+                yield (self.format_tensor_name(T.A_GEN_CODE_ATTN_K, il), k)
+                yield (self.format_tensor_name(T.A_GEN_CODE_ATTN_V, il), v)
+                return
+            tensor, zero_centered = self._LAYER_MAP[key]
+            if zero_centered:
+                data_torch = data_torch.float() + 1
+            yield (self.format_tensor_name(tensor, il), data_torch)
+            return
+
+        raise ValueError(f"unexpected tensor {name}")
+
+    def generate_extra_tensors(self) -> Iterable[tuple[str, Tensor]]:
+        # only the DAC decoder, the encoder is for voice cloning (not supported yet)
+        sd = torch.load(self.dir_model / "codec.pth", map_location="cpu", mmap=True, weights_only=True)
+        sd = sd.get("state_dict", sd)
+        T = gguf.MODEL_TENSOR
+
+        def get(name: str) -> Tensor:
+            return sd[name].float()
+
+        def conv(prefix: str) -> Tensor:
+            # fold weight norm
+            if prefix + ".conv.weight" in sd:
+                return get(prefix + ".conv.weight")
+            g = get(prefix + ".conv.parametrizations.weight.original0")
+            v = get(prefix + ".conv.parametrizations.weight.original1")
+            return g * v / v.norm(dim=(1, 2), keepdim=True)
+
+        def quant_table(prefix: str) -> Tensor:
+            # fold out_proj into the codebook, so decoding is a row lookup
+            g = get(prefix + ".out_proj.weight_g")
+            v = get(prefix + ".out_proj.weight_v")
+            w = (g * v / v.norm(dim=(1, 2), keepdim=True)).squeeze(-1)
+            return get(prefix + ".codebook.weight") @ w.T + get(prefix + ".out_proj.bias")
+
+        def snake(alpha: Tensor) -> tuple[Tensor, Tensor]:
+            alpha = alpha.flatten()
+            return alpha, 1.0 / (alpha + 1e-9)
+
+        q = "quantizer."
+        yield (self.format_tensor_name(T.A_GEN_WAV_QUANT_FIRST_CB), quant_table(q + "semantic_quantizer.quantizers.0"))
+        n_rest = sum(1 for k in sd if k.startswith(q + "quantizer.quantizers.") and k.endswith(".codebook.weight"))
+        yield (self.format_tensor_name(T.A_GEN_WAV_QUANT_REST_CB),
+               torch.stack([quant_table(f"{q}quantizer.quantizers.{i}") for i in range(n_rest)]))
+
+        # post_module uses interleaved RoPE, permute q/k rows so NEOX RoPE gives the same result
+        tfm = q + "post_module."
+        n_head, d_head = 16, 64
+        order = torch.cat([torch.arange(0, d_head, 2), torch.arange(1, d_head, 2)])
+
+        def neox_rows(w: Tensor) -> Tensor:
+            return w.reshape(n_head, d_head, -1)[:, order].reshape(w.shape)
+
+        n_tfm = sum(1 for k in sd if k.startswith(tfm + "layers.") and k.endswith(".ffn_norm.weight"))
+        for il in range(n_tfm):
+            p = f"{tfm}layers.{il}."
+            wq, wk, wv = get(p + "attention.wqkv.weight").chunk(3, dim=0)
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_ATTN_Q, il), neox_rows(wq))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_ATTN_K, il), neox_rows(wk))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_ATTN_V, il), wv)
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_ATTN_OUT, il), get(p + "attention.wo.weight"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_ATTN_NORM, il), get(p + "attention_norm.weight"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_ATTN_SCALE, il), get(p + "attention_layer_scale.gamma"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_FFN_NORM, il), get(p + "ffn_norm.weight"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_FFN_GATE, il), get(p + "feed_forward.w1.weight"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_FFN_UP, il), get(p + "feed_forward.w3.weight"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_FFN_DOWN, il), get(p + "feed_forward.w2.weight"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_TFM_FFN_SCALE, il), get(p + "ffn_layer_scale.gamma"))
+        yield (self.format_tensor_name(T.A_GEN_WAV_TFM_OUTPUT_NORM), get(tfm + "norm.weight"))
+
+        # upsample: 2x (ConvTranspose1d + ConvNeXt)
+        for il in range(2):
+            p = f"{q}upsample.{il}."
+            yield (self.format_tensor_name(T.A_GEN_WAV_UP_CONV, il, suffix=".weight"), conv(p + "0"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_UP_CONV, il, suffix=".bias"), get(p + "0.conv.bias"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_UP_DWCONV, il, suffix=".weight"), conv(p + "1.dwconv"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_UP_DWCONV, il, suffix=".bias"), get(p + "1.dwconv.conv.bias"))
+            for key, tensor in (("norm", T.A_GEN_WAV_UP_NORM), ("pwconv1", T.A_GEN_WAV_UP_PW1), ("pwconv2", T.A_GEN_WAV_UP_PW2)):
+                yield (self.format_tensor_name(tensor, il, suffix=".weight"), get(f"{p}1.{key}.weight"))
+                yield (self.format_tensor_name(tensor, il, suffix=".bias"), get(f"{p}1.{key}.bias"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_UP_GAMMA, il, suffix=""), get(p + "1.gamma"))
+
+        # DAC: entry conv, 4x (Snake, ConvTranspose1d, 3 residual units), Snake, post conv
+        d = "decoder.model."
+        yield (self.format_tensor_name(T.A_GEN_WAV_DAC_ENTRY, suffix=".weight"), conv(d + "0"))
+        yield (self.format_tensor_name(T.A_GEN_WAV_DAC_ENTRY, suffix=".bias"), get(d + "0.conv.bias"))
+        for bid in range(4):
+            p = f"{d}{bid + 1}.block."
+            a, b = snake(get(p + "0.alpha"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_DAC_UP_SNAKE, bid, suffix=".alpha"), a)
+            yield (self.format_tensor_name(T.A_GEN_WAV_DAC_UP_SNAKE, bid, suffix=".beta"), b)
+            yield (self.format_tensor_name(T.A_GEN_WAV_DAC_UP_CONV, bid, suffix=".weight"), conv(p + "1"))
+            yield (self.format_tensor_name(T.A_GEN_WAV_DAC_UP_CONV, bid, suffix=".bias"), get(p + "1.conv.bias"))
+            for xid in range(3):
+                r = f"{p}{xid + 2}.block."
+                for act, tensor in (("0", T.A_GEN_WAV_DAC_RES_ACT1), ("2", T.A_GEN_WAV_DAC_RES_ACT2)):
+                    a, b = snake(get(r + act + ".alpha"))
+                    name = gguf.TENSOR_NAMES[tensor].format(bid=bid, xid=xid)
+                    yield (name + ".alpha", a)
+                    yield (name + ".beta", b)
+                for cv, tensor in (("1", T.A_GEN_WAV_DAC_RES_CONV1), ("3", T.A_GEN_WAV_DAC_RES_CONV2)):
+                    name = gguf.TENSOR_NAMES[tensor].format(bid=bid, xid=xid)
+                    yield (name + ".weight", conv(r + cv))
+                    yield (name + ".bias", get(r + cv + ".conv.bias"))
+        a, b = snake(get(d + "5.alpha"))
+        yield (self.format_tensor_name(T.A_GEN_WAV_DAC_POST_SNAKE, suffix=".alpha"), a)
+        yield (self.format_tensor_name(T.A_GEN_WAV_DAC_POST_SNAKE, suffix=".beta"), b)
+        yield (self.format_tensor_name(T.A_GEN_WAV_DAC_POST_CONV, suffix=".weight"), conv(d + "6"))
+        yield (self.format_tensor_name(T.A_GEN_WAV_DAC_POST_CONV, suffix=".bias"), get(d + "6.conv.bias"))
@@ -4,7 +4,7 @@ project("ggml" C CXX ASM)

 ### GGML Version
 set(GGML_VERSION_MAJOR 0)
-set(GGML_VERSION_MINOR 28)
+set(GGML_VERSION_MINOR 29)
 set(GGML_VERSION_PATCH 0)
 set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")

@@ -29,7 +29,7 @@ void ggml_quantize_mat_q8_0_4x8(const float * GGML_RESTRICT x, void * GGML_RESTR
    assert(k % QK8_0 == 0);
    const int nb = k / QK8_0;

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
    block_q8_0x4 * GGML_RESTRICT y = (block_q8_0x4 *) vy;
    const size_t vl_calc = __riscv_vsetvl_e32m8(QK8_0);
    const size_t vl_save = __riscv_vsetvl_e64m2(4);
@@ -345,7 +345,7 @@ static inline int32x4_t ggml_nvfp4_dot8(const int8x8_t q4_lo, const int8x8_t q8_
 #include <immintrin.h>
 #endif

-#ifdef __riscv_v_intrinsic
+#if defined(__riscv_v)
 #include <riscv_vector.h>
 #endif

@@ -746,7 +746,7 @@ static void ggml_init_arm_arch_features(void) {}
 #endif
 #endif // __ARM_ARCH

-#if defined(__riscv) && defined(__riscv_v_intrinsic)
+#if defined(__riscv) && defined(__riscv_v)
 #include <riscv_vector.h>
 static void ggml_init_riscv_arch_features(void) {
    ggml_riscv_arch_features.rvv_vlen = __riscv_vlenb();
@@ -3556,7 +3556,7 @@ void ggml_cpu_fp16_to_fp32(const ggml_fp16_t * x, float * y, int64_t n) {
        _mm_storeu_ps(y + i, y_vec);
    }

-#elif defined(__riscv_v_intrinsic) && defined(__riscv_zvfhmin)
+#elif defined(__riscv_v) && defined(__riscv_zvfhmin)
    // calculate step size
    const int epr = __riscv_vsetvlmax_e16m2();
    const int step = epr * 2;
@@ -3626,7 +3626,7 @@ void ggml_cpu_bf16_to_fp32(const ggml_bf16_t * x, float * y, int64_t n) {
                                        (const __m128i *)(x + i))),
                                16)));
    }
-#elif defined(__riscv_v_intrinsic) && defined(__riscv_zvfbfmin)
+#elif defined(__riscv_v) && defined(__riscv_zvfbfmin)
    // calculate step size
    const int epr = __riscv_vsetvlmax_e16m2();
    const int step = epr * 2;
@@ -3746,7 +3746,7 @@ int ggml_cpu_has_arm_fma(void) {
 }

 int ggml_cpu_has_riscv_v(void) {
-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
    return 1;
 #else
    return 0;
@@ -3754,7 +3754,7 @@ int ggml_cpu_has_riscv_v(void) {
 }

 int ggml_cpu_get_rvv_vlen(void) {
-#if defined(__riscv) && defined(__riscv_v_intrinsic)
+#if defined(__riscv) && defined(__riscv_v)
    return ggml_riscv_arch_features.rvv_vlen;
 #else
    return 0;
@@ -69,7 +69,7 @@
 #define VECTOR_REGISTERS 16
 #endif

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
 #define LMUL 4
 #endif

@@ -180,7 +180,7 @@ inline float32x4_t madd(float32x4_t a, float32x4_t b, float32x4_t c) {
 }
 #endif

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
 template <> inline vfloat32m1_t madd(vfloat32m1_t a, vfloat32m1_t b, vfloat32m1_t c) {
    return __riscv_vfmacc_vv_f32m1(c, a, b, __riscv_vsetvlmax_e32m1());
 }
@@ -277,7 +277,7 @@ inline float hsum(__m512 x) {
 }
 #endif // __AVX512F__

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
 inline float hsum(vfloat32m1_t x) {
    return __riscv_vfmv_f_s_f32m1_f32(
        __riscv_vfredusum_vs_f32m1_f32m1(x, __riscv_vfmv_v_f_f32m1(0, 1), __riscv_vsetvlmax_e32m1()));
@@ -458,7 +458,7 @@ template <> inline __m512bh load_partial(const ggml_bf16_t *p, int n) {
 #endif
 #endif

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
 template <> inline vfloat32m1_t load(const float *p) {
    return __riscv_vle32_v_f32m1(p, __riscv_vsetvlmax_e32m1());
 }
@@ -503,7 +503,7 @@ template <> inline vbfloat16m4_t load(const ggml_bf16_t *p) {
 }
 #endif

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
 template <typename T> T set_zero();

 template <> inline vfloat32m1_t set_zero() {
@@ -520,7 +520,7 @@ template <> inline vfloat32m8_t set_zero() {
 }
 #endif

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
 template <typename T> size_t vlmax() {
    if constexpr (std::is_same_v<T, vfloat32m1_t>) { return  __riscv_vsetvlmax_e32m1(); }
    else if constexpr (std::is_same_v<T, vfloat32m2_t>) { return  __riscv_vsetvlmax_e32m2(); }
@@ -732,7 +732,7 @@ class tinyBLAS {
    const int64_t ldc;
 };

-#if defined(__riscv_v_intrinsic)
+#if defined(__riscv_v)
 template <typename D, typename V, typename TA, typename TB, typename TC>
 class tinyBLAS_RVV {
  public:
@@ -3958,7 +3958,7 @@ bool llamafile_sgemm(const struct ggml_compute_params * params, int64_t m, int64
            params->ith, params->nth};
        tb.matmul(m, n);
        return true;
-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)
    #if LMUL == 1
        tinyBLAS_RVV<vfloat32m1_t, vfloat32m1_t, float, float, float> tb{ params,
            k, (const float *)A, lda,
@@ -10114,7 +10114,7 @@ static void ggml_compute_forward_ssm_scan_f32(
                        }

                        sumf = GGML_F32xt_REDUCE_ONE(sum);
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
                        // todo: RVV implementation
                        const int np = 0;
    #else
@@ -11778,7 +11778,7 @@ static void ggml_compute_forward_rwkv_wkv7_f32(
    int64_t h_stride_2d = head_size * head_size;

    #if defined(GGML_SIMD)
-        #if defined(__ARM_FEATURE_SVE) || defined(__riscv_v_intrinsic)
+        #if defined(__ARM_FEATURE_SVE) || defined(__riscv_v)
            // scalar Route to scalar implementation       //TODO: Write SVE code and RVV code
            for (int64_t t = 0; t < T; t++) {
                int64_t t_offset = t * t_stride;
@@ -5,7 +5,7 @@
 #include "simd-mappings.h"

 // TODO: add support for sizeless vector types
-#if defined(GGML_SIMD) && !defined(__ARM_FEATURE_SVE) && !defined(__riscv_v_intrinsic)
+#if defined(GGML_SIMD) && !defined(__ARM_FEATURE_SVE) && !defined(__riscv_v)

 // TODO: untested on avx512
 // These are in units of GGML_F32_EPR
@@ -149,7 +149,7 @@ static void simd_gemm(
        C += N;
    }
 }
-#elif defined(GGML_SIMD) && defined(__riscv_v_intrinsic)
+#elif defined(GGML_SIMD) && defined(__riscv_v)
 // RM accumulators + 1 B vector = RM + 1 <= 8  =>  RM <= 7
 // Microkernel: C[RM x vl] += A[RM x K] * B[K x N]
 template <int RM>
@@ -14,7 +14,9 @@
 #include <arm_neon.h>
 #endif

-#if defined(__riscv_v_intrinsic)
+// __riscv_v_intrinsic reports compiler API support and can be defined without
+// enabling V; use the ISA feature macro to select RVV code.
+#if defined(__riscv_v)
 #include <riscv_vector.h>
 #endif

@@ -1306,7 +1308,7 @@ static inline void __lzs_f16cx4_store(ggml_fp16_t * x, float32x4_t v_y) {
 #define GGML_BF16_FMA_HI(acc, x, y) \
    (acc) = GGML_F32x4_FMA((acc), GGML_BF16_TO_F32_HI(x), GGML_BF16_TO_F32_HI(y))

-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)

 // compatible with vlen >= 128

@@ -84,7 +84,7 @@ void ggml_vec_dot_f32(int n, float * GGML_RESTRICT s, size_t bs, const float * G
        }
        // reduce sum1,sum2 to sum1
        GGML_F32_VEC_REDUCE(sumf, sum1, sum2, sum3, sum4, sum5, sum6, sum7, sum8);
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
        int vl = __riscv_vsetvlmax_e32m8();
        vfloat32m1_t vs = __riscv_vfmv_v_f_f32m1(0.0f, 1);
        vfloat32m8_t vsum;
@@ -195,7 +195,7 @@ void ggml_vec_dot_bf16(int n, float * GGML_RESTRICT s, size_t bs, ggml_bf16_t *
    sumf += (ggml_float)_mm_cvtss_f32(g);

 #undef LOAD
-#elif defined(__riscv_v_intrinsic) && defined(__riscv_zvfbfwma)
+#elif defined(__riscv_v) && defined(__riscv_zvfbfwma)
    size_t vl = __riscv_vsetvlmax_e32m4();

    // initialize accumulators to all zeroes
@@ -318,7 +318,7 @@ void ggml_vec_dot_f16(int n, float * GGML_RESTRICT s, size_t bs, ggml_fp16_t * G
        sum1_hi = svadd_f32_m(DEFAULT_PG32, sum1_hi, sum3_hi);

        sumf = ggml_sve_sum_f32x2(sum1_lo, sum1_hi);
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
        #if defined(__riscv_zvfh)
            int vl = __riscv_vsetvlmax_e32m2();
            vfloat32m1_t vs = __riscv_vfmv_v_f_f32m1(0.0f, 1);
@@ -405,7 +405,7 @@ void ggml_vec_silu_f32(const int n, float * y, const float * x) {
    for (; i + 3 < n; i += 4) {
        vec_xst(ggml_v_silu(vec_xl(0, x + i)), 0, y + i);
    }
-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)
    for (int vl; i < n; i += vl) {
        vl = __riscv_vsetvl_e32m2(n - i);
        vfloat32m2_t vx = __riscv_vle32_v_f32m2(&x[i], vl);
@@ -446,7 +446,7 @@ void ggml_vec_swiglu_f32(const int n, float * y, const float * x, const float *
    for (; i + 3 < n; i += 4) {
        vec_xst(vec_mul(ggml_v_silu(vec_xl(0, x + i)), vec_xl(0, g + i)), 0, y + i);
    }
-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)
    for (int vl; i < n; i += vl) {
        vl = __riscv_vsetvl_e32m2(n - i);
        vfloat32m2_t vx = __riscv_vle32_v_f32m2(&x[i], vl);
@@ -602,7 +602,7 @@ ggml_float ggml_vec_cvar_f32(const int n, float * y, const float * x, const floa
        val = vec_mul(val, val);
        sum += (ggml_float)vec_hsum_f32x4(val);
    }
-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)
    vfloat64m1_t vsum = __riscv_vfmv_v_f_f64m1(0, 1);
    for (int vl; i < n; i += vl) {
        vl = __riscv_vsetvl_e32m2(n - i);
@@ -681,7 +681,7 @@ ggml_float ggml_vec_soft_max_f32(const int n, float * y, const float * x, float
        vec_xst(v_val, 0, y + i);
        sum += (ggml_float)vec_hsum_f32x4(v_val);
    }
-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)
    vfloat64m1_t vsum = __riscv_vfmv_v_f_f64m1(0, 1);
    for (int avl; i < n; i += avl) {
        avl = __riscv_vsetvl_e32m2(n - i);
@@ -206,7 +206,7 @@ inline static void ggml_vec_dot_f16_unroll(const int n, const int xs, float * GG
        sumf[0] = ggml_sve_sum_f32x2(sum_0_lo, sum_0_hi);
        sumf[1] = ggml_sve_sum_f32x2(sum_1_lo, sum_1_hi);
        np = n;
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
        #if defined(__riscv_zvfh)
            size_t vl = __riscv_vsetvlmax_e32m4();

@@ -396,7 +396,7 @@ inline static void ggml_vec_mad_f32(const int n, float * GGML_RESTRICT y, const

            svst1_f32(pg, y + np2, ay1);
        }
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
        for (int i = 0, avl; i < n; i += avl) {
            avl = __riscv_vsetvl_e32m8(n - i);
            vfloat32m8_t ax = __riscv_vle32_v_f32m8(&x[i], avl);
@@ -513,7 +513,7 @@ inline static void ggml_vec_mad_f16(const int n, ggml_fp16_t * GGML_RESTRICT y,
        svst1_f16(pg, (__fp16 *)(y + np2), hy);
    }
    np = n;
-#elif defined(__riscv_v_intrinsic) // implies __riscv_v_intrinsic
+#elif defined(__riscv_v)
    #if defined (__riscv_zvfh)
        const ggml_fp16_t s = GGML_CPU_FP32_TO_FP16(v);
        const _Float16 scale = *(const _Float16*)(&s);
@@ -599,7 +599,7 @@ inline static void ggml_vec_mad_f32_unroll(const int n, const int xs, const int
                y[i] += x[k][i]*v[k][0];
            }
        }
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
        for (int i = 0, avl; i < n; i += avl) {
            avl = __riscv_vsetvl_e32m8(n - i);
            vfloat32m8_t ay = __riscv_vle32_v_f32m8(&y[i], avl);
@@ -660,7 +660,7 @@ inline static void ggml_vec_mad1_f32(const int n, float * y, const float * x, co
        for (int i = 0; i < n; ++i) {
            y[i] = x[i]*s + b;
        }
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
        for (int i = 0, avl; i < n; i += avl) {
            avl = __riscv_vsetvl_e32m8(n - i);
            vfloat32m8_t ax = __riscv_vle32_v_f32m8(&x[i], avl);
@@ -729,7 +729,7 @@ inline static void ggml_vec_scale_f32(const int n, float * y, const float   v) {
            ay1 = svmul_f32_m(pg, ay1, vx);
            svst1_f32(pg, y + i, ay1);
        }
-    #elif defined(__riscv_v_intrinsic)
+    #elif defined(__riscv_v)
        for (int i = 0, avl; i < n; i += avl) {
            avl = __riscv_vsetvl_e32m8(n - i);
            vfloat32m8_t ay = __riscv_vle32_v_f32m8(&y[i], avl);
@@ -793,7 +793,7 @@ inline static void ggml_vec_scale_f16(const int n, ggml_fp16_t * y, const float
        svst1_f16(pg, (__fp16 *)(y + np), out);
    }
    np = n;
-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)
    #if defined(__riscv_zvfh)
        const ggml_fp16_t s = GGML_CPU_FP32_TO_FP16(v);
        const _Float16 scale = *(const _Float16*)(&s);
@@ -1411,7 +1411,7 @@ inline static float32x4_t ggml_v_gelu_quick(float32x4_t x) {
    return vec_div(x, vec_add(one, exp_cx));
 }

-#elif defined(__riscv_v_intrinsic)
+#elif defined(__riscv_v)

 // adapted from arm limited optimized routine
 // the maximum error is 1.45358 plus 0.5 ulps
@@ -1469,7 +1469,7 @@ inline static vfloat32m2_t ggml_v_gelu_quick_m2(vfloat32m2_t x, int vl) {
    return __riscv_vfdiv_vv_f32m2(x, __riscv_vfadd_vf_f32m2(exp_cx, 1.0f, vl), vl);
 }

-#endif // __ARM_NEON / __AVX2__ / __SSE2__ / __VXE__ / __riscv_v_intrinsic
+#endif // __ARM_NEON / __AVX2__ / __SSE2__ / __VXE__ / __riscv_v

 inline static void ggml_vec_silu_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
    for (int i = 0; i < n; ++i) {
@@ -135,6 +135,7 @@ size_t opt_mbuf_total  = 0;                        // total buffer space limit (
 int    opt_mm_select     = 2; // 2 = HMX -> HVX -> CPU, 1 = HVX -> CPU, 0 = CPU (unsupported)
 int    opt_fa_select     = 2; // 2 = HMX -> HVX -> CPU, 1 = HVX -> CPU, 0 = CPU (unsupported)
 int    opt_fa_split      = 1; // 1 = head-split + team-split when n_kv_heads < N_cores (default), 0 = off
+int    opt_li_select     = 2; // 2 = HMX -> HVX, 1 = HVX, 0 = CPU (unsupported)
 int    opt_gdn_select    = 2; // 2 = HMX -> HVX, 1 = HVX, 0 = CPU (unsupported)
 int    opt_ar_select     = 2; // 2 = fused ALLREDUCE+ADD (default), 1 = unfused ALLREDUCE, 0 = fallback to CPY+FENCE
 int    opt_ar_scatter    = 1; // 1 = reduce-scatter the fused ALLREDUCE+ADD (default), 0 = full reduction
@@ -3409,6 +3410,7 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) {
    const char * str_mm_select = getenv("GGML_HEXAGON_MM_SELECT");
    const char * str_fa_select = getenv("GGML_HEXAGON_FA_SELECT");
    const char * str_fa_split  = getenv("GGML_HEXAGON_FA_SPLIT");
+    const char * str_li_select = getenv("GGML_HEXAGON_LI_SELECT");
    const char * str_gdn_select = getenv("GGML_HEXAGON_GDN_SELECT");
    const char * str_ar_select = getenv("GGML_HEXAGON_AR_SELECT");
    const char * str_ar_scatter = getenv("GGML_HEXAGON_AR_SCATTER");
@@ -3475,6 +3477,7 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) {
    opt_mm_select = str_mm_select ? atoi(str_mm_select)                   : opt_mm_select;
    opt_fa_select = str_fa_select ? atoi(str_fa_select)                   : opt_fa_select;
    opt_fa_split  = str_fa_split  ? atoi(str_fa_split)                   : opt_fa_split;
+    opt_li_select = str_li_select ? atoi(str_li_select)                   : opt_li_select;
    opt_gdn_select    = str_gdn_select    ? atoi(str_gdn_select)          : opt_gdn_select;
    opt_ar_select     = str_ar_select     ? atoi(str_ar_select)           : opt_ar_select;
    opt_ar_scatter    = str_ar_scatter    ? atoi(str_ar_scatter)          : opt_ar_scatter;
@@ -29,6 +29,7 @@ extern size_t opt_mbuf_total;
 extern int    opt_mm_select;
 extern int    opt_fa_select;
 extern int    opt_fa_split;
+extern int    opt_li_select;
 extern int    opt_gdn_select;
 extern int    opt_ar_select;
 extern int    opt_ar_scatter;
@@ -64,6 +64,7 @@ htp_op_code htp_op_remap(const struct ggml_tensor * t) {
        case GGML_OP_CONV_2D_DW:      return HTP_OP_CONV_2D_DW;
        case GGML_OP_ROLL:            return HTP_OP_ROLL;
        case GGML_OP_UPSCALE:         return HTP_OP_UPSCALE;
+        case GGML_OP_LIGHTNING_INDEXER: return HTP_OP_LIGHTNING_INDEXER;

        case GGML_OP_UNARY:
            switch (ggml_get_unary_op(t)) {
@@ -4708,6 +4709,233 @@ static inline bool htp_kparams_solve_dsv4_hc_post(
    return htp_kparams_solve_dsv4_hc_post_impl(sess, op->src[0], op->src[1], op->src[2], op->src[3], op, kparams);
 }

+static bool htp_lightning_indexer_can_collapse(
+    const struct ggml_tensor * q,
+    const struct ggml_tensor * k,
+    const struct ggml_tensor * w,
+    const struct ggml_tensor * m,
+    const struct ggml_tensor * dst
+) {
+    if (!q || !k || !w || !m || !dst) {
+        return false;
+    }
+
+    // 2D K normalization: ne[1] is n_kv, ne[2] is 1
+    const bool can_collapse_k = (k->ne[1] > 1 && k->ne[2] == 1);
+
+    // Stream-to-token collapse across q, w, dst (and m)
+    bool can_collapse_stream = false;
+    if (q->ne[3] > 1 && k->ne[3] == 1) {
+        const bool q_ok = (q->nb[3] == (size_t) q->ne[2] * q->nb[2]);
+        const bool dst_ok = (dst->ne[3] == q->ne[3] && dst->ne[2] == 1 && dst->ne[1] == q->ne[2] &&
+                             dst->nb[3] == (size_t) dst->ne[1] * dst->nb[1]);
+        const bool w_ok = (w->ne[3] == q->ne[3] && w->ne[2] == 1 && w->ne[1] == q->ne[2] &&
+                           w->nb[3] == (size_t) w->ne[1] * w->nb[1]) ||
+                          (w->ne[3] == 1 && w->ne[1] == q->ne[2] * q->ne[3]);
+        const bool m_ok = (m->ne[3] == q->ne[3] && m->ne[2] == 1 && m->ne[1] == q->ne[2] &&
+                           m->nb[3] == (size_t) m->ne[1] * m->nb[1]) ||
+                          (m->ne[3] == 1 && m->ne[1] == q->ne[2] * q->ne[3]);
+        can_collapse_stream = q_ok && dst_ok && w_ok && m_ok;
+    }
+
+    return can_collapse_k || can_collapse_stream;
+}
+
+static void htp_lightning_indexer_collapse_tensors(
+    const struct ggml_tensor * q,
+    const struct ggml_tensor * k,
+    const struct ggml_tensor * w,
+    const struct ggml_tensor * m,
+    const struct ggml_tensor * dst,
+    struct ggml_tensor * q_c,
+    struct ggml_tensor * k_c,
+    struct ggml_tensor * w_c,
+    struct ggml_tensor * m_c,
+    struct ggml_tensor * dst_c
+) {
+    *q_c   = *q;
+    *k_c   = *k;
+    *w_c   = *w;
+    *m_c   = *m;
+    *dst_c = *dst;
+
+    if (k->ne[1] > 1 && k->ne[2] == 1) {
+        k_c->ne[2] = k->ne[1];
+        k_c->ne[1] = 1;
+        k_c->nb[1] = k->nb[1];
+        k_c->nb[2] = k->nb[1];
+        k_c->nb[3] = k->nb[2];
+    }
+
+    if (q->ne[3] > 1 && k->ne[3] == 1) {
+        const int64_t n_tokens_collapsed = (int64_t) q->ne[2] * q->ne[3];
+
+        q_c->ne[2] = n_tokens_collapsed;
+        q_c->ne[3] = 1;
+        q_c->nb[3] = (size_t) q_c->ne[2] * q_c->nb[2];
+
+        if (w->ne[3] == q->ne[3]) {
+            w_c->ne[1] = n_tokens_collapsed;
+            w_c->ne[3] = 1;
+            w_c->nb[2] = (size_t) w_c->ne[1] * w_c->nb[1];
+            w_c->nb[3] = w_c->nb[2];
+        }
+
+        dst_c->ne[1] = n_tokens_collapsed;
+        dst_c->ne[3] = 1;
+        dst_c->nb[2] = (size_t) dst_c->ne[1] * dst_c->nb[1];
+        dst_c->nb[3] = dst_c->nb[2];
+
+        if (m->ne[3] == q->ne[3] && m->ne[1] == q->ne[2]) {
+            m_c->ne[1] = n_tokens_collapsed;
+            m_c->ne[3] = 1;
+            m_c->nb[2] = (size_t) m_c->ne[1] * m_c->nb[1];
+            m_c->nb[3] = m_c->nb[2];
+        }
+    }
+}
+
+static bool htp_kparams_solve_lightning_indexer_impl(
+    const struct ggml_hexagon_session * sess,
+    const struct ggml_tensor * q,
+    const struct ggml_tensor * k,
+    const struct ggml_tensor * w,
+    const struct ggml_tensor * m,
+    const struct ggml_tensor * op,
+    struct htp_lid_kernel_params * kparams
+) {
+    if (opt_li_select < 1) {
+        return false;
+    }
+
+    if (k->type != GGML_TYPE_F32 && k->type != GGML_TYPE_F16 && k->type != GGML_TYPE_Q8_0 && k->type != GGML_TYPE_Q4_0) {
+        return false;
+    }
+
+    if (q->ne[0] % 32 != 0 ||
+        !ggml_is_contiguous_rows(q) || !ggml_is_contiguous_rows(k) ||
+        !ggml_is_contiguous_rows(w) || !ggml_is_contiguous_rows(m) ||
+        !ggml_is_contiguous_rows(op)) {
+        return false;
+    }
+
+    const uint32_t n_kv_blocks = (k->ne[2] + HTP_LID_KV_BLOCK - 1) / HTP_LID_KV_BLOCK;
+    if (n_kv_blocks == 0 || q->ne[2] == 0) {
+        return false;
+    }
+
+    const uint32_t D          = q->ne[0];
+    const uint32_t n_head     = q->ne[1];
+    const uint32_t n_tokens   = q->ne[2];
+    const size_t   k_row_size = ggml_row_size(k->type, D);
+
+    kparams->k_row_size = (uint32_t) k_row_size;
+    kparams->div_n_head = init_fastdiv_values(n_head);
+    kparams->div_m_ne3  = init_fastdiv_values((uint32_t) m->ne[3]);
+
+    if (opt_li_select >= 2 && htp_session_n_hmx(sess) > 0 && D % HMX_FP16_TILE_N_COLS == 0) {
+        const uint32_t n_threads = htp_session_n_threads(sess);
+        const uint32_t n_stream  = q->ne[3];
+        const uint32_t n_dot     = D / HMX_FP16_TILE_N_COLS;
+
+        auto vtcm_usage = [&](uint32_t n_tok, uint32_t kv_block, uint32_t n_qsets) {
+            return htp_lid_hmx_compute_vtcm_usage(D, n_head, n_tok, kv_block, n_qsets, n_threads, k_row_size);
+        };
+
+        uint32_t n_tok = (std::min)(n_tokens, (uint32_t) HTP_LID_HMX_TOKENS_MAX);
+        while (n_tok > 32 && vtcm_usage(n_tok, 64, 1) > htp_session_vtcm_size(sess)) {
+            n_tok /= 2;
+        }
+
+        uint32_t n_qsets = n_tok == n_tokens ? (std::min)(n_stream, (uint32_t) HTP_LID_HMX_MAX_QSETS) : 1;
+        while (n_qsets > 1 && vtcm_usage(n_tok, 64, n_qsets) > htp_session_vtcm_size(sess)) {
+            n_qsets /= 2;
+        }
+
+        uint32_t n_qbuf = 1;
+        if (n_qsets == 1 && n_head >= 16 && n_stream * ((n_tokens + n_tok - 1) / n_tok) > 1) {
+            uint32_t nt2 = n_tok;
+            while (nt2 > 32 && vtcm_usage(nt2, 64, 2) > htp_session_vtcm_size(sess)) {
+                nt2 /= 2;
+            }
+            const size_t q_bytes = (size_t) nt2 * n_head * D * sizeof(float);
+            const size_t k_bytes = (size_t) k->ne[2] * ggml_row_size(k->type, D);
+            if (vtcm_usage(nt2, 64, 2) <= htp_session_vtcm_size(sess) && q_bytes >= k_bytes) {
+                n_tok  = nt2;
+                n_qbuf = 2;
+            }
+        }
+
+        const uint32_t n_ch   = (n_tok * n_head + 31) / 32;
+        const uint32_t kv_seg = hex_round_up(k->ne[2], 64);
+        const uint32_t n_rows = n_qsets * kv_seg;
+        const uint32_t bc_max = n_rows <= 512 ? n_rows : (std::max)(256u, n_rows / 4);
+        uint32_t kv_block = 64;
+        while (kv_block < HTP_LID_HMX_KV_BLOCK_MAX && kv_block < bc_max && (kv_block / 32) * n_ch * (n_dot + 1) < 256 &&
+               vtcm_usage(n_tok, kv_block * 2, n_qsets * n_qbuf) <= htp_session_vtcm_size(sess)) {
+            kv_block *= 2;
+        }
+
+        const size_t vtcm_size = vtcm_usage(n_tok, kv_block, n_qsets * n_qbuf);
+        if (vtcm_size <= htp_session_vtcm_size(sess)) {
+            const uint32_t n_tt = (n_tokens + n_tok - 1) / n_tok;
+            kparams->kernel_type       = HTP_LID_KERNEL_HMX;
+            kparams->n_threads         = n_threads;
+            kparams->n_tokens_tile     = n_tok;
+            kparams->n_kv_blocks       = (n_rows + kv_block - 1) / kv_block;
+            kparams->blocks_per_thread = 0;
+            kparams->kv_block          = kv_block;
+            kparams->n_qsets           = n_qsets;
+            kparams->n_qbuf            = n_qbuf;
+            kparams->kv_seg            = kv_seg;
+            kparams->n_tt              = n_tt;
+            kparams->vtcm_size         = (uint32_t) vtcm_size;
+            kparams->div_kv_seg        = init_fastdiv_values(kv_seg);
+            kparams->div_n_tt          = init_fastdiv_values(n_tt);
+            kparams->div_n_ch          = init_fastdiv_values(n_ch);
+            return true;
+        }
+    }
+
+    const uint32_t n_threads = (std::min)(htp_session_n_threads(sess), n_kv_blocks);
+    uint32_t n_tok = (std::min)((uint32_t) q->ne[2], (uint32_t) HTP_LID_TOKENS_MAX);
+    while (n_tok > 1 && htp_lid_compute_vtcm_usage(D, n_head, n_tok, n_threads, k_row_size) > htp_session_vtcm_size(sess)) {
+        n_tok /= 2;
+    }
+
+    const size_t vtcm_size = htp_lid_compute_vtcm_usage(D, n_head, n_tok, n_threads, k_row_size);
+    if (vtcm_size > htp_session_vtcm_size(sess)) {
+        return false;
+    }
+
+    kparams->kernel_type       = HTP_LID_KERNEL_HVX;
+    kparams->n_threads         = n_threads;
+    kparams->n_tokens_tile     = n_tok;
+    kparams->n_kv_blocks       = n_kv_blocks;
+    kparams->blocks_per_thread = (n_kv_blocks + n_threads - 1) / n_threads;
+    kparams->kv_block          = HTP_LID_KV_BLOCK;
+    kparams->n_qsets           = 1;
+    kparams->n_qbuf            = 1;
+    kparams->kv_seg            = 0;
+    kparams->n_tt              = 0;
+    kparams->vtcm_size         = (uint32_t) vtcm_size;
+
+    return true;
+}
+
+static bool htp_kparams_solve_lightning_indexer(
+    const struct ggml_hexagon_session * sess,
+    const struct ggml_tensor * op,
+    struct htp_lid_kernel_params * kparams
+) {
+    if (htp_lightning_indexer_can_collapse(op->src[0], op->src[1], op->src[2], op->src[3], op)) {
+        struct ggml_tensor q_c, k_c, w_c, m_c, d_c;
+        htp_lightning_indexer_collapse_tensors(op->src[0], op->src[1], op->src[2], op->src[3], op,
+                                               &q_c, &k_c, &w_c, &m_c, &d_c);
+        return htp_kparams_solve_lightning_indexer_impl(sess, &q_c, &k_c, &w_c, &m_c, &d_c, kparams);
+    }
+    return htp_kparams_solve_lightning_indexer_impl(sess, op->src[0], op->src[1], op->src[2], op->src[3], op, kparams);
+}
 template <auto SolveFn>
 static inline bool htp_kparams_invoke(
    const struct ggml_hexagon_session * sess,
@@ -4861,6 +5089,28 @@ void htp_opnode_collapse(const struct ggml_hexagon_session * sess, htp_opnode &
        }
        return;
    }
+
+    if (node.opcode == HTP_OP_LIGHTNING_INDEXER) {
+        if (node.inputs.size() < 4 || !node.inputs[0] || !node.inputs[1] || !node.inputs[2] || !node.inputs[3] || node.outputs.empty() || !node.outputs[0]) {
+            return;
+        }
+        const struct ggml_tensor * q   = node.inputs[0];
+        const struct ggml_tensor * k   = node.inputs[1];
+        const struct ggml_tensor * w   = node.inputs[2];
+        const struct ggml_tensor * m   = node.inputs[3];
+        const struct ggml_tensor * dst = node.outputs[0];
+
+        if (htp_lightning_indexer_can_collapse(q, k, w, m, dst)) {
+            struct ggml_tensor q_c, k_c, w_c, m_c, d_c;
+            htp_lightning_indexer_collapse_tensors(q, k, w, m, dst, &q_c, &k_c, &w_c, &m_c, &d_c);
+            node.inputs[0]  = node.add_dummy(q_c);
+            node.inputs[1]  = node.add_dummy(k_c);
+            node.inputs[2]  = node.add_dummy(w_c);
+            node.inputs[3]  = node.add_dummy(m_c);
+            node.outputs[0] = node.add_dummy(d_c);
+        }
+        return;
+    }
 }

 bool htp_kparams_solve(
@@ -4901,6 +5151,11 @@ bool htp_kparams_solve(
        const ggml_tensor * src3 = node.inputs.size() > 3 ? node.inputs[3] : nullptr;
        return htp_kparams_solve_dsv4_hc_post_impl(sess, src0, src1, src2, src3, dst, node.kparams<htp_dsv4_hc_post_kernel_params>());
    }
+    if (node.opcode == HTP_OP_LIGHTNING_INDEXER) {
+        const ggml_tensor * src2 = node.inputs.size() > 2 ? node.inputs[2] : nullptr;
+        const ggml_tensor * src3 = node.inputs.size() > 3 ? node.inputs[3] : nullptr;
+        return htp_kparams_solve_lightning_indexer_impl(sess, src0, src1, src2, src3, dst, node.kparams<htp_lid_kernel_params>());
+    }

    return htp_kparams_solve(sess, node.node, node.kernel_params);
 }
@@ -5005,6 +5260,9 @@ bool htp_kparams_solve(
        case HTP_OP_ROLL:
            return htp_kparams_invoke<htp_kparams_solve_roll>(sess, op, kparams);

+        case HTP_OP_LIGHTNING_INDEXER:
+            return htp_kparams_invoke<htp_kparams_solve_lightning_indexer>(sess, op, kparams);
+
        case HTP_OP_DIAG:
            return htp_kparams_invoke<htp_kparams_solve_diag>(sess, op, kparams);

@@ -18,6 +18,7 @@
 #include "htp/get-rows-ops.h"
 #include "htp/glu-ops.h"
 #include "htp/im2col-ops.h"
+#include "htp/lightning-indexer-ops.h"
 #include "htp/matmul-ops.h"
 #include "htp/pad-ops.h"
 #include "htp/pool-ops.h"
@@ -11,6 +11,7 @@
 #include "htp/gated-delta-net-ops.h"
 #include "htp/get-rows-ops.h"
 #include "htp/glu-ops.h"
+#include "htp/lightning-indexer-ops.h"
 #include "htp/matmul-ops.h"
 #include "htp/pad-ops.h"
 #include "htp/repeat-ops.h"
@@ -197,6 +198,12 @@ static std::string format_kernel_params(const htp_opnode & node) {
        const char * ktype_str = kparams->kernel_type == HTP_DSV4_HC_POST_KERNEL_COMB ? "comb" : "identity";
        snprintf(str, max_size, "%s nth %u vtcm %u", ktype_str,
                 (unsigned int) kparams->n_threads, (unsigned int) kparams->vtcm_size);
+    } else if (node.opcode == HTP_OP_LIGHTNING_INDEXER) {
+        const auto * kparams = node.kparams<htp_lid_kernel_params>();
+        snprintf(str, max_size, "%s tok %u bc %u nth %u vtcm %u",
+                 kparams->kernel_type == HTP_LID_KERNEL_HMX ? "hmx" : "hvx",
+                 (unsigned int) kparams->n_tokens_tile, (unsigned int) kparams->kv_block,
+                 (unsigned int) kparams->n_threads, (unsigned int) kparams->vtcm_size);
    } else if (node.opcode == HTP_OP_MUL || node.opcode == HTP_OP_ADD || node.opcode == HTP_OP_ADD_ID ||
               node.opcode == HTP_OP_SUB || node.opcode == HTP_OP_DIV) {
        const auto * kparams = node.kparams<htp_binary_kernel_params>();
@@ -51,6 +51,7 @@ add_library(${HTP_LIB} SHARED
    upscale-ops.c
    allreduce-ops.c
    dsv4-hc-ops.c
+    lightning-indexer-ops.c
 )

 target_compile_definitions(${HTP_LIB} PRIVATE
@@ -203,5 +203,6 @@ int op_pool_1d(struct htp_ops_context * octx);
 int op_allreduce(struct htp_ops_context * octx);
 int op_roll(struct htp_ops_context * octx);
 int op_upscale(struct htp_ops_context * octx);
+int op_lightning_indexer(struct htp_ops_context * octx);

 #endif /* HTP_CTX_H */
@@ -120,6 +120,7 @@ enum htp_op_code {
    HTP_OP_DSV4_HC_PRE,
    HTP_OP_DSV4_HC_POST,
    HTP_OP_TOP_K_MASK,
+    HTP_OP_LIGHTNING_INDEXER,

    HTP_OP_INVALID
 };
@@ -0,0 +1,213 @@
+#ifndef HTP_LIGHTNING_INDEXER_OPS_H
+#define HTP_LIGHTNING_INDEXER_OPS_H
+
+#include <stdint.h>
+#include <stddef.h>
+#include <stdbool.h>
+
+#include "hex-common.h"
+#include "hex-fastdiv.h"
+#include "htp-vtcm.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#define HTP_LID_KV_BLOCK   64 // kv rows per block (HVX and HMX Bc)
+#define HTP_LID_TOKENS_MAX 32 // max tokens per Q tile (HVX splat path)
+#define HTP_LID_HMX_TOKENS_MAX 512 // HMX can keep more tokens; VTCM decides
+#define HTP_LID_HMX_KV_BLOCK_MAX 2048
+#define HTP_LID_HMX_MAX_QSETS    8    // streams in one HMX pipeline run
+#define HTP_LID_N_BUF      4  // K / mask / output buffers per thread (two blocks computed, two in flight)
+
+#ifndef HMX_FP16_TILE_N_COLS
+#define HMX_FP16_TILE_N_COLS 32
+#endif
+#ifndef HMX_FP16_TILE_SIZE
+#define HMX_FP16_TILE_SIZE 2048
+#endif
+
+enum htp_lid_kernel_type {
+    HTP_LID_KERNEL_HVX = 0,
+    HTP_LID_KERNEL_HMX = 1,
+};
+
+struct htp_lid_kernel_params {
+    uint8_t  kernel_type;       // enum htp_lid_kernel_type
+    uint32_t n_threads;
+    uint32_t n_tokens_tile;     // tokens per Q tile
+    uint32_t n_kv_blocks;       // HVX: per stream, HMX: per run of n_qsets sets
+    uint32_t blocks_per_thread;
+    uint32_t vtcm_size;
+    uint32_t k_row_size;        // K row bytes
+    uint32_t kv_block;          // HMX: kv rows per block (multiple of 64)
+    uint32_t n_qsets;           // HMX: Q sets kept in VTCM (streams run in one pipeline)
+    uint32_t n_qbuf;            // HMX: Q buffers, 2: the next run's Q is prepped during the current run
+    uint32_t kv_seg;            // HMX: rows per set, n_kv padded to 64
+    uint32_t n_tt;              // HMX: token tiles per stream
+
+    struct fastdiv_values div_n_head;
+    struct fastdiv_values div_m_ne3;  // mask streams (broadcast)
+    struct fastdiv_values div_kv_seg; // HMX
+    struct fastdiv_values div_n_tt;   // HMX
+    struct fastdiv_values div_n_ch;   // HMX: channel tiles per Q set
+};
+
+#if defined(__cplusplus)
+static_assert(sizeof(struct htp_lid_kernel_params) <= 128, "htp_lid_kernel_params is too large");
+#else
+_Static_assert(sizeof(struct htp_lid_kernel_params) <= 128, "htp_lid_kernel_params is too large");
+#endif
+
+// VTCM layout shared by host (fit check) and device (buffer offsets).
+// Shared regions: Q splats (one f16 vector per q element) and head weight splats for the current token tile.
+// Per-thread regions: Q prep staging, K / mask / output buffers, two transposed K blocks and transpose scratch.
+struct htp_lid_vtcm_layout {
+    size_t off_q;
+    size_t off_w;
+    size_t off_qs;
+    size_t off_k;
+    size_t off_kt;
+    size_t off_m;
+    size_t off_o;
+
+    size_t size_q_row;   // f16 Q row, padded
+    size_t size_qs_row;  // f32 Q staging row, padded
+    size_t size_k_row;   // f16 K row, padded
+    size_t size_k_raw_row;
+    size_t size_m_row;   // f16 mask row of one kv block, padded
+    size_t size_o_row;   // f32 output row of one kv block, padded
+
+    size_t size_qs;      // per thread
+    size_t size_k_buf;   // one K buffer
+    size_t size_kt;      // per thread: two transposed K blocks, two 64x64 scratch tiles, Q8_0 block scales
+    size_t size_m_buf;   // one mask buffer
+    size_t size_o_buf;   // one output buffer
+
+    size_t total_bytes;
+};
+
+static inline void htp_lid_vtcm_layout_build(struct htp_lid_vtcm_layout * L, size_t D, size_t n_head, size_t n_tok, size_t n_threads, size_t k_row_size) {
+    L->size_q_row  = hex_round_up(D * sizeof(uint16_t), 128);
+    L->size_qs_row = hex_round_up(D * sizeof(float), 128);
+    L->size_k_row  = hex_round_up(D * sizeof(uint16_t), 128);
+    L->size_k_raw_row = L->size_k_row > hex_round_up(k_row_size, 128) ? L->size_k_row : hex_round_up(k_row_size, 128);
+    L->size_m_row  = hex_round_up(HTP_LID_KV_BLOCK * sizeof(uint16_t), 128);
+    L->size_o_row  = hex_round_up(HTP_LID_KV_BLOCK * sizeof(float), 128);
+
+    L->size_qs    = 2 * (n_head * L->size_qs_row + hex_round_up(n_head * sizeof(float), 128)) + L->size_q_row;
+    L->size_k_buf = HTP_LID_KV_BLOCK * L->size_k_raw_row;
+    L->size_kt    = (2 * hex_round_up(D, HTP_LID_KV_BLOCK) + 2 * HTP_LID_KV_BLOCK + D / 32) * 128;
+    L->size_m_buf = n_tok * L->size_m_row;
+    L->size_o_buf = n_tok * L->size_o_row;
+
+    size_t off = 0;
+    VTCM_LAYOUT_ALLOC(off, off_q,  n_tok * n_head * D * 128);
+    VTCM_LAYOUT_ALLOC(off, off_w,  n_tok * n_head * 128);
+    VTCM_LAYOUT_ALLOC(off, off_qs, n_threads * L->size_qs);
+    VTCM_LAYOUT_ALLOC(off, off_k,  n_threads * HTP_LID_N_BUF * L->size_k_buf);
+    VTCM_LAYOUT_ALLOC(off, off_kt, n_threads * L->size_kt);
+    VTCM_LAYOUT_ALLOC(off, off_m,  n_threads * HTP_LID_N_BUF * L->size_m_buf);
+    VTCM_LAYOUT_ALLOC(off, off_o,  n_threads * HTP_LID_N_BUF * L->size_o_buf);
+
+    L->total_bytes = off;
+}
+
+static inline size_t htp_lid_compute_vtcm_usage(size_t D, size_t n_head, size_t n_tok, size_t n_threads, size_t k_row_size) {
+    struct htp_lid_vtcm_layout L;
+    htp_lid_vtcm_layout_build(&L, D, n_head, n_tok, n_threads, k_row_size);
+    return L.total_bytes;
+}
+
+// HMX path. K is the activation (kv rows), Q the weights (one output channel per (token, head)).
+// mm1: S = relu(K Q^T) * |w| per channel (HMX convert scale + pos clamp)
+// mm2: O = S G, G[(t,h), t] = sign(w[t,h]). O is [kv, token], f16.
+// Q weights, G, scales and w are per Q set (one per stream when all tokens fit one tile).
+struct htp_lid_hmx_vtcm_layout {
+    size_t off_qw;       // Q weight tiles [n_qsets][n_ch][n_dot]
+    size_t off_g;        // G weight tiles [n_qsets][n_ch]
+    size_t off_sc;       // per channel scales, 256 bytes per channel tile, [n_qsets][n_ch]
+    size_t off_sc_id;    // identity scales
+    size_t off_s;        // S tiles of one block and token tile, [n_rt][n_sg]
+    size_t off_w;        // w of the token tile, [n_qsets][n_tok][n_head] f32
+    size_t off_qs;       // per thread: f32 and f16 staging for 32 Q rows, gather scratch
+    size_t off_k_raw[2]; // K rows, [Bc][size_k_row]
+    size_t off_k_t[2];   // K tiles [n_rt][n_dot]
+    size_t off_o[2];     // O tiles [n_rt][n_u]
+    size_t off_m[4];     // mask rows [n_tok][size_m_row]
+    size_t off_out[2];   // f32 output rows [n_tok][size_o_row]
+
+    size_t size_qs_row;  // f32 Q row, padded
+    size_t size_qh_row;  // f16 Q row, padded
+    size_t size_k_row;   // f16 K row, padded
+    size_t size_k_raw_row;
+    size_t size_m_row;
+    size_t size_o_row;
+    size_t size_qs;      // per thread
+    size_t size_w;       // per Q set
+
+    uint32_t n_ch;       // channel tiles: ceil(n_tok * n_head / 32)
+    uint32_t n_u;        // token tiles: ceil(n_tok / 32)
+    uint32_t n_dot;      // D / 32
+    uint32_t n_rt;       // kv row tiles per block: Bc / 32
+    uint32_t n_sg;       // S tiles per (row tile, token tile): min(n_head, n_ch)
+
+    size_t total_bytes;
+};
+
+static inline void htp_lid_hmx_vtcm_layout_build(struct htp_lid_hmx_vtcm_layout * L,
+                                                  size_t D, size_t n_head, size_t n_tok, size_t Bc, size_t n_qsets, size_t n_threads, size_t k_row_size) {
+    L->n_ch  = (uint32_t) ((n_tok * n_head + 31) / 32);
+    L->n_u   = (uint32_t) ((n_tok + 31) / 32);
+    L->n_dot = (uint32_t) (D / 32);
+    L->n_rt  = (uint32_t) (Bc / 32);
+    L->n_sg  = (uint32_t) (n_head < L->n_ch ? n_head : L->n_ch);
+
+    L->size_qs_row = hex_round_up(D * sizeof(float), 128);
+    L->size_qh_row = hex_round_up(D * sizeof(uint16_t), 128);
+    L->size_k_row  = hex_round_up(D * sizeof(uint16_t), 128);
+    L->size_k_raw_row = L->size_k_row > hex_round_up(k_row_size, 128) ? L->size_k_row : hex_round_up(k_row_size, 128);
+    L->size_m_row  = hex_round_up(Bc * sizeof(uint16_t), 128);
+    L->size_o_row  = hex_round_up(Bc * sizeof(float), 128);
+    L->size_qs     = 2 * 32 * L->size_qs_row + 32 * L->size_qh_row + hex_round_up(Bc / 64 * 128, 128);
+    L->size_w      = hex_round_up(n_tok * n_head * sizeof(float), 128);
+
+    const size_t tile = HMX_FP16_TILE_SIZE;
+
+    // tile regions first: they need 2 KB alignment
+    size_t off = 0;
+    VTCM_LAYOUT_ALLOC(off, off_qw,     n_qsets * L->n_ch * L->n_dot * tile);
+    VTCM_LAYOUT_ALLOC(off, off_g,      n_qsets * L->n_ch * tile);
+    VTCM_LAYOUT_ALLOC(off, off_s,      (size_t) L->n_rt * L->n_sg * tile);
+    VTCM_LAYOUT_ALLOC(off, off_k_t[0], (size_t) L->n_rt * L->n_dot * tile);
+    VTCM_LAYOUT_ALLOC(off, off_k_t[1], (size_t) L->n_rt * L->n_dot * tile);
+    VTCM_LAYOUT_ALLOC(off, off_o[0],   (size_t) L->n_rt * L->n_u * tile);
+    VTCM_LAYOUT_ALLOC(off, off_o[1],   (size_t) L->n_rt * L->n_u * tile);
+    VTCM_LAYOUT_ALLOC(off, off_sc,     n_qsets * L->n_ch * 256);
+    VTCM_LAYOUT_ALLOC(off, off_sc_id,  256);
+    VTCM_LAYOUT_ALLOC(off, off_w,      n_qsets * L->size_w);
+    VTCM_LAYOUT_ALLOC(off, off_qs,     n_threads * L->size_qs);
+    for (int i = 0; i < 2; i++) {
+        VTCM_LAYOUT_ALLOC(off, off_k_raw[i], Bc * L->size_k_raw_row);
+    }
+    for (int i = 0; i < 4; i++) {
+        VTCM_LAYOUT_ALLOC(off, off_m[i], n_tok * L->size_m_row);
+    }
+    for (int i = 0; i < 2; i++) {
+        VTCM_LAYOUT_ALLOC(off, off_out[i], n_tok * L->size_o_row);
+    }
+
+    L->total_bytes = off;
+}
+
+static inline size_t htp_lid_hmx_compute_vtcm_usage(size_t D, size_t n_head, size_t n_tok, size_t Bc, size_t n_qsets, size_t n_threads, size_t k_row_size) {
+    struct htp_lid_hmx_vtcm_layout L;
+    htp_lid_hmx_vtcm_layout_build(&L, D, n_head, n_tok, Bc, n_qsets, n_threads, k_row_size);
+    return L.total_bytes;
+}
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif /* HTP_LIGHTNING_INDEXER_OPS_H */
@@ -1056,6 +1056,9 @@ static int execute_op(struct htp_ops_context * octx) {
        case HTP_OP_GATED_DELTA_NET:
            return op_gated_delta_net(octx);

+        case HTP_OP_LIGHTNING_INDEXER:
+            return op_lightning_indexer(octx);
+
        case HTP_OP_TRI:
            return op_unary(octx);

@@ -887,6 +887,7 @@ struct ggml_backend_opencl_context {
    cl_kernel kernel_cpy_f32_f32_pack = nullptr;
    cl_kernel kernel_cpy_f32_f32_flat = nullptr;
    cl_kernel kernel_mul_mat_f32_f32;
+    cl_kernel kernel_mul_mv_f32_f32_r2 = nullptr;
    cl_kernel kernel_fwht[2] = {};
    size_t fwht_workgroup_size[2] = {};
    cl_kernel kernel_mul_mat_f16_f16;
@@ -1027,10 +1028,13 @@ struct ggml_backend_opencl_context {
    cl_kernel kernel_mul_mv_q2_0_f32_flat;
    cl_kernel kernel_mul_mm_q2_0_f32_l4_lm;
    cl_kernel kernel_gemv_noshuffle_q2_0_f32 = nullptr;
+    cl_kernel kernel_gemv_noshuffle_q2_0_f32_wide = nullptr;
+    size_t q2_0_mv_workgroup_size = 512;
+    size_t q2_0_mv_workgroup_size_wide = 512;
    cl_kernel kernel_gemm_noshuffle_q2_0_f32 = nullptr;
    cl_kernel kernel_gemm_noshuffle_q2_0_q8_1_dp4a = nullptr;
    cl_kernel kernel_gemm_noshuffle_q2_0_q8_1_dp4a_narrow = nullptr;
-    int q2_0_mv_nsg = 8;
+    int q2_0_mv_nsg = 32;
    cl_kernel kernel_mul_mv_q2_0_f32;
    cl_kernel kernel_convert_block_iq1_m, kernel_restore_block_iq1_m;
    cl_kernel kernel_mul_mv_iq1_m_f32_flat;
@@ -3591,6 +3595,17 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
            build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);

        CL_CHECK((backend_ctx->kernel_mul_mat_f32_f32 = clCreateKernel(backend_ctx->program_mul_mv_f32_f32, "kernel_mul_mat_f32_f32", &err), err));
+        if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::A8X) {
+            cl_kernel kernel = clCreateKernel(backend_ctx->program_mul_mv_f32_f32, "kernel_mul_mv_f32_f32_r2", &err);
+            CL_CHECK(err);
+            size_t max_workgroup;
+            CL_CHECK(clGetKernelWorkGroupInfo(kernel, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_workgroup), &max_workgroup, nullptr));
+            if (max_workgroup >= 64) {
+                backend_ctx->kernel_mul_mv_f32_f32_r2 = kernel;
+            } else {
+                CL_CHECK(clReleaseKernel(kernel));
+            }
+        }
        GGML_LOG_CONT(".");
    }

@@ -5076,25 +5091,30 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
 #endif
        if (const char * env = getenv("GGML_OPENCL_Q2_0_MV_NSG")) {
            const int nsg = atoi(env);
-            if (nsg == 1 || nsg == 2 || nsg == 4 || nsg == 8) {
+            if (nsg == 1 || nsg == 2 || nsg == 4 || nsg == 8 || nsg == 16 || nsg == 32) {
                backend_ctx->q2_0_mv_nsg = nsg;
            }
        }
-        for (;;) {
-            const std::string opts = compile_opts + " -DQ2_0_MV_NSG=" + std::to_string(backend_ctx->q2_0_mv_nsg);
+        for (int wide = 0; wide < 2; ++wide) {
+            const std::string opts = compile_opts + " -DQ2_0_MV_WIDE=" + std::to_string(wide);
            cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), opts);
            cl_kernel kernel = clCreateKernel(prog, "kernel_gemv_noshuffle_q2_0_f32", &err);
            CL_CHECK(err);
            CL_CHECK(clReleaseProgram(prog));
            size_t max_workgroup;
            CL_CHECK(clGetKernelWorkGroupInfo(kernel, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_workgroup), &max_workgroup, nullptr));
-            if (64*(size_t)backend_ctx->q2_0_mv_nsg <= max_workgroup) {
+            size_t workgroup = 512;
+            while (workgroup > max_workgroup) {
+                workgroup /= 2;
+            }
+            GGML_ASSERT(workgroup >= 64);
+            if (wide) {
+                backend_ctx->kernel_gemv_noshuffle_q2_0_f32_wide = kernel;
+                backend_ctx->q2_0_mv_workgroup_size_wide = workgroup;
+            } else {
                backend_ctx->kernel_gemv_noshuffle_q2_0_f32 = kernel;
-                break;
+                backend_ctx->q2_0_mv_workgroup_size = workgroup;
            }
-            CL_CHECK(clReleaseKernel(kernel));
-            GGML_ASSERT(backend_ctx->q2_0_mv_nsg > 1);
-            backend_ctx->q2_0_mv_nsg /= 2;
        }
        GGML_LOG_CONT(".");
    }
@@ -24248,7 +24268,15 @@ static void ggml_cl_mul_mat_q2_0_f32_adreno(ggml_backend_t backend, const ggml_t
    const int k = src0->ne[0];

    if (n == 1 || !ggml_is_contiguous(src1) || offset1 % backend_ctx->alignment != 0) {
-        cl_kernel gemv = backend_ctx->kernel_gemv_noshuffle_q2_0_f32;
+        const bool wide = m >= 16384 && m % 8 == 0;
+        cl_kernel gemv = wide ? backend_ctx->kernel_gemv_noshuffle_q2_0_f32_wide : backend_ctx->kernel_gemv_noshuffle_q2_0_f32;
+        const size_t workgroup = wide ? backend_ctx->q2_0_mv_workgroup_size_wide : backend_ctx->q2_0_mv_workgroup_size;
+        const size_t rows = wide ? 8 : 4;
+        size_t nsg = std::min((size_t)backend_ctx->q2_0_mv_nsg, wide ? size_t(8) : m >= 4096 ? size_t(16) : size_t(32));
+        while (nsg > (size_t)k/64 || nsg*16 > workgroup) {
+            nsg /= 2;
+        }
+        const size_t lanes = workgroup/nsg;
        const int stride_d = dst->ne[0];
        const int stride_b = src1->nb[1]/sizeof(float);
        CL_CHECK(clSetKernelArg(gemv,  0, sizeof(cl_mem),    &extra0->qs));
@@ -24261,9 +24289,8 @@ static void ggml_cl_mul_mat_q2_0_f32_adreno(ggml_backend_t backend, const ggml_t
        CL_CHECK(clSetKernelArg(gemv,  7, sizeof(int),       &m));
        CL_CHECK(clSetKernelArg(gemv,  8, sizeof(int),       &stride_b));
        CL_CHECK(clSetKernelArg(gemv,  9, sizeof(int),       &stride_d));
-        const size_t nsg = backend_ctx->q2_0_mv_nsg;
-        size_t global[] = { (size_t)CEIL_DIV(m, 128)*64, (size_t)n*nsg, 1 };
-        size_t local[] = { 64, nsg, 1 };
+        size_t global[] = { (size_t)CEIL_DIV(m, rows*lanes)*lanes, (size_t)n*nsg, 1 };
+        size_t local[] = { lanes, nsg, 1 };
        backend_ctx->enqueue_ndrange_kernel(gemv, 3, global, local, dst);
        return;
    }
@@ -30881,9 +30908,15 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
    // Override with GGML_OPENCL_A7X_F32_LM_BYPASS=0.
    static const char * a7x_f32lm_env    = getenv("GGML_OPENCL_A7X_F32_LM_BYPASS");
    static const bool   a7x_f32lm_bypass = (a7x_f32lm_env == nullptr || a7x_f32lm_env[0] != '0');
+    // The 64-column tile wastes work on small F32 matrices, including block rotations.
+    const bool f32_use_mv = (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E ||
+                            (backend_ctx->kernel_mul_mv_f32_f32_r2 != nullptr && ne00 <= 512)) &&
+                           src0t == GGML_TYPE_F32 && ne00 >= 256 && ne00 <= 1024 && ne00 % 16 == 0 && ne01 <= 512 && ne11 <= 16 &&
+                           ne02 == 1 && ne12 == 1 && ne13 == 1 &&
+                           ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && (offset0 | offset1) % 16 == 0;
    if (src1t == GGML_TYPE_F32 &&
        ne00 % 16 == 0 &&
-        ne11 > 1 &&
+        ne11 > 1 && !f32_use_mv &&
        !(a7x_f32lm_bypass && src0t == GGML_TYPE_F32 && ne11 > 8 &&
          backend_ctx->adreno_gen == ADRENO_GPU_GEN::A7X)) {
        switch(src0t) {
@@ -32072,6 +32105,9 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
            GGML_ASSERT(src1t == GGML_TYPE_F32);
            kernel = backend_ctx->kernel_mul_mat_f32_f32;
            nrows = 4;
+            if (f32_use_mv && ne11 > 1 && backend_ctx->kernel_mul_mv_f32_f32_r2 != nullptr) {
+                kernel = backend_ctx->kernel_mul_mv_f32_f32_r2;
+            }

            if (backend_ctx->gpu_family == INTEL) {
                nth0 = 32;
@@ -33833,8 +33869,9 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
            backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
        } else {
            int64_t ny = (ne11 + nrows - 1)/nrows;
+            int64_t nx = kernel == backend_ctx->kernel_mul_mv_f32_f32_r2 ? CEIL_DIV(ne01, 2) : ne01;

-            size_t global_work_size[] = {(size_t)ne01*nth0, (size_t)ny*nth1, (size_t)ne12*ne13};
+            size_t global_work_size[] = {(size_t)nx*nth0, (size_t)ny*nth1, (size_t)ne12*ne13};
            size_t local_work_size[] = {(size_t)nth0, (size_t)nth1, 1};

            backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
@@ -1,7 +1,19 @@
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable

-#ifndef Q2_0_MV_NSG
-#define Q2_0_MV_NSG 8
+#if Q2_0_MV_WIDE
+#define Q2_0_MV_ROWS 8
+#define Q2_0_MV_LOAD_D vload8
+#define Q2_0_MV_STORE vstore8
+#define Q2_0_MV_FLOAT convert_float8
+typedef float8 q2_0_mv_acc_t;
+typedef uchar8 q2_0_mv_quant_t;
+#else
+#define Q2_0_MV_ROWS 4
+#define Q2_0_MV_LOAD_D vload4
+#define Q2_0_MV_STORE vstore4
+#define Q2_0_MV_FLOAT convert_float4
+typedef float4 q2_0_mv_acc_t;
+typedef uchar4 q2_0_mv_quant_t;
 #endif

 kernel void kernel_gemv_noshuffle_q2_0_f32(
@@ -19,46 +31,47 @@ kernel void kernel_gemv_noshuffle_q2_0_f32(
    src1 = (global const float *)((global const char *)src1 + offset1);
    dst = (global float *)((global char *)dst + offsetd);

-    const uint m = (uint)ne01;
-    const uint mh = m >> 1;
+    const uint mr = (uint)ne01/Q2_0_MV_ROWS;
    const uint lid = get_local_id(0);
    const uint sgi = get_local_id(1);
+    const uint lanes = get_local_size(0);
+    const uint nsg = get_local_size(1);
    const uint col = get_group_id(1);
-    const uint j = get_group_id(0)*64u + lid;
-    const uint row = j << 1;
+    const uint j = get_group_id(0)*lanes + lid;
    global const float * y = src1 + (ulong)col * (uint)ne10;

-    float2 sum = 0.0f;
-    if (j < mh) {
-        global const ushort * qs = (global const ushort *)src0_qs;
-        for (uint ib = sgi; ib < (uint)ne00/64; ib += Q2_0_MV_NSG) {
-            const float2 d = convert_float2(vload2(j + ib*mh, src0_d));
-            float2 acc = 0.0f;
+    q2_0_mv_acc_t sum = 0.0f;
+    if (j < mr) {
+        for (uint ib = sgi; ib < (uint)ne00/64; ib += nsg) {
+            const q2_0_mv_acc_t d = Q2_0_MV_FLOAT(Q2_0_MV_LOAD_D(j + ib*mr, src0_d));
+            q2_0_mv_acc_t acc = 0.0f;
            for (uint u = 0; u < 16; ++u) {
                const uint group = ib*16 + u;
-                const uint q = qs[j + group*mh];
+#if Q2_0_MV_WIDE
+                const uchar8 q = as_uchar8(((global const uint2 *)src0_qs)[j + group*mr]);
+#else
+                const uchar4 q = as_uchar4(((global const uint *)src0_qs)[j + group*mr]);
+#endif
                const float4 yv = vload4(group, y);
-                const int4 v0 = convert_int4(((uint4)(q) >> (uint4)(0, 2, 4, 6)) & (uint4)(3)) - (int4)(1);
-                const int4 v1 = convert_int4(((uint4)(q) >> (uint4)(8, 10, 12, 14)) & (uint4)(3)) - (int4)(1);
-                acc.s0 += dot(yv, convert_float4(v0));
-                acc.s1 += dot(yv, convert_float4(v1));
+                acc += (Q2_0_MV_FLOAT( q                          & (q2_0_mv_quant_t)(3)) - 1.0f)*yv.s0;
+                acc += (Q2_0_MV_FLOAT((q >> (q2_0_mv_quant_t)(2)) & (q2_0_mv_quant_t)(3)) - 1.0f)*yv.s1;
+                acc += (Q2_0_MV_FLOAT((q >> (q2_0_mv_quant_t)(4)) & (q2_0_mv_quant_t)(3)) - 1.0f)*yv.s2;
+                acc += (Q2_0_MV_FLOAT( q >> (q2_0_mv_quant_t)(6))                          - 1.0f)*yv.s3;
            }
            sum += d*acc;
        }
    }

-#if Q2_0_MV_NSG > 1
-    local float2 part[Q2_0_MV_NSG][64];
-    part[sgi][lid] = sum;
+    local q2_0_mv_acc_t part[512];
+    part[sgi*lanes + lid] = sum;
    barrier(CLK_LOCAL_MEM_FENCE);
    if (sgi != 0) {
        return;
    }
-    for (uint s = 1; s < Q2_0_MV_NSG; ++s) {
-        sum += part[s][lid];
+    for (uint s = 1; s < nsg; ++s) {
+        sum += part[s*lanes + lid];
    }
-#endif
-    if (j < mh) {
-        vstore2(sum, 0, dst + (ulong)col*(uint)ne0 + row);
+    if (j < mr) {
+        Q2_0_MV_STORE(sum, 0, dst + (ulong)col*(uint)ne0 + j*Q2_0_MV_ROWS);
    }
 }
@@ -116,3 +116,70 @@ kernel void kernel_mul_mat_f32_f32(
        }
    }
 }
+
+// Each group computes two rows and four columns, with 32 lanes per row.
+kernel void kernel_mul_mv_f32_f32_r2(
+        global const char * src0,
+        ulong offset0,
+        global const char * src1,
+        ulong offset1,
+        global float * dst,
+        ulong offsetd,
+        int ne00,
+        int ne01,
+        int ne02,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        int ne10,
+        int ne11,
+        int ne12,
+        ulong nb10,
+        ulong nb11,
+        ulong nb12,
+        ulong nb13,
+        int ne0,
+        int ne1,
+        int r2,
+        int r3
+) {
+    src0 += offset0;
+    src1 += offset1;
+    dst = (global float *)((global char *)dst + offsetd);
+
+    const int lid = get_local_id(0);
+    const int lane = lid % 32;
+    const int row = get_group_id(0)*2 + lid/32;
+    const int col = get_group_id(1)*4;
+    const int im = get_group_id(2);
+    const int i12 = im % ne12;
+    const int i13 = im / ne12;
+    global const float4 * x = (global const float4 *)(src0 + min(row, ne01 - 1)*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03);
+    global const char * y = src1 + col*nb11 + i12*nb12 + i13*nb13;
+
+    float4 sum = 0.0f;
+    for (int i = lane; i < ne00/4; i += 32) {
+        const float4 xv = x[i];
+        if (col + 0 < ne11) sum.s0 += dot(xv, ((global const float4 *)(y + 0*nb11))[i]);
+        if (col + 1 < ne11) sum.s1 += dot(xv, ((global const float4 *)(y + 1*nb11))[i]);
+        if (col + 2 < ne11) sum.s2 += dot(xv, ((global const float4 *)(y + 2*nb11))[i]);
+        if (col + 3 < ne11) sum.s3 += dot(xv, ((global const float4 *)(y + 3*nb11))[i]);
+    }
+
+    local float4 part[64];
+    part[lid] = sum;
+    barrier(CLK_LOCAL_MEM_FENCE);
+    for (int s = 16; s > 0; s /= 2) {
+        if (lane < s) {
+            part[lid] += part[lid + s];
+        }
+        barrier(CLK_LOCAL_MEM_FENCE);
+    }
+    if (lane == 0 && row < ne01) {
+        if (col + 0 < ne11) dst[im*ne1*ne0 + (col + 0)*ne0 + row] = part[lid].s0;
+        if (col + 1 < ne11) dst[im*ne1*ne0 + (col + 1)*ne0 + row] = part[lid].s1;
+        if (col + 2 < ne11) dst[im*ne1*ne0 + (col + 2)*ne0 + row] = part[lid].s2;
+        if (col + 3 < ne11) dst[im*ne1*ne0 + (col + 3)*ne0 + row] = part[lid].s3;
+    }
+}
@@ -182,7 +182,6 @@ struct webgpu_capabilities {
    uint32_t subgroup_size     = 0;
    uint32_t min_subgroup_size = 0;
    uint32_t max_subgroup_size = 0;
-    size_t   memset_bytes_per_thread;
 };

 // Stores global webgpu members
@@ -256,8 +255,6 @@ struct webgpu_context_struct {
    wgpu::ComputePassEncoder active_compute_pass;
    bool                     batch_compute_passes = true;

-    size_t memset_bytes_per_thread;
-
 #ifdef GGML_WEBGPU_GPU_PROFILE
    // Profiling: per-shader GPU time in ms
    std::unordered_map<std::string, double> shader_gpu_time_ms;
@@ -672,38 +669,54 @@ static void ggml_backend_webgpu_buffer_memset(webgpu_global_context & ctx,
                                              size_t                  size) {
    std::lock_guard<std::recursive_mutex> lock(ctx->mutex);

-    std::vector<uint32_t>             params  = { (uint32_t) offset, (uint32_t) size, value };
-    std::vector<wgpu::BindGroupEntry> entries = { ggml_webgpu_make_bind_group_entry(0, buf, 0, buf.GetSize()) };
-    size_t                            bytes_per_wg =
-        ctx->capabilities.limits.maxComputeInvocationsPerWorkgroup * ctx->capabilities.memset_bytes_per_thread;
-    uint32_t wg_x = CEIL_DIV(size + 3, bytes_per_wg);
+    const size_t max_bind_size = ctx->capabilities.limits.maxStorageBufferBindingSize;
+    const size_t max_wg_size   = ctx->capabilities.limits.maxComputeInvocationsPerWorkgroup;

-    ctx->queue.WriteBuffer(ctx->memset_params_buf, 0, params.data(), params.size() * sizeof(uint32_t));
+    while (size > 0) {
+        const size_t align_offset = offset & ~(ctx->capabilities.limits.minStorageBufferOffsetAlignment - 1);
+        const size_t lead         = offset - align_offset;
+        const size_t chunk_size   = std::min(size, max_bind_size - lead);
+        const size_t bind_size    = ROUNDUP_POW2(chunk_size + lead, WEBGPU_STORAGE_BUF_BINDING_MULT);

-    wgpu::BindGroupEntry params_entry = {};
-    params_entry.binding              = 1;
-    params_entry.buffer               = ctx->memset_params_buf;
-    params_entry.offset               = 0;
-    params_entry.size                 = WEBGPU_PARAMS_BUF_SIZE_BYTES;
-    entries.push_back(params_entry);
+        std::vector<uint32_t>             params  = { (uint32_t) lead, (uint32_t) chunk_size, value };
+        std::vector<wgpu::BindGroupEntry> entries = { ggml_webgpu_make_bind_group_entry(0, buf, align_offset,
+                                                                                        bind_size) };
+        uint32_t                          wg_x;
+        uint32_t                          wg_y;
+        uint32_t                          n_values = CEIL_DIV(chunk_size + lead, sizeof(uint32_t));
+        uint32_t                          total_wg = CEIL_DIV(n_values, max_wg_size);
+        compute_2d_workgroups(total_wg, ctx->capabilities.limits.maxComputeWorkgroupsPerDimension, wg_x, wg_y);

-    wgpu::BindGroupDescriptor bind_group_desc;
-    bind_group_desc.layout     = ctx->memset_pipeline.pipeline.GetBindGroupLayout(0);
-    bind_group_desc.entryCount = entries.size();
-    bind_group_desc.entries    = entries.data();
-    bind_group_desc.label      = ctx->memset_pipeline.name.c_str();
-    wgpu::BindGroup bind_group = ctx->device.CreateBindGroup(&bind_group_desc);
+        ctx->queue.WriteBuffer(ctx->memset_params_buf, 0, params.data(), params.size() * sizeof(uint32_t));

-    wgpu::CommandEncoder     encoder = ctx->device.CreateCommandEncoder();
-    wgpu::ComputePassEncoder pass    = encoder.BeginComputePass();
-    pass.SetPipeline(ctx->memset_pipeline.pipeline);
-    pass.SetBindGroup(0, bind_group);
-    pass.DispatchWorkgroups(wg_x, 1, 1);
-    pass.End();
+        wgpu::BindGroupEntry params_entry = {};
+        params_entry.binding              = 1;
+        params_entry.buffer               = ctx->memset_params_buf;
+        params_entry.offset               = 0;
+        params_entry.size                 = WEBGPU_PARAMS_BUF_SIZE_BYTES;
+        entries.push_back(params_entry);

-    wgpu::CommandBuffer              command  = encoder.Finish();
-    std::vector<wgpu::CommandBuffer> commands = { command };
-    ctx->queue.Submit(commands.size(), commands.data());
+        wgpu::BindGroupDescriptor bind_group_desc;
+        bind_group_desc.layout     = ctx->memset_pipeline.pipeline.GetBindGroupLayout(0);
+        bind_group_desc.entryCount = entries.size();
+        bind_group_desc.entries    = entries.data();
+        bind_group_desc.label      = ctx->memset_pipeline.name.c_str();
+        wgpu::BindGroup bind_group = ctx->device.CreateBindGroup(&bind_group_desc);
+
+        wgpu::CommandEncoder     encoder = ctx->device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass    = encoder.BeginComputePass();
+        pass.SetPipeline(ctx->memset_pipeline.pipeline);
+        pass.SetBindGroup(0, bind_group);
+        pass.DispatchWorkgroups(wg_x, wg_y, 1);
+        pass.End();
+
+        wgpu::CommandBuffer              command  = encoder.Finish();
+        std::vector<wgpu::CommandBuffer> commands = { command };
+        ctx->queue.Submit(commands.size(), commands.data());
+
+        offset += chunk_size;
+        size -= chunk_size;
+    }
 }

 /** End WebGPU Actions */
@@ -4006,12 +4019,10 @@ static size_t ggml_backend_webgpu_buffer_type_get_alignment(ggml_backend_buffer_
    return dev_ctx->webgpu_global_ctx->capabilities.limits.minStorageBufferOffsetAlignment;
 }

-// maxBufferSize might be larger, but you can't bind more than
-// maxStorageBufferBindingSize to a single binding.
 static size_t ggml_backend_webgpu_buffer_type_get_max_size(ggml_backend_buffer_type_t buft) {
    ggml_backend_webgpu_device_context * dev_ctx =
        static_cast<ggml_backend_webgpu_device_context *>(buft->device->context);
-    return dev_ctx->webgpu_global_ctx->capabilities.limits.maxStorageBufferBindingSize;
+    return dev_ctx->webgpu_global_ctx->capabilities.limits.maxBufferSize;
 }

 static size_t ggml_backend_webgpu_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft,
@@ -4176,17 +4187,9 @@ static ggml_guid_t ggml_backend_webgpu_guid(void) {
 }

 static void ggml_webgpu_init_memset_pipeline(webgpu_global_context & ctx) {
-    // we use the maximum workgroup size for the memset pipeline
-    size_t max_threads = ctx->capabilities.limits.maxComputeInvocationsPerWorkgroup *
-                         ctx->capabilities.limits.maxComputeWorkgroupsPerDimension;
-    // Size the bytes_per_thread so that the largest buffer size can be handled
-    ctx->capabilities.memset_bytes_per_thread =
-        CEIL_DIV(ctx->capabilities.limits.maxStorageBufferBindingSize, max_threads);
-    std::vector<wgpu::ConstantEntry> constants(2);
+    std::vector<wgpu::ConstantEntry> constants(1);
    constants[0].key     = "wg_size";
    constants[0].value   = ctx->capabilities.limits.maxComputeInvocationsPerWorkgroup;
-    constants[1].key     = "bytes_per_thread";
-    constants[1].value   = ctx->capabilities.memset_bytes_per_thread;
    ctx->memset_pipeline = ggml_webgpu_create_pipeline(ctx->device, wgsl_memset, "memset", constants);
 }

@@ -11,30 +11,33 @@ struct Params {
 var<uniform> params: Params;

 override wg_size: u32;
-override bytes_per_thread: u32;

@compute @workgroup_size(wg_size)
-fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-    let i = gid.x * bytes_per_thread;
+fn main(@builtin(workgroup_id) wid: vec3<u32>,
+        @builtin(num_workgroups) num_wg: vec3<u32>,
+        @builtin(local_invocation_id) lid: vec3<u32>) {
+
+    let i = (wid.x + wid.y * num_wg.x) * wg_size + lid.x;
    let start = params.offset;
    let end = params.offset + params.size;
+    let id = (start >> 2u) + i;
+    let id_u32 = id * 4;

-    for (var j: u32 = 0u; j < bytes_per_thread; j += 4) {
-        let byte_index = start + i + j;
-        if (byte_index + 4 <= end) {
-            output_buffer[byte_index >> 2] = params.value;
-        } else {
-            // Handle tail (unaligned)
-            for (var k: u32 = 0; k < 4; k++) {
-                let idx = byte_index + k;
-                if (idx < end) {
-                    let word_idx = idx >> 2;
-                    let bit_offset = (idx & 3) * 8u;
-                    let mask = ~(0xffu << bit_offset);
-                    let existing = output_buffer[word_idx];
-                    output_buffer[word_idx] = (existing & mask) | (params.value & (0xffu << bit_offset));
-                }
-            }
-        }
+    if (id_u32 >= end) {
+        return;
+    }
+
+    let lo = max(start, id_u32) - id_u32;
+    let hi = min(end, id_u32 + 4u) - id_u32;
+    var mask = 0xFFFFFFFFu << (8u * lo);
+    if (hi < 4u) {
+        mask &= (1u << (8u * hi)) - 1u;
+    }
+
+    if (mask == 0xFFFFFFFFu) {
+        output_buffer[id] = params.value;
+    } else {
+        let existing = output_buffer[id];
+        output_buffer[id] = (existing & ~(mask)) | (params.value & mask);
    }
 }
@@ -485,27 +485,25 @@ void ggml_fp16_to_fp32_row(const ggml_fp16_t * x, float * y, int64_t n) {
 }

 void ggml_fp32_to_fp16_row(const float * x, ggml_fp16_t * y, int64_t n) {
-    int i = 0;
-    for (; i < n; ++i) {
+    for (int64_t i = 0; i < n; i++) {
        y[i] = GGML_FP32_TO_FP16(x[i]);
    }
 }

 void ggml_bf16_to_fp32_row(const ggml_bf16_t * x, float * y, int64_t n) {
-    int i = 0;
-    for (; i < n; ++i) {
+    for (int64_t i = 0; i < n; i++) {
        y[i] = GGML_BF16_TO_FP32(x[i]);
    }
 }

 void ggml_fp32_to_bf16_row_ref(const float * x, ggml_bf16_t * y, int64_t n) {
-    for (int i = 0; i < n; i++) {
+    for (int64_t i = 0; i < n; i++) {
        y[i] = ggml_compute_fp32_to_bf16(x[i]);
    }
 }

 void ggml_fp32_to_bf16_row(const float * x, ggml_bf16_t * y, int64_t n) {
-  int i = 0;
+    int64_t i = 0;
 #if defined(__AVX512BF16__)
  // subnormals are flushed to zero on this platform
  for (; i + 32 <= n; i += 32) {
@@ -1788,6 +1786,7 @@ static struct ggml_tensor * ggml_new_tensor_impl(
        enum   ggml_type      type,
        int                   n_dims,
        const int64_t       * ne,
+        const size_t        * nb,
        struct ggml_tensor  * view_src,
        size_t                view_offs) {

@@ -1814,8 +1813,6 @@ static struct ggml_tensor * ggml_new_tensor_impl(
        data_size *= ne[i];
    }

-    GGML_ASSERT(view_src == NULL || data_size == 0 || data_size + view_offs <= ggml_nbytes(view_src));
-
    void * data = view_src != NULL ? view_src->data : NULL;
    if (data != NULL) {
        data = (char *) data + view_offs;
@@ -1859,12 +1856,23 @@ static struct ggml_tensor * ggml_new_tensor_impl(
        result->ne[i] = ne[i];
    }

-    result->nb[0] = ggml_type_size(type);
-    result->nb[1] = result->nb[0]*(result->ne[0]/ggml_blck_size(type));
-    for (int i = 2; i < GGML_MAX_DIMS; i++) {
-        result->nb[i] = result->nb[i - 1]*result->ne[i - 1];
+    // a view built with strides takes them from the caller, any other tensor is contiguous
+    if (nb != NULL) {
+        for (int i = 0; i < GGML_MAX_DIMS; i++) {
+            result->nb[i] = nb[i];
+        }
+    } else {
+        result->nb[0] = ggml_type_size(type);
+        result->nb[1] = result->nb[0]*(result->ne[0]/ggml_blck_size(type));
+        for (int i = 2; i < GGML_MAX_DIMS; i++) {
+            result->nb[i] = result->nb[i - 1]*result->ne[i - 1];
+        }
    }

+    // a view stays inside its source, over the bytes its strides actually reach
+    const size_t view_size = ggml_nbytes(result);
+    GGML_ASSERT(view_src == NULL || view_size == 0 || view_size + view_offs <= ggml_nbytes(view_src));
+
    ctx->n_objects++;

    return result;
@@ -1875,7 +1883,7 @@ struct ggml_tensor * ggml_new_tensor(
        enum   ggml_type      type,
        int                   n_dims,
        const int64_t       * ne) {
-    return ggml_new_tensor_impl(ctx, type, n_dims, ne, NULL, 0);
+    return ggml_new_tensor_impl(ctx, type, n_dims, ne, NULL, NULL, 0);
 }

 struct ggml_tensor * ggml_new_tensor_1d(
@@ -1992,13 +2000,9 @@ struct ggml_tensor * ggml_format_name(struct ggml_tensor * tensor, const char *
 struct ggml_tensor * ggml_view_tensor(
        struct ggml_context * ctx,
        struct ggml_tensor  * src) {
-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, src->type, GGML_MAX_DIMS, src->ne, src, 0);
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, src->type, GGML_MAX_DIMS, src->ne, src->nb, src, 0);
    ggml_format_name(result, "%s (view)", src->name);

-    for (int i = 0; i < GGML_MAX_DIMS; i++) {
-        result->nb[i] = src->nb[i];
-    }
-
    return result;
 }

@@ -2948,7 +2952,7 @@ static struct ggml_tensor * ggml_glu_impl(
    }

    int64_t ne[GGML_MAX_DIMS] = { a->ne[0] / 2 }; for (int i = 1; i < GGML_MAX_DIMS; i++) ne[i] = a->ne[i];
-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, GGML_MAX_DIMS, b ? a->ne : ne, NULL, 0);
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, GGML_MAX_DIMS, b ? a->ne : ne, NULL, NULL, 0);

    ggml_set_op_params_i32(result, 0, (int32_t) op);
    ggml_set_op_params_i32(result, 1, (int32_t) swapped);
@@ -3757,7 +3761,7 @@ struct ggml_tensor * ggml_reshape(
    // as only the shape of b is relevant, and not its memory layout, b is allowed to be non contiguous.
    GGML_ASSERT(ggml_nelements(a) == ggml_nelements(b));

-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, GGML_MAX_DIMS, b->ne, a, 0);
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, GGML_MAX_DIMS, b->ne, NULL, a, 0);
    ggml_format_name(result, "%s (reshaped)", a->name);

    result->op     = GGML_OP_RESHAPE;
@@ -3774,7 +3778,7 @@ struct ggml_tensor * ggml_reshape_1d(
    GGML_ASSERT(ggml_nelements(a) == ne0);

    const int64_t ne[1] = { ne0 };
-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 1, ne, a, 0);
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 1, ne, NULL, a, 0);
    ggml_format_name(result, "%s (reshaped)", a->name);

    result->op     = GGML_OP_RESHAPE;
@@ -3792,7 +3796,7 @@ struct ggml_tensor * ggml_reshape_2d(
    GGML_ASSERT(ggml_nelements(a) == ne0*ne1);

    const int64_t ne[2] = { ne0, ne1 };
-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 2, ne, a, 0);
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 2, ne, NULL, a, 0);
    ggml_format_name(result, "%s (reshaped)", a->name);

    result->op     = GGML_OP_RESHAPE;
@@ -3811,7 +3815,7 @@ struct ggml_tensor * ggml_reshape_3d(
    GGML_ASSERT(ggml_nelements(a) == ne0*ne1*ne2);

    const int64_t ne[3] = { ne0, ne1, ne2 };
-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 3, ne, a, 0);
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 3, ne, NULL, a, 0);
    ggml_format_name(result, "%s (reshaped)", a->name);

    result->op     = GGML_OP_RESHAPE;
@@ -3831,7 +3835,7 @@ struct ggml_tensor * ggml_reshape_4d(
    GGML_ASSERT(ggml_nelements(a) == ne0*ne1*ne2*ne3);

    const int64_t ne[4] = { ne0, ne1, ne2, ne3 };
-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 4, ne, a, 0);
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, 4, ne, NULL, a, 0);
    ggml_format_name(result, "%s (reshaped)", a->name);

    result->op     = GGML_OP_RESHAPE;
@@ -3845,8 +3849,12 @@ static struct ggml_tensor * ggml_view_impl(
        struct ggml_tensor  * a,
        int                   n_dims,
        const int64_t       * ne,
+        const size_t        * nb,
        size_t                offset) {
-    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, n_dims, ne, a, offset);
+    // a strided view counts its strides from elements that are contiguous along dim 0 of its source
+    GGML_ASSERT(nb == NULL || a->nb[0] == ggml_type_size(a->type));
+
+    struct ggml_tensor * result = ggml_new_tensor_impl(ctx, a->type, n_dims, ne, nb, a, offset);
    ggml_format_name(result, "%s (view)", a->name);

    ggml_set_op_params(result, &offset, sizeof(offset));
@@ -3864,7 +3872,7 @@ struct ggml_tensor * ggml_view_1d(
        struct ggml_tensor  * a,
        int64_t               ne0,
        size_t                offset) {
-    struct ggml_tensor * result = ggml_view_impl(ctx, a, 1, &ne0, offset);
+    struct ggml_tensor * result = ggml_view_impl(ctx, a, 1, &ne0, NULL, offset);

    return result;
 }
@@ -3879,12 +3887,9 @@ struct ggml_tensor * ggml_view_2d(
        size_t                nb1,
        size_t                offset) {
    const int64_t ne[2] = { ne0, ne1 };
+    const size_t  nb[4] = { ggml_type_size(a->type), nb1, nb1*ne1, nb1*ne1 };

-    struct ggml_tensor * result = ggml_view_impl(ctx, a, 2, ne, offset);
-
-    result->nb[1] = nb1;
-    result->nb[2] = result->nb[1]*ne1;
-    result->nb[3] = result->nb[2];
+    struct ggml_tensor * result = ggml_view_impl(ctx, a, 2, ne, nb, offset);

    return result;
 }
@@ -3901,12 +3906,9 @@ struct ggml_tensor * ggml_view_3d(
        size_t                nb2,
        size_t                offset) {
    const int64_t ne[3] = { ne0, ne1, ne2 };
+    const size_t  nb[4] = { ggml_type_size(a->type), nb1, nb2, nb2*ne2 };

-    struct ggml_tensor * result = ggml_view_impl(ctx, a, 3, ne, offset);
-
-    result->nb[1] = nb1;
-    result->nb[2] = nb2;
-    result->nb[3] = result->nb[2]*ne2;
+    struct ggml_tensor * result = ggml_view_impl(ctx, a, 3, ne, nb, offset);

    return result;
 }
@@ -3925,12 +3927,9 @@ struct ggml_tensor * ggml_view_4d(
        size_t                nb3,
        size_t                offset) {
    const int64_t ne[4] = { ne0, ne1, ne2, ne3 };
+    const size_t  nb[4] = { ggml_type_size(a->type), nb1, nb2, nb3 };

-    struct ggml_tensor * result = ggml_view_impl(ctx, a, 4, ne, offset);
-
-    result->nb[1] = nb1;
-    result->nb[2] = nb2;
-    result->nb[3] = nb3;
+    struct ggml_tensor * result = ggml_view_impl(ctx, a, 4, ne, nb, offset);

    return result;
 }
@@ -1463,6 +1463,7 @@ struct gguf_writer_base {
    // we bet on devirtualization
    virtual void write(int8_t val) = 0;
    virtual void write(const std::vector<int8_t> & val) = 0;
+    virtual void write(const void * data, size_t size) = 0;
    virtual void write_tensor_data(const struct gguf_tensor_info & info, size_t offset_data, size_t alignment) = 0;

    template <typename T>
@@ -1579,21 +1580,31 @@ struct gguf_writer_buf final : public gguf_writer_base {
        written_bytes += val.size();
    }

+    void write(const void * data, const size_t size) override {
+        buf.insert(buf.end(), (const int8_t *) data, (const int8_t *) data + size);
+        written_bytes += size;
+    }
+
    void write_tensor_data(const struct gguf_tensor_info & info, const size_t offset_data, const size_t alignment) override {
        GGML_ASSERT(buf.size() - offset_data == info.offset);

        GGML_ASSERT(ggml_is_contiguous(&info.t));
-        const size_t offset = buf.size();
        const size_t nbytes = ggml_nbytes(&info.t);

-        buf.resize(offset + nbytes);
-        if (info.t.buffer) {
-            ggml_backend_tensor_get(&info.t, buf.data() + offset, 0, nbytes);
+        if (info.t.buffer && ggml_backend_buffer_is_host(info.t.buffer)) {
+            GGML_ASSERT(info.t.data != NULL && "tensor not allocated");
+            write(info.t.data, nbytes);
        } else {
-            GGML_ASSERT(info.t.data);
-            memcpy(buf.data() + offset, info.t.data, nbytes);
+            const size_t offset = buf.size();
+            buf.resize(offset + nbytes);
+            if (info.t.buffer) {
+                ggml_backend_tensor_get(&info.t, buf.data() + offset, 0, nbytes);
+            } else {
+                GGML_ASSERT(info.t.data);
+                memcpy(buf.data() + offset, info.t.data, nbytes);
+            }
+            written_bytes += nbytes;
        }
-        written_bytes += nbytes;

        pad(alignment);
    }
@@ -1624,20 +1635,33 @@ struct gguf_writer_file final : public gguf_writer_base {
        }
    }

+    void write(const void * data, const size_t size) override {
+        const auto ret = fwrite(data, 1, size, file);
+        written_bytes += size;
+        if (ret != size) {
+            throw std::runtime_error("unexpected fwrite number of bytes written, '" + std::to_string(ret) + "' instead of '" + std::to_string(size) + "'");
+        }
+    }
+
    void write_tensor_data(const struct gguf_tensor_info & info, const size_t offset_data, const size_t alignment) override {
        GGML_ASSERT(written_bytes - offset_data == info.offset);

        GGML_ASSERT(ggml_is_contiguous(&info.t));
        const size_t nbytes = ggml_nbytes(&info.t);

-        std::vector<int8_t> buf(nbytes);
-        if (info.t.buffer) {
-            ggml_backend_tensor_get(&info.t, buf.data(), 0, nbytes);
+        if (info.t.buffer && ggml_backend_buffer_is_host(info.t.buffer)) {
+            GGML_ASSERT(info.t.data != NULL && "tensor not allocated");
+            write(info.t.data, nbytes);
        } else {
-            GGML_ASSERT(info.t.data);
-            memcpy(buf.data(), info.t.data, nbytes);
+            std::vector<int8_t> buf(nbytes);
+            if (info.t.buffer) {
+                ggml_backend_tensor_get(&info.t, buf.data(), 0, nbytes);
+            } else {
+                GGML_ASSERT(info.t.data);
+                memcpy(buf.data(), info.t.data, nbytes);
+            }
+            write(buf);
        }
-        write(buf);

        pad(alignment);
    }
@@ -1295,6 +1295,11 @@ class MODEL_TENSOR(IntEnum):
    A_GEN_CODE_ATTN_GATE   = auto() # lfm3a: per-head sigmoid gate on the attention output
    A_GEN_CODE_HEAD_NORM   = auto() # lfm3a: per-codebook norm before the output head, merged 2D [n_codebooks, dim]
    A_GEN_CODE_MASK_EMBD   = auto() # lfm3a: MaskGIT mask embedding
+    A_GEN_CODE_ENORM       = auto() # drama3: per-layer norm of the code embedding
+    A_GEN_CODE_HNORM       = auto() # drama3: per-layer norm of the input hidden state
+    A_GEN_CODE_EH_PROJ     = auto() # drama3: per-layer [code embd; hidden] -> hidden projection
+    A_GEN_CODE_BLK_EMBD    = auto() # drama3: per-layer code embedding table
+    A_GEN_CODE_BLK_OUT_NORM = auto() # drama3: per-layer output norm
    # qwen3tts code2wav: RVQ codes -> raw PCM
    A_GEN_WAV_QUANT_FIRST_IN       = auto() # semantic RVQ, in_proj (1x1 conv, loaded as 2D)
    A_GEN_WAV_QUANT_FIRST_OUT      = auto() # semantic RVQ, out_proj
@@ -2198,6 +2203,11 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
    MODEL_TENSOR.A_GEN_CODE_ATTN_GATE:      "a.gen.code.blk.{bid}.attn_gate",
    MODEL_TENSOR.A_GEN_CODE_HEAD_NORM:      "a.gen.code.head_norm",
    MODEL_TENSOR.A_GEN_CODE_MASK_EMBD:      "a.gen.code.mask_embd",
+    MODEL_TENSOR.A_GEN_CODE_ENORM:          "a.gen.code.blk.{bid}.enorm",
+    MODEL_TENSOR.A_GEN_CODE_HNORM:          "a.gen.code.blk.{bid}.hnorm",
+    MODEL_TENSOR.A_GEN_CODE_EH_PROJ:        "a.gen.code.blk.{bid}.eh_proj",
+    MODEL_TENSOR.A_GEN_CODE_BLK_EMBD:       "a.gen.code.blk.{bid}.embd",
+    MODEL_TENSOR.A_GEN_CODE_BLK_OUT_NORM:   "a.gen.code.blk.{bid}.out_norm",
    MODEL_TENSOR.A_GEN_WAV_QUANT_FIRST_IN:  "a.gen.wav.quant.first.in_proj",
    MODEL_TENSOR.A_GEN_WAV_QUANT_FIRST_OUT: "a.gen.wav.quant.first.out_proj",
    MODEL_TENSOR.A_GEN_WAV_QUANT_FIRST_CB:  "a.gen.wav.quant.first.codebook",
@@ -2615,6 +2625,11 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
        MODEL_TENSOR.A_GEN_CODE_ATTN_GATE,
        MODEL_TENSOR.A_GEN_CODE_HEAD_NORM,
        MODEL_TENSOR.A_GEN_CODE_MASK_EMBD,
+        MODEL_TENSOR.A_GEN_CODE_ENORM,
+        MODEL_TENSOR.A_GEN_CODE_HNORM,
+        MODEL_TENSOR.A_GEN_CODE_EH_PROJ,
+        MODEL_TENSOR.A_GEN_CODE_BLK_EMBD,
+        MODEL_TENSOR.A_GEN_CODE_BLK_OUT_NORM,
        MODEL_TENSOR.A_GEN_WAV_QUANT_FIRST_IN,
        MODEL_TENSOR.A_GEN_WAV_QUANT_FIRST_OUT,
        MODEL_TENSOR.A_GEN_WAV_QUANT_FIRST_CB,
@@ -6919,6 +6934,7 @@ class VisionProjectorType:
    CHATTERBOX_SPKENC = "chatterbox_spkenc" # audio: CAMPPlus speaker encoder
    CHATTERBOX = "chatterbox" # audio generation: s3gen flow matching + HiFT vocoder
    SOPRANO = "soprano" # audio generation: Vocos decoder
+    DRAMA3_GEN = "drama3_gen" # audio generation: codebook tower + Fish DAC decoder
    HUNYUANVL      = "hunyuanvl"
    PARAKEET       = "parakeet"  # audio
    MINIMAXM3      = "minimax_m3"
@@ -221,10 +221,12 @@ class SpecialVocab:
                if isinstance(merges, list) and merges:
                    if isinstance(merges[0], str):
                        self.merges = merges
-                    elif isinstance(merges[0], list) and len(merges[0]) == 2 and isinstance(merges[0][0], str):
+                    elif isinstance(merges[0], list):
                        # New format since transformers 4.45 to support spaces in merges
                        # ref: https://github.com/ggml-org/llama.cpp/issues/9692
                        # TODO: internally store as the new format instead of converting to old
+                        if not all(isinstance(pair, list) and len(pair) == 2 and all(isinstance(s, str) for s in pair) for pair in merges):
+                            raise ValueError("Unknown tokenizer merges format")
                        if any(' ' in s for pair in merges for s in pair):
                            logger.warning(f'Spaces in merges detected, encoding as {chr(ord(" ") + 256)!r}')
                        self.merges = [
@@ -33,6 +33,7 @@ MANAGED_ENV_NAMES = (
    "GGML_HEXAGON_MM_SELECT",
    "GGML_HEXAGON_FA_SELECT",
    "GGML_HEXAGON_FA_HEAD_SPLIT",
+    "GGML_HEXAGON_LI_SELECT",
    "GGML_HEXAGON_GDN_SELECT",
    "GGML_HEXAGON_AR_SELECT",
    "GGML_HEXAGON_AR_SCATTER",
@@ -235,6 +236,7 @@ def main():
    parser.add_argument("--hex-mm-select", help="Select MUL_MAT and MUL_MAT_ID kernel (GGML_HEXAGON_MM_SELECT) 2:HMX,1:HVX,0:disable")
    parser.add_argument("--hex-fa-select", help="Select Flash Attention kernel (GGML_HEXAGON_FA_SELECT) 2:HMX,1:HVX,0:disable")
    parser.add_argument("--hex-fa-head-split", help="Enable (1) or disable (0) head-parallel flash_attn partitioning (GGML_HEXAGON_FA_HEAD_SPLIT)")
+    parser.add_argument("--hex-li-select", help="Select Lightning Indexer kernel (GGML_HEXAGON_LI_SELECT) 2:HMX,1:HVX,0:disable")
    parser.add_argument("--hex-gdn-select", help="Select Gated Delta Net kernel (GGML_HEXAGON_GDN_SELECT) 2:HMX,1:HVX,0:disable")
    parser.add_argument("--hex-ar-select", help="Select All-Reduce kernel (GGML_HEXAGON_AR_SELECT) 1:enable,0:disable")
    parser.add_argument("--hex-ar-scatter", help="Enable (1) or disable (0) reduce-scatter for fused ALLREDUCE+ADD (GGML_HEXAGON_AR_SCATTER)")
@@ -388,6 +390,7 @@ def main():
    set_env("GGML_HEXAGON_MM_SELECT", args.hex_mm_select)
    set_env("GGML_HEXAGON_FA_SELECT", args.hex_fa_select)
    set_env("GGML_HEXAGON_FA_HEAD_SPLIT", args.hex_fa_head_split)
+    set_env("GGML_HEXAGON_LI_SELECT", args.hex_li_select)
    set_env("GGML_HEXAGON_GDN_SELECT", args.hex_gdn_select)
    set_env("GGML_HEXAGON_AR_SELECT", args.hex_ar_select)
    set_env("GGML_HEXAGON_AR_SCATTER", args.hex_ar_scatter)
@@ -1 +1 @@
-facdc44a67da17df4433adcc468b424f48e9ded2
+57e2c2ca2a7cae756b1c4dfafa0cdfef9051d234
@@ -14,6 +14,7 @@
 #include "llama-sampler.h"
 #include "llama.h"

+#include <algorithm>
 #include <cinttypes>
 #include <cmath>
 #include <cstring>
@@ -332,6 +333,18 @@ llama_context::llama_context(
    cparams.flash_attn = params.flash_attn_type != LLAMA_FLASH_ATTN_TYPE_DISABLED;
    cparams.auto_fa    = params.flash_attn_type == LLAMA_FLASH_ATTN_TYPE_AUTO;

+    // the KV cache is shared with ctx_other, so follow its flash attention setting instead of resolving it here:
+    // the probe disables it when the layers and the shared KV are on different devices, which breaks a quantized V cache
+    if (model.arch == LLM_ARCH_GEMMA4_ASSISTANT) {
+        const bool flash_attn_other = cparams.ctx_other->get_cparams().flash_attn;
+        if (!cparams.auto_fa && cparams.flash_attn != flash_attn_other) {
+            LLAMA_LOG_WARN("%s: flash_attn = %d does not match the context the KV cache is shared with, using %d\n",
+                    __func__, cparams.flash_attn, flash_attn_other);
+        }
+        cparams.flash_attn = flash_attn_other;
+        cparams.auto_fa    = false;
+    }
+
    cparams.fused_gdn_ar = true;
    cparams.fused_gdn_ch = true;
    cparams.auto_fgdn    = false;
@@ -459,6 +472,22 @@ llama_context::llama_context(
            backends.emplace_back(backend);
        }

+        // tensors of ctx_other (e.g. a shared KV cache) can be on devices that this model does not use
+        if (cparams.ctx_other) {
+            for (const auto & dev_other : llama_get_model(cparams.ctx_other)->devices) {
+                const bool found = std::any_of(model.devices.begin(), model.devices.end(),
+                        [&](const llama_device & dev) { return dev.dev == dev_other.dev; });
+                if (found) {
+                    continue;
+                }
+                ggml_backend_t backend = ggml_backend_dev_init(dev_other.dev, nullptr);
+                if (backend == nullptr) {
+                    throw std::runtime_error(format("failed to initialize %s backend", ggml_backend_dev_name(dev_other.dev)));
+                }
+                backends.emplace_back(backend);
+            }
+        }
+
        // add ACCEL backends (such as BLAS)
        for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
            ggml_backend_dev_t dev = ggml_backend_dev_get(i);
@@ -509,12 +509,11 @@ ggml_tensor * llm_build_delta_net_base::build_conv_state(

        // window w starts at column n - (Kw - 1) + w (ascending time) and is rollback slot Kw - 1 - w
        // slots reaching back past this ubatch are left unwritten, like the ssm snapshots in build_recurrent_attn
-        // the windows overlap, so the view is created with one window and widened after the size check
+        // consecutive windows start one element apart, so they overlap
        ggml_tensor * windows = ggml_view_4d(ctx0, conv_input,
-                conv_kernel_size - 1, 1, conv_channels, n_seqs,
+                conv_kernel_size - 1, Kw, conv_channels, n_seqs,
                ggml_element_size(conv_input), conv_input->nb[1], conv_input->nb[2],
                ggml_row_size(conv_input->type, n - (Kw - 1)));
-        windows->ne[1] = Kw;

        // [d-1, slot, channel, seq] view of the cells of this ubatch
        ggml_tensor * slots = ggml_view_4d(ctx0, conv_states_all,
@@ -12795,6 +12795,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
        test_cases.emplace_back(new test_dsv4_hc_post(2560, n_tokens, true));       // synthetic, Qwen4-Exp kernel without the gate chain
    }

+    // Qwen3.8-Flash-Next pooled indexer prefill.
+    test_cases.emplace_back(new test_lightning_indexer(128, 4, 1088, 1024, 1, 1, GGML_TYPE_F32));
+    test_cases.emplace_back(new test_lightning_indexer(128, 4, 1088, 1024, 1, 1, GGML_TYPE_Q4_0));
    // SWIGLU at a 27B-class FFN width, fused [gate|up] vs split operands
    // note: same bytes either way, so a backend that indexes them differently shows it here
    for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16}) {
@@ -153,6 +153,14 @@ static const std::vector<std::string> vendors = {
    "BlackSheep-24B.Q8_0.gguf",
 };

+// a tag must match the full token: F16 ends inside BF16, Q4_XS inside IQ4_XS
+static const std::vector<std::string> twins = {
+    "model-BF16.gguf",
+    "model-F16.gguf",
+    "model-IQ4_XS.gguf",
+    "model-Q4_XS.gguf",
+};
+
 // every speculative sidecar type at the same quant
 static const std::vector<std::string> quad = {
    "model-Q8_0.gguf",
@@ -270,6 +278,23 @@ static const plan_case plan_cases[] = {
     "TheDrummer_Model-24B-v4.1-Q8_0.gguf", {"TheDrummer_Model-24B-v4.1-Q8_0.gguf"},
     "", "", "", "", ""},

+    // the tag picks the file carrying it as a full token, never a longer one
+    {"twins f16 tag", twins, "test/repo:F16", "", false, false,
+     "model-F16.gguf", {"model-F16.gguf"},
+     "", "", "", "", ""},
+
+    {"twins bf16 tag", twins, "test/repo:BF16", "", false, false,
+     "model-BF16.gguf", {"model-BF16.gguf"},
+     "", "", "", "", ""},
+
+    {"twins q4_xs tag", twins, "test/repo:Q4_XS", "", false, false,
+     "model-Q4_XS.gguf", {"model-Q4_XS.gguf"},
+     "", "", "", "", ""},
+
+    {"twins iq4_xs tag", twins, "test/repo:IQ4_XS", "", false, false,
+     "model-IQ4_XS.gguf", {"model-IQ4_XS.gguf"},
+     "", "", "", "", ""},
+
    // every sidecar type resolves at the tag
    {"quad exact tag", quad, "test/repo:Q8_0", "", true, false,
     "model-Q8_0.gguf", {"model-Q8_0.gguf"},
@@ -868,7 +893,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "MTP/Qwen3.8-27B-Uncensored-OrcaRouter-MTP-Q8_0.gguf", 1, "", "", "", "", ""},
      {"BF16", "BF16/Qwen3.8-27B-Uncensored-OrcaRouter-FP8-derived-BF16-00001-of-00002.gguf", 2, "", "", "", "", ""},
      {"Q5_K_M", "Qwen3.8-27B-Uncensored-OrcaRouter-Q5_K_M.gguf", 1, "", "", "", "", ""},
-      {"F16", "BF16/Qwen3.8-27B-Uncensored-OrcaRouter-FP8-derived-BF16-00001-of-00002.gguf", 2, "", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_0", "Q5_K_M", "Q4_K_M", "BF16"}},

    {"ChrisColeTech/qwen-image-edit-uncensored-GGUF",
@@ -980,7 +1005,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Qwen3.5-9B-The-Defiant-Fable-Uncnr-Heretic-NEO-MAX-MTP-Q8_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"IQ2_M", "Qwen3.5-9B-The-Defiant-Fable-Uncnr-Heretic-NEO-MAX-IQ2_M.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "Qwen3.5-9B-The-Defiant-Fable-Uncnr-Heretic-plusIQ-NEO-MAX-MTP-bf16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "Qwen3.5-9B-The-Defiant-Fable-Uncnr-Heretic-plusIQ-NEO-MAX-MTP-bf16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_0", "BF16", "Q4_K_S", "Q4_K_M", "IQ2_M"}},

    {"DavidAU/Qwen3.6-27B-Heretic-Uncensored-FINETUNE-NEO-CODE-Di-IMatrix-MAX-GGUF",
@@ -1044,7 +1069,7 @@ static const real_repo real_repos[] = {
      {"Q4_K_M", "Qwythos-27B-MTP-Q4_K_M.gguf", 1, "mmproj-Qwythos-27B-F16.gguf", "", "", "", ""},
      {"Q8_0", "Qwythos-27B-MTP-Q8_0.gguf", 1, "mmproj-Qwythos-27B-F16.gguf", "", "", "", ""},
      {"BF16", "Qwythos-27B-BF16.gguf", 1, "mmproj-Qwythos-27B-F16.gguf", "", "", "", ""},
-      {"F16", "Qwythos-27B-BF16.gguf", 1, "mmproj-Qwythos-27B-F16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_0", "Q5_K_M", "Q4_K_M", "BF16"}},

    {"empero-ai/Qwythos-9B-v2-GGUF",
@@ -1199,6 +1224,20 @@ static const real_repo real_repos[] = {
      {"BF16", "", 0, "", "", "", "eagle3-gpt-oss-120b-BF16.gguf", ""}},
     {"MXFP4"}},

+    // F16 and BF16 in one repo, the F16 tag must not pick the BF16 file
+    {"ggml-org/Llama-2-7B-GGUF",
+     {{"Llama-2-7B-BF16.gguf"},
+      {"Llama-2-7B-F16.gguf"},
+      {"Llama-2-7B-Q4_0.gguf"},
+      {"Llama-2-7B-Q8_0.gguf"}},
+     {
+      {"", "Llama-2-7B-Q8_0.gguf", 1, "", "", "", "", ""},
+      {"Q8_0", "Llama-2-7B-Q8_0.gguf", 1, "", "", "", "", ""},
+      {"Q4_0", "Llama-2-7B-Q4_0.gguf", 1, "", "", "", "", ""},
+      {"F16", "Llama-2-7B-F16.gguf", 1, "", "", "", "", ""},
+      {"BF16", "Llama-2-7B-BF16.gguf", 1, "", "", "", "", ""}},
+     {"F16", "Q4_0", "Q8_0", "BF16"}},
+
    {"ggml-org/MiMo-V2.6-Distill-Qwen-9B-GGUF",
     {{"MiMo-V2.6-Distill-Qwen-9B-Q8_0.gguf"},
      {"mmproj-MiMo-V2.6-Distill-Qwen-9B-Q8_0.gguf"}},
@@ -1412,7 +1451,7 @@ static const real_repo real_repos[] = {
      {"Q4_K_M", "Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-Q4_K_M.gguf", 1, "mmproj-Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""},
      {"Q8_0", "Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-Q8_0.gguf", 1, "mmproj-Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""},
      {"BF16", "Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-BF16.gguf", 1, "mmproj-Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""},
-      {"F16", "Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-BF16.gguf", 1, "mmproj-Qwen3.5-27B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_0", "Q4_K_M", "Q3_K_M", "BF16"}},

    {"HauhauCS/Qwen3.5-2B-Uncensored-HauhauCS-Aggressive",
@@ -1426,7 +1465,7 @@ static const real_repo real_repos[] = {
      {"Q4_K_M", "Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-Q4_K_M.gguf", 1, "mmproj-Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""},
      {"Q8_0", "Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-Q8_0.gguf", 1, "mmproj-Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""},
      {"BF16", "Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-BF16.gguf", 1, "mmproj-Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""},
-      {"F16", "Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-BF16.gguf", 1, "mmproj-Qwen3.5-2B-Uncensored-HauhauCS-Aggressive-f16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_0", "Q6_K", "Q4_K_M", "BF16"}},

    {"HauhauCS/Qwen3.5-4B-Uncensored-HauhauCS-Aggressive",
@@ -1622,7 +1661,7 @@ static const real_repo real_repos[] = {
      {"IQ2_M", "Qwen3.8-27B-Uncensored-IQ2_M.gguf", 1, "mmproj-Qwen3.8-27B-Uncensored-F16.gguf", "", "", "", ""},
      {"Q4_0", "Qwen3.8-27B-Uncensored-draft-Q4_0.gguf", 1, "mmproj-Qwen3.8-27B-Uncensored-F16.gguf", "", "", "", ""},
      {"F16", "", 0, "", "", "", "", ""},
-      {"mtp", "Qwen3.8-27B-Uncensored-noMTP-IQ2_M.gguf", 1, "mmproj-Qwen3.8-27B-Uncensored-F16.gguf", "", "", "", ""}},
+      {"mtp", "", 0, "", "", "", "", ""}},
     {"Q8_0", "Q5_K_M", "Q4_K_M", "IQ2_M", "Q4_0"}},

    {"kai-os/Carnice-V2-27b-GGUF",
@@ -1903,7 +1942,7 @@ static const real_repo real_repos[] = {
      {"Q4_K_M", "Huihui-Qwen3-VL-4B-Instruct-abliterated-Q4_K_M.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"Q8_0", "Huihui-Qwen3-VL-4B-Instruct-abliterated-Q8_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "Huihui-Qwen3-VL-4B-Instruct-abliterated-BF16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "Huihui-Qwen3-VL-4B-Instruct-abliterated-BF16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_0", "Q4_K_M", "IQ4_XS", "BF16"}},

    {"OBLITERATUS/Gemma-4-12B-OBLITERATED",
@@ -2291,7 +2330,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "embeddinggemma-2-Q8_0.gguf", 1, "mmproj-Q8_0.gguf", "", "", "", ""},
      {"BF16", "embeddinggemma-2-BF16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q6_K_XL", "embeddinggemma-2-UD-Q6_K_XL.gguf", 1, "mmproj-Q8_0.gguf", "", "", "", ""},
-      {"F16", "embeddinggemma-2-BF16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q6_K_XL", "Q4_K_XL", "Q8_0", "BF16"}},

    {"unsloth/GLM-4.7-Flash-GGUF",
@@ -2320,7 +2359,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Q8_0/GLM-5.3-Flash-Q8_0-00001-of-00008.gguf", 8, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/GLM-5.3-Flash-BF16-00001-of-00014.gguf", 14, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q6_K_XL", "UD-Q6_K_XL/GLM-5.3-Flash-UD-Q6_K_XL-00001-of-00007.gguf", 7, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/GLM-5.3-Flash-BF16-00001-of-00014.gguf", 14, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q6_K_XL", "IQ4_XS", "Q8_0", "BF16"}},

    {"unsloth/GLM-5.3-GGUF",
@@ -2359,7 +2398,7 @@ static const real_repo real_repos[] = {
      {"", "BF16/Kimi-K2.6-BF16-00001-of-00046.gguf", 46, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/Kimi-K2.6-BF16-00001-of-00046.gguf", 46, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q8_K_XL", "UD-Q8_K_XL/Kimi-K2.6-UD-Q8_K_XL-00001-of-00014.gguf", 14, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/Kimi-K2.6-BF16-00001-of-00046.gguf", 46, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_K_XL", "Q4_K_XL", "BF16"}},

    {"unsloth/Kimi-K2.7-Code-GGUF",
@@ -2405,7 +2444,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "LFM2.5-VL-3B-Q8_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "LFM2.5-VL-3B-BF16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q8_K_XL", "LFM2.5-VL-3B-UD-Q8_K_XL.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "LFM2.5-VL-3B-BF16.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_K_XL", "Q8_0", "Q5_K_S", "Q4_K_M", "BF16"}},

    {"unsloth/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-GGUF",
@@ -2449,7 +2488,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Qwen3-VL-30B-A3B-Instruct-Q8_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/Qwen3-VL-30B-A3B-Instruct-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-TQ1_0", "Qwen3-VL-30B-A3B-Instruct-UD-TQ1_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/Qwen3-VL-30B-A3B-Instruct-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"TQ1_0", "Q8_0", "Q6_K", "Q4_K_M", "BF16"}},

    {"unsloth/Qwen3.5-27B-GGUF",
@@ -2467,7 +2506,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Qwen3.5-27B-Q8_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/Qwen3.5-27B-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q8_K_XL", "Qwen3.5-27B-UD-Q8_K_XL.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/Qwen3.5-27B-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_K_XL", "Q8_0", "Q6_K", "Q4_K_M", "BF16"}},

    {"unsloth/Qwen3.5-35B-A3B-GGUF",
@@ -2485,7 +2524,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Qwen3.5-35B-A3B-Q8_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/Qwen3.5-35B-A3B-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q8_K_XL", "Qwen3.5-35B-A3B-UD-Q8_K_XL.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/Qwen3.5-35B-A3B-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_K_XL", "IQ3_S", "Q8_0", "Q4_K_M", "BF16"}},

    {"unsloth/Qwen3.5-397B-A17B-GGUF",
@@ -2503,7 +2542,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Q8_0/Qwen3.5-397B-A17B-Q8_0-00001-of-00010.gguf", 10, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/Qwen3.5-397B-A17B-BF16-00001-of-00017.gguf", 17, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q8_K_XL", "UD-Q8_K_XL/Qwen3.5-397B-A17B-UD-Q8_K_XL-00001-of-00010.gguf", 10, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/Qwen3.5-397B-A17B-BF16-00001-of-00017.gguf", 17, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_K_XL", "IQ3_XXS", "Q8_0", "Q4_K_M", "BF16"}},

    {"unsloth/Qwen3.6-35B-A3B-MTP-GGUF",
@@ -2519,7 +2558,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Qwen3.6-35B-A3B-Q8_0.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/Qwen3.6-35B-A3B-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q8_K_XL", "Qwen3.6-35B-A3B-UD-Q8_K_XL.gguf", 1, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/Qwen3.6-35B-A3B-BF16-00001-of-00002.gguf", 2, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_K_XL", "Q3_K_M", "Q8_0", "BF16"}},

    {"unsloth/Qwen3.8-Flash-Next-GGUF",
@@ -2557,7 +2596,7 @@ static const real_repo real_repos[] = {
      {"Q8_0", "Q8_0/Step-3.7-Flash-Q8_0-00001-of-00006.gguf", 6, "mmproj-BF16.gguf", "", "", "", ""},
      {"BF16", "BF16/Step-3.7-Flash-BF16-00001-of-00008.gguf", 8, "mmproj-BF16.gguf", "", "", "", ""},
      {"UD-Q8_K_XL", "UD-Q8_K_XL/Step-3.7-Flash-UD-Q8_K_XL-00001-of-00006.gguf", 6, "mmproj-BF16.gguf", "", "", "", ""},
-      {"F16", "BF16/Step-3.7-Flash-BF16-00001-of-00008.gguf", 8, "mmproj-BF16.gguf", "", "", "", ""}},
+      {"F16", "", 0, "", "", "", "", ""}},
     {"Q8_K_XL", "Q3_K_M", "Q8_0", "BF16"}},

    {"vmlinux/Qwen3.8-Flash-Next-Uncensored-Gufo-Q4Mix-GGUF",
@@ -182,6 +182,8 @@ struct lora_merge_ctx {
    ggml_type get_out_tensor_type(struct ggml_tensor * t) {
        if (t->type == GGML_TYPE_F32) {
            return GGML_TYPE_F32;
+        } else if (t->type == GGML_TYPE_BF16) {
+            return GGML_TYPE_BF16;
        } else {
            return GGML_TYPE_F16;
        }
@@ -409,8 +411,7 @@ struct lora_merge_ctx {

 static void print_usage(int, char ** argv) {
    printf("\nexample usage:\n");
-    printf("\n  %s -m base-model.gguf --lora lora-file.gguf -o merged-model-f16.gguf\n", argv[0]);
-    printf("\nNOTE: output model is F16\n");
+    printf("\n  %s -m base-model.gguf --lora lora-file.gguf -o merged-model.gguf\n", argv[0]);
    printf("\n");
 }

@@ -70,6 +70,7 @@ add_library(mtmd
            models/mimo-audio.cpp
            models/qwen3tts-spkenc.cpp
            models/qwen3tts-gen.cpp
+            models/drama3-gen.cpp
            models/pockettts-seanet.cpp
            models/pockettts-spkenc.cpp
            models/pockettts-gen.cpp
@@ -261,6 +261,13 @@
 #define TN_A_GEN_CODE_MASK_EMBD "a.gen.code.mask_embd"
 #define TN_ATTN_GATE            "%s.blk.%d.attn_gate.%s"

+// drama3 tower, the rest of the block uses the generic names with prefix "a.gen.code"
+#define TN_A_GEN_CODE_ENORM        "a.gen.code.blk.%d.enorm.%s"
+#define TN_A_GEN_CODE_HNORM        "a.gen.code.blk.%d.hnorm.%s"
+#define TN_A_GEN_CODE_EH_PROJ      "a.gen.code.blk.%d.eh_proj.%s"
+#define TN_A_GEN_CODE_BLK_EMBD     "a.gen.code.blk.%d.embd.%s"
+#define TN_A_GEN_CODE_BLK_OUT_NORM "a.gen.code.blk.%d.out_norm.%s"
+
 // qwen3tts code2wav (RVQ codes -> raw PCM)
 // pre_transformer layers use the generic TN_ATTN_*/TN_FFN_*/TN_LN_*/TN_LS_* macros, prefix "a.gen.wav.tfm"
 #define TN_A_GEN_WAV_QUANT_FIRST_IN  "a.gen.wav.quant.first.in_proj.%s"
@@ -595,6 +602,7 @@ enum projector_type {
    PROJECTOR_TYPE_POCKETTTS_GEN,
    PROJECTOR_TYPE_SOPRANO,
    PROJECTOR_TYPE_LFM3A_GEN,
+    PROJECTOR_TYPE_DRAMA3_GEN,
    PROJECTOR_TYPE_MUSE_GLIMMER,
    PROJECTOR_TYPE_SENSENOVA_U1,
    PROJECTOR_TYPE_COHERE2V,
@@ -679,6 +687,7 @@ static std::map<projector_type, std::string> PROJECTOR_TYPE_NAMES = {
    { PROJECTOR_TYPE_POCKETTTS_GEN,     "pockettts_gen"},
    { PROJECTOR_TYPE_SOPRANO,           "soprano"},
    { PROJECTOR_TYPE_LFM3A_GEN,         "lfm3a_gen"},
+    { PROJECTOR_TYPE_DRAMA3_GEN,        "drama3_gen"},
    { PROJECTOR_TYPE_MUSE_GLIMMER,      "muse-glimmer"},
    { PROJECTOR_TYPE_SENSENOVA_U1,      "sensenova_u1"},
    { PROJECTOR_TYPE_COHERE2V,          "cohere2v"},
@@ -841,6 +841,15 @@ struct clip_chatterbox {
    ggml_tensor * spk_affine_b = nullptr;
 };

+// drama3 tower: extra tensors around the generic block
+struct clip_drama3_layer {
+    ggml_tensor * enorm_w    = nullptr;
+    ggml_tensor * hnorm_w    = nullptr;
+    ggml_tensor * eh_proj_w  = nullptr; // [code embd; hidden] -> hidden
+    ggml_tensor * embd_w     = nullptr;
+    ggml_tensor * out_norm_w = nullptr;
+};
+
 struct clip_vocos {
    ggml_tensor * input_w        = nullptr;
    ggml_tensor * input_b        = nullptr;
@@ -1107,6 +1116,7 @@ struct clip_model {

    // qwen3tts code2wav: RVQ codes -> raw PCM
    clip_code2wav c2w;
+    std::vector<clip_drama3_layer> drama3_layers;

    // pocket-tts: SEANet stack, shared by the encoder (speaker path) and the decoder (gen path)
    clip_seanet seanet;
@@ -1237,6 +1237,15 @@ static std::unique_ptr<clip_graph> clip_get_graph_builder(clip_ctx * ctx, const
                }
                builder = std::make_unique<clip_graph_lfm3a_gen>(ctx, img, gen_process, all_known, n_frames);
            } break;
+        case PROJECTOR_TYPE_DRAMA3_GEN:
+            {
+                const auto gen_process = params ? params->gen_process : CLIP_GEN_PROCESS_GEN_CODE;
+                const int64_t n_cb = ctx->model.gen_code_embd_w->ne[2];
+                const int n_frames = params && params->codes ? (int) (params->codes->size() / n_cb) : 1;
+                builder = std::make_unique<clip_graph_drama3_gen>(ctx, img, gen_process, params ? params->gen_state : nullptr, n_frames,
+                                                                  params ? params->top_k : 30, params ? params->top_p : 0.9f,
+                                                                  params ? params->temp : 1.0f);
+            } break;
        case PROJECTOR_TYPE_QWEN3TTS_GEN:
            {
                const auto  gen_process = params ? params->gen_process : CLIP_GEN_PROCESS_GEN_CODE;
@@ -2257,6 +2266,22 @@ struct clip_model_loader {
                        // matches the reference decoder's sliding_window (speech_tokenizer/config.json)
                        hparams.wav_tfm_swa = 72;
                    } break;
+                case PROJECTOR_TYPE_DRAMA3_GEN:
+                    {
+                        // Fish DAC decoder, codec.pth has no config
+                        hparams.audio_sample_rate    = 44100;
+                        hparams.wav_tfm_n_layer      = 8;
+                        hparams.wav_tfm_n_embd       = 1024;
+                        hparams.wav_tfm_n_ff         = 3072;
+                        hparams.wav_tfm_n_head       = 16;
+                        hparams.wav_tfm_n_head_kv    = 16;
+                        hparams.wav_tfm_eps          = 1e-5f;
+                        hparams.wav_tfm_rope_theta   = 10000.0f;
+                        hparams.wav_tfm_swa          = 128;
+                        hparams.wav_upsample_n_block = 2;
+                        hparams.wav_dac_n_block      = 4;
+                        hparams.wav_dac_n_res        = 3;
+                    } break;
                case PROJECTOR_TYPE_POCKETTTS_SPKENC:
                case PROJECTOR_TYPE_POCKETTTS_GEN:
                    {
@@ -2785,6 +2810,92 @@ struct clip_model_loader {
            }
        }

+        // code2wav: RVQ codes -> raw PCM, drama3 has no quant/pre_conv/tfm projections
+        auto load_code2wav = [&](bool with_proj) {
+            auto & c2w = model.c2w;
+
+            c2w.quant_first_in_w  = get_tensor(string_format(TN_A_GEN_WAV_QUANT_FIRST_IN,  "weight"), with_proj);
+            c2w.quant_first_out_w = get_tensor(string_format(TN_A_GEN_WAV_QUANT_FIRST_OUT, "weight"), with_proj);
+            c2w.quant_first_cb_w  = get_tensor(string_format(TN_A_GEN_WAV_QUANT_FIRST_CB,  "weight"));
+            c2w.quant_rest_in_w   = get_tensor(string_format(TN_A_GEN_WAV_QUANT_REST_IN,   "weight"), with_proj);
+            c2w.quant_rest_out_w  = get_tensor(string_format(TN_A_GEN_WAV_QUANT_REST_OUT,  "weight"), with_proj);
+            c2w.quant_rest_cb_w   = get_tensor(string_format(TN_A_GEN_WAV_QUANT_REST_CB,   "weight"));
+
+            c2w.pre_conv_w = get_tensor(string_format(TN_A_GEN_WAV_PRE_CONV, "weight"), with_proj);
+            c2w.pre_conv_b = get_tensor(string_format(TN_A_GEN_WAV_PRE_CONV, "bias"),   with_proj);
+
+            c2w.tfm_in_proj_w     = get_tensor(string_format(TN_A_GEN_WAV_TFM_IN_PROJ,  "weight"), with_proj);
+            c2w.tfm_in_proj_b     = get_tensor(string_format(TN_A_GEN_WAV_TFM_IN_PROJ,  "bias"), with_proj);
+            c2w.tfm_out_proj_w    = get_tensor(string_format(TN_A_GEN_WAV_TFM_OUT_PROJ, "weight"), with_proj);
+            c2w.tfm_out_proj_b    = get_tensor(string_format(TN_A_GEN_WAV_TFM_OUT_PROJ, "bias"), with_proj);
+            c2w.tfm_output_norm_w = get_tensor(string_format(TN_A_GEN_WAV_TFM_OUT_NORM, "weight"));
+
+            // loaded manually, the generic model.layers loop is taken by code_predictor
+            c2w.tfm_layers.resize(hparams.wav_tfm_n_layer);
+            for (int il = 0; il < hparams.wav_tfm_n_layer; il++) {
+                auto & layer = c2w.tfm_layers[il];
+                const char * p = "a.gen.wav.tfm";
+                layer.q_w      = get_tensor(string_format(TN_ATTN_Q,      p, il, "weight"));
+                layer.k_w      = get_tensor(string_format(TN_ATTN_K,      p, il, "weight"));
+                layer.v_w      = get_tensor(string_format(TN_ATTN_V,      p, il, "weight"));
+                layer.o_w      = get_tensor(string_format(TN_ATTN_OUTPUT, p, il, "weight"));
+                layer.ln_1_w   = get_tensor(string_format(TN_LN_1,        p, il, "weight"));
+                layer.ln_2_w   = get_tensor(string_format(TN_LN_2,        p, il, "weight"));
+                layer.ls_1_w   = get_tensor(string_format(TN_LS_1,        p, il, "weight"));
+                layer.ls_2_w   = get_tensor(string_format(TN_LS_2,        p, il, "weight"));
+                layer.ff_gate_w = get_tensor(string_format(TN_FFN_GATE,   p, il, "weight"));
+                layer.ff_up_w   = get_tensor(string_format(TN_FFN_UP,     p, il, "weight"));
+                layer.ff_down_w = get_tensor(string_format(TN_FFN_DOWN,   p, il, "weight"));
+            }
+
+            // upsample: 2x (causal ConvTranspose1d + ConvNeXt block)
+            c2w.upsample.resize(hparams.wav_upsample_n_block);
+            for (int il = 0; il < hparams.wav_upsample_n_block; il++) {
+                auto & up = c2w.upsample[il];
+                up.conv_w   = get_tensor(string_format(TN_A_GEN_WAV_UP_CONV,   il, "weight"));
+                up.conv_b   = get_tensor(string_format(TN_A_GEN_WAV_UP_CONV,   il, "bias"));
+                up.dwconv_w = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "weight"));
+                up.dwconv_b = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "bias"));
+                up.norm_w   = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM,   il, "weight"));
+                up.norm_b   = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM,   il, "bias"));
+                up.pw1_w    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1,    il, "weight"));
+                up.pw1_b    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1,    il, "bias"));
+                up.pw2_w    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2,    il, "weight"));
+                up.pw2_b    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2,    il, "bias"));
+                up.gamma    = get_tensor(string_format(TN_A_GEN_WAV_UP_GAMMA,  il));
+            }
+
+            // DAC decoder: conv_pre + n upsample blocks (each with n_res residual units) + conv_post
+            c2w.dac_entry_w = get_tensor(string_format(TN_A_GEN_WAV_DAC_ENTRY, "weight"));
+            c2w.dac_entry_b = get_tensor(string_format(TN_A_GEN_WAV_DAC_ENTRY, "bias"));
+
+            c2w.dac.resize(hparams.wav_dac_n_block);
+            for (int il = 0; il < hparams.wav_dac_n_block; il++) {
+                auto & blk = c2w.dac[il];
+                blk.snake_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_SNAKE, il, "alpha"));
+                blk.snake_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_SNAKE, il, "beta"));
+                blk.conv_w      = get_tensor(string_format(TN_A_GEN_WAV_DAC_CONV,  il, "weight"));
+                blk.conv_b      = get_tensor(string_format(TN_A_GEN_WAV_DAC_CONV,  il, "bias"));
+
+                blk.res.resize(hparams.wav_dac_n_res);
+                for (int ir = 0; ir < hparams.wav_dac_n_res; ir++) {
+                    auto & res = blk.res[ir];
+                    res.act1_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT1,  il, ir, "alpha"));
+                    res.act1_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT1,  il, ir, "beta"));
+                    res.conv1_w    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV1, il, ir, "weight"));
+                    res.conv1_b    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV1, il, ir, "bias"));
+                    res.act2_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT2,  il, ir, "alpha"));
+                    res.act2_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT2,  il, ir, "beta"));
+                    res.conv2_w    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV2, il, ir, "weight"));
+                    res.conv2_b    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV2, il, ir, "bias"));
+                }
+            }
+
+            c2w.dac_post_snake_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_SNAKE, "alpha"));
+            c2w.dac_post_snake_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_SNAKE, "beta"));
+            c2w.dac_post_conv_w      = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_CONV,  "weight"));
+            c2w.dac_post_conv_b      = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_CONV,  "bias"));
+        };

        switch (model.proj_type) {
            case PROJECTOR_TYPE_MLP:
@@ -3656,91 +3767,25 @@ struct clip_model_loader {
                    model.gen_code_norm_w     = get_tensor(string_format(TN_A_GEN_CODE_NORM,     "weight"));

                    // code2wav: RVQ codes -> raw PCM, lives in the same ctx as code_predictor
-                    {
-                        auto & c2w = model.c2w;
-
-                        c2w.quant_first_in_w  = get_tensor(string_format(TN_A_GEN_WAV_QUANT_FIRST_IN,  "weight"));
-                        c2w.quant_first_out_w = get_tensor(string_format(TN_A_GEN_WAV_QUANT_FIRST_OUT, "weight"));
-                        c2w.quant_first_cb_w  = get_tensor(string_format(TN_A_GEN_WAV_QUANT_FIRST_CB,  "weight"));
-                        c2w.quant_rest_in_w   = get_tensor(string_format(TN_A_GEN_WAV_QUANT_REST_IN,   "weight"));
-                        c2w.quant_rest_out_w  = get_tensor(string_format(TN_A_GEN_WAV_QUANT_REST_OUT,  "weight"));
-                        c2w.quant_rest_cb_w   = get_tensor(string_format(TN_A_GEN_WAV_QUANT_REST_CB,   "weight"));
-
-                        c2w.pre_conv_w = get_tensor(string_format(TN_A_GEN_WAV_PRE_CONV, "weight"));
-                        c2w.pre_conv_b = get_tensor(string_format(TN_A_GEN_WAV_PRE_CONV, "bias"));
-
-                        c2w.tfm_in_proj_w     = get_tensor(string_format(TN_A_GEN_WAV_TFM_IN_PROJ,  "weight"));
-                        c2w.tfm_in_proj_b     = get_tensor(string_format(TN_A_GEN_WAV_TFM_IN_PROJ,  "bias"));
-                        c2w.tfm_out_proj_w    = get_tensor(string_format(TN_A_GEN_WAV_TFM_OUT_PROJ, "weight"));
-                        c2w.tfm_out_proj_b    = get_tensor(string_format(TN_A_GEN_WAV_TFM_OUT_PROJ, "bias"));
-                        c2w.tfm_output_norm_w = get_tensor(string_format(TN_A_GEN_WAV_TFM_OUT_NORM, "weight"));
-
-                        // loaded manually, the generic model.layers loop is taken by code_predictor
-                        c2w.tfm_layers.resize(hparams.wav_tfm_n_layer);
-                        for (int il = 0; il < hparams.wav_tfm_n_layer; il++) {
-                            auto & layer = c2w.tfm_layers[il];
-                            const char * p = "a.gen.wav.tfm";
-                            layer.q_w      = get_tensor(string_format(TN_ATTN_Q,      p, il, "weight"));
-                            layer.k_w      = get_tensor(string_format(TN_ATTN_K,      p, il, "weight"));
-                            layer.v_w      = get_tensor(string_format(TN_ATTN_V,      p, il, "weight"));
-                            layer.o_w      = get_tensor(string_format(TN_ATTN_OUTPUT, p, il, "weight"));
-                            layer.ln_1_w   = get_tensor(string_format(TN_LN_1,        p, il, "weight"));
-                            layer.ln_2_w   = get_tensor(string_format(TN_LN_2,        p, il, "weight"));
-                            layer.ls_1_w   = get_tensor(string_format(TN_LS_1,        p, il, "weight"));
-                            layer.ls_2_w   = get_tensor(string_format(TN_LS_2,        p, il, "weight"));
-                            layer.ff_gate_w = get_tensor(string_format(TN_FFN_GATE,   p, il, "weight"));
-                            layer.ff_up_w   = get_tensor(string_format(TN_FFN_UP,     p, il, "weight"));
-                            layer.ff_down_w = get_tensor(string_format(TN_FFN_DOWN,   p, il, "weight"));
-                        }
-
-                        // upsample: 2x (causal ConvTranspose1d + ConvNeXt block)
-                        c2w.upsample.resize(hparams.wav_upsample_n_block);
-                        for (int il = 0; il < hparams.wav_upsample_n_block; il++) {
-                            auto & up = c2w.upsample[il];
-                            up.conv_w   = get_tensor(string_format(TN_A_GEN_WAV_UP_CONV,   il, "weight"));
-                            up.conv_b   = get_tensor(string_format(TN_A_GEN_WAV_UP_CONV,   il, "bias"));
-                            up.dwconv_w = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "weight"));
-                            up.dwconv_b = get_tensor(string_format(TN_A_GEN_WAV_UP_DWCONV, il, "bias"));
-                            up.norm_w   = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM,   il, "weight"));
-                            up.norm_b   = get_tensor(string_format(TN_A_GEN_WAV_UP_NORM,   il, "bias"));
-                            up.pw1_w    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1,    il, "weight"));
-                            up.pw1_b    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW1,    il, "bias"));
-                            up.pw2_w    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2,    il, "weight"));
-                            up.pw2_b    = get_tensor(string_format(TN_A_GEN_WAV_UP_PW2,    il, "bias"));
-                            up.gamma    = get_tensor(string_format(TN_A_GEN_WAV_UP_GAMMA,  il));
-                        }
-
-                        // DAC decoder: conv_pre + n upsample blocks (each with n_res residual units) + conv_post
-                        c2w.dac_entry_w = get_tensor(string_format(TN_A_GEN_WAV_DAC_ENTRY, "weight"));
-                        c2w.dac_entry_b = get_tensor(string_format(TN_A_GEN_WAV_DAC_ENTRY, "bias"));
-
-                        c2w.dac.resize(hparams.wav_dac_n_block);
-                        for (int il = 0; il < hparams.wav_dac_n_block; il++) {
-                            auto & blk = c2w.dac[il];
-                            blk.snake_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_SNAKE, il, "alpha"));
-                            blk.snake_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_SNAKE, il, "beta"));
-                            blk.conv_w      = get_tensor(string_format(TN_A_GEN_WAV_DAC_CONV,  il, "weight"));
-                            blk.conv_b      = get_tensor(string_format(TN_A_GEN_WAV_DAC_CONV,  il, "bias"));
-
-                            blk.res.resize(hparams.wav_dac_n_res);
-                            for (int ir = 0; ir < hparams.wav_dac_n_res; ir++) {
-                                auto & res = blk.res[ir];
-                                res.act1_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT1,  il, ir, "alpha"));
-                                res.act1_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT1,  il, ir, "beta"));
-                                res.conv1_w    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV1, il, ir, "weight"));
-                                res.conv1_b    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV1, il, ir, "bias"));
-                                res.act2_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT2,  il, ir, "alpha"));
-                                res.act2_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_ACT2,  il, ir, "beta"));
-                                res.conv2_w    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV2, il, ir, "weight"));
-                                res.conv2_b    = get_tensor(string_format(TN_A_GEN_WAV_DAC_RES_CONV2, il, ir, "bias"));
-                            }
-                        }
-
-                        c2w.dac_post_snake_alpha = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_SNAKE, "alpha"));
-                        c2w.dac_post_snake_beta  = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_SNAKE, "beta"));
-                        c2w.dac_post_conv_w      = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_CONV,  "weight"));
-                        c2w.dac_post_conv_b      = get_tensor(string_format(TN_A_GEN_WAV_DAC_POST_CONV,  "bias"));
+                    load_code2wav(true);
+                } break;
+            case PROJECTOR_TYPE_DRAMA3_GEN:
+                {
+                    // generic a.gen.code.blk.* tensors are loaded by the loop above
+                    model.drama3_layers.resize(hparams.n_layer);
+                    for (int il = 0; il < hparams.n_layer; ++il) {
+                        auto & tl = model.drama3_layers[il];
+                        model.layers[il].attn_gate_w = get_tensor(string_format(TN_ATTN_GATE, prefix, il, "weight"));
+                        tl.enorm_w    = get_tensor(string_format(TN_A_GEN_CODE_ENORM,        il, "weight"));
+                        tl.hnorm_w    = get_tensor(string_format(TN_A_GEN_CODE_HNORM,        il, "weight"));
+                        tl.eh_proj_w  = get_tensor(string_format(TN_A_GEN_CODE_EH_PROJ,      il, "weight"));
+                        tl.embd_w     = get_tensor(string_format(TN_A_GEN_CODE_BLK_EMBD,     il, "weight"));
+                        tl.out_norm_w = get_tensor(string_format(TN_A_GEN_CODE_BLK_OUT_NORM, il, "weight"));
                    }
+                    model.gen_code_embd_w     = get_tensor(string_format(TN_A_GEN_CODE_EMBD,     "weight")); // merged 3D, one table per codebook
+                    model.gen_code_head_w     = get_tensor(string_format(TN_A_GEN_CODE_HEAD,     "weight")); // merged 3D, codebooks 1..
+                    model.gen_code_out_embd_w = get_tensor(string_format(TN_A_GEN_CODE_OUT_EMBD, "weight")); // backbone rows of the semantic band
+                    load_code2wav(false);
                } break;
            case PROJECTOR_TYPE_CHATTERBOX:
                {
@@ -5520,6 +5565,7 @@ int clip_n_output_tokens(const clip_ctx * ctx, const clip_image_f32 * img) {
            } break;
        case PROJECTOR_TYPE_QWEN3TTS_GEN:
        case PROJECTOR_TYPE_LFM3A_GEN:
+        case PROJECTOR_TYPE_DRAMA3_GEN:
            {
                // one hidden-state vector fed back to the talker per call
                n_patches = 1;
@@ -5587,6 +5633,7 @@ bool clip_image_batch_encode(clip_ctx * ctx, int n_threads, const clip_image_f32
 static std::vector<c2w_state_slot> list_gen_state_slots(const clip_hparams & hparams, const clip_model & model) {
    switch (model.proj_type) {
        case PROJECTOR_TYPE_QWEN3TTS_GEN:  return list_c2w_state_slots(hparams, model);
+        case PROJECTOR_TYPE_DRAMA3_GEN:    return list_drama3_state_slots(hparams, model);
        case PROJECTOR_TYPE_POCKETTTS_GEN: return list_pockettts_state_slots(hparams, model);
        case PROJECTOR_TYPE_LFM3A_GEN:     return list_lfm3a_gen_state_slots(hparams, model);
        default:                           return {};
@@ -5772,6 +5819,14 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
        LOG_ERR("%s: audio gen requires a gen state\n", __func__);
        return false;
    }
+    if (ctx->proj_type() == PROJECTOR_TYPE_DRAMA3_GEN && params->gen_process == CLIP_GEN_PROCESS_GEN_CODE) {
+        float pos = 0.0f;
+        ggml_backend_tensor_get(params->gen_state->slots[0], &pos, 0, sizeof(float));
+        if ((int) pos >= clip_graph_drama3_gen::n_pos_max) {
+            LOG_ERR("%s: the codebook tower holds at most %d rows\n", __func__, clip_graph_drama3_gen::n_pos_max);
+            return false;
+        }
+    }

    // build the inference graph
    ggml_backend_sched_reset(ctx->sched.get());
@@ -5845,6 +5900,10 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
        size_t i = 0;
        for (const auto & slot : list_gen_state_slots(hparams, model)) {
            GGML_ASSERT(i < slots.size());
+            if (slot.in_place) {
+                i++;
+                continue;
+            }
            ggml_backend_tensor_copy(slots[i++], get_inp_tensor(("state_in_" + slot.name).c_str()));
        }
    };
@@ -6900,6 +6959,58 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
                    }
                }
            } break;
+        case PROJECTOR_TYPE_DRAMA3_GEN:
+            {
+                const int64_t n_cb      = model.gen_code_embd_w->ne[2];
+                const int64_t n_vocab   = model.gen_code_embd_w->ne[1];
+                if (params->gen_process == CLIP_GEN_PROCESS_GEN_WAV) {
+                    // frame-major to group-major, clamp: codec residual codebooks are smaller than the tower vocab
+                    GGML_ASSERT(params->codes != nullptr && params->codes->size() % n_cb == 0);
+                    const int64_t n_frames = (int64_t) params->codes->size() / n_cb;
+                    const int64_t n_first  = model.c2w.quant_first_cb_w->ne[1];
+                    const int64_t n_rest   = model.c2w.quant_rest_cb_w->ne[1];
+                    std::vector<int32_t> codes(params->codes->size());
+                    for (int64_t f = 0; f < n_frames; f++) {
+                        for (int64_t g = 0; g < n_cb; g++) {
+                            const int64_t n = g == 0 ? n_first : n_rest;
+                            codes[g * n_frames + f] = std::clamp<int32_t>((*params->codes)[f * n_cb + g], 0, (int32_t) n - 1);
+                        }
+                    }
+                    set_input_i32("inp_codes", codes);
+                    set_gen_state_in();
+                } else {
+                    // code0 < 0: next row is text (prompt), the tower only fills its KV
+                    // codes: input frame codes for a speech row, empty for a text row
+                    const bool next_speech = params->code0 >= 0;
+                    const bool cur_speech  = params->codes && (int64_t) params->codes->size() == n_cb;
+                    if (params->code0 >= n_vocab) {
+                        LOG_ERR("%s: code0 out of range (%d, vocab %lld)\n", __func__, params->code0, (long long) n_vocab);
+                        return false;
+                    }
+                    std::vector<int32_t> code0 = { std::max(params->code0, 0) };
+                    set_input_i32("inp_code0", code0);
+
+                    std::vector<int32_t> extra(n_cb - 2);
+                    for (int64_t g = 2; g < n_cb; g++) {
+                        const int32_t c = cur_speech ? std::clamp<int32_t>((*params->codes)[g], 0, (int32_t) n_vocab - 1) : 0;
+                        extra[g - 2] = c + (int32_t) (g * n_vocab);
+                    }
+                    set_input_i32("inp_extra", extra);
+                    set_input_f32("inp_scale", { next_speech ? 1.0f : 0.0f, cur_speech ? 1.0f : 0.0f });
+
+                    float pos = 0.0f;
+                    ggml_backend_tensor_get(params->gen_state->slots[0], &pos, 0, sizeof(float));
+                    std::vector<int32_t> inp_pos = { (int32_t) pos };
+                    set_input_i32("inp_pos", inp_pos);
+
+                    std::uniform_real_distribution<float> dist(0.0f, 1.0f);
+                    std::vector<float> r(model.gen_code_head_w->ne[2]);
+                    for (auto & v : r) {
+                        v = dist(params->gen_state->rng);
+                    }
+                    set_input_f32("inp_rand", r);
+                }
+            } break;
        case PROJECTOR_TYPE_QWEN3TTS_GEN:
            {
                if (params->gen_process == CLIP_GEN_PROCESS_GEN_WAV) {
@@ -7700,13 +7811,15 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
            clip_graph_soprano::decode_spectrum(model, spectrum, out_audio);
        }

-        // drop the tail audio that comes from the code-0 rear padding
-        const int64_t n_codes    = params->codes ? model.gen_code_head_w->ne[2] + 1 : 0;
-        const int64_t n_frames_w = hparams.wav_tfm_swa;
-        const int64_t n_frames   = params->codes ? (int64_t) params->codes->size() / n_codes : n_frames_w;
-        if (n_frames < n_frames_w) {
-            const size_t hop = out_audio.size() / n_frames_w;
-            out_audio.resize((size_t) n_frames * hop);
+        // drop the tail audio that comes from the code-0 rear padding, drama3 has no padding
+        if (model.proj_type != PROJECTOR_TYPE_DRAMA3_GEN) {
+            const int64_t n_codes    = params->codes ? model.gen_code_head_w->ne[2] + 1 : 0;
+            const int64_t n_frames_w = hparams.wav_tfm_swa;
+            const int64_t n_frames   = params->codes ? (int64_t) params->codes->size() / n_codes : n_frames_w;
+            if (n_frames < n_frames_w) {
+                const size_t hop = out_audio.size() / n_frames_w;
+                out_audio.resize((size_t) n_frames * hop);
+            }
        }
    }
    if (params->enc_state != nullptr && model.proj_type == PROJECTOR_TYPE_LFM3A) {
@@ -7744,6 +7857,10 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
        const auto & slots = params->gen_state->slots;
        size_t i = 0;
        for (const auto & slot : list_gen_state_slots(hparams, model)) {
+            if (slot.in_place) {
+                i++;
+                continue;
+            }
            ggml_tensor * t = ggml_graph_get_tensor(gf, ("state_out_" + slot.name).c_str());
            if (t == nullptr) {
                GGML_ABORT("graph has no \"state_out_%s\" tensor", slot.name.c_str());
@@ -7924,6 +8041,7 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) {
            return ctx->model.mm_fc_w->ne[2];
        case PROJECTOR_TYPE_QWEN3TTS_GEN:
        case PROJECTOR_TYPE_LFM3A_GEN:
+        case PROJECTOR_TYPE_DRAMA3_GEN:
            return ctx->model.gen_code_out_embd_w->ne[0];
        case PROJECTOR_TYPE_POCKETTTS_SPKENC:
            return ctx->model.spk_proj_w->ne[1];
@@ -0,0 +1,256 @@
+#include "models.h"
+
+#include <cmath>
+#include <string>
+
+// tower RoPE: interleaved, first 64 dims of each head
+static constexpr int   DRAMA3_N_ROT      = 64;
+static constexpr float DRAMA3_ROPE_THETA = 1e7f;
+
+std::vector<c2w_state_slot> list_drama3_state_slots(const clip_hparams & hparams, const clip_model & model) {
+    std::vector<c2w_state_slot> slots;
+    slots.push_back({"tower_pos", 1, 1, true});
+    for (size_t il = 0; il < model.layers.size(); il++) {
+        const int64_t n_kv = model.layers[il].k_w->ne[1];
+        slots.push_back({"tower_k_" + std::to_string(il), n_kv, clip_graph_drama3_gen::n_pos_max, true});
+        slots.push_back({"tower_v_" + std::to_string(il), n_kv, clip_graph_drama3_gen::n_pos_max, true});
+    }
+    for (auto & slot : list_c2w_state_slots(hparams, model)) {
+        slots.push_back(slot);
+    }
+    return slots;
+}
+
+// h_out = out_norm(Block(eh_proj([enorm(mask * E[code] + extra); hnorm(h_in)])))
+ggml_tensor * clip_graph_drama3_gen::tower_layer(ggml_tensor * x, int il, ggml_tensor * code, ggml_tensor * mask, ggml_tensor * extra,
+                                                 ggml_tensor * inp_pos, ggml_tensor * k_slot, ggml_tensor * v_slot, int pos) {
+    const auto & layer = model.layers[il];
+    const auto & tl    = model.drama3_layers[il];
+    const int     n_head    = hparams.n_head;
+    const int     n_head_kv = hparams.n_head_kv;
+    const int64_t d_head    = layer.q_w->ne[1] / n_head;
+    const int64_t n_kv      = d_head * n_head_kv;
+    const float   eps       = hparams.eps;
+
+    ggml_tensor * e = ggml_mul(ctx0, ggml_get_rows(ctx0, tl.embd_w, code), mask);
+    if (extra) {
+        e = ggml_add(ctx0, e, extra);
+    }
+    e = ggml_mul(ctx0, ggml_rms_norm(ctx0, e, eps), tl.enorm_w);
+    ggml_tensor * h = ggml_mul(ctx0, ggml_rms_norm(ctx0, x, eps), tl.hnorm_w);
+    ggml_tensor * mixed = ggml_mul_mat(ctx0, tl.eh_proj_w, ggml_concat(ctx0, e, h, 0));
+
+    ggml_tensor * cur = ggml_mul(ctx0, ggml_rms_norm(ctx0, mixed, eps), layer.ln_1_w);
+    ggml_tensor * q    = ggml_reshape_3d(ctx0, ggml_mul_mat(ctx0, layer.q_w, cur), d_head, n_head, 1);
+    ggml_tensor * k    = ggml_reshape_3d(ctx0, ggml_mul_mat(ctx0, layer.k_w, cur), d_head, n_head_kv, 1);
+    ggml_tensor * v    = ggml_mul_mat(ctx0, layer.v_w, cur);
+    ggml_tensor * gate = ggml_mul_mat(ctx0, layer.attn_gate_w, cur);
+
+    q = ggml_mul(ctx0, ggml_rms_norm(ctx0, q, eps), layer.q_norm);
+    k = ggml_mul(ctx0, ggml_rms_norm(ctx0, k, eps), layer.k_norm);
+    q = ggml_rope_ext(ctx0, q, inp_pos, nullptr, DRAMA3_N_ROT, GGML_ROPE_TYPE_NORMAL, 0, DRAMA3_ROPE_THETA, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
+    k = ggml_rope_ext(ctx0, k, inp_pos, nullptr, DRAMA3_N_ROT, GGML_ROPE_TYPE_NORMAL, 0, DRAMA3_ROPE_THETA, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
+
+    // write K/V of this row to the slot, attention uses the cached rows + this row
+    ggml_tensor * k_new = ggml_reshape_2d(ctx0, k, n_kv, 1);
+    ggml_build_forward_expand(gf, ggml_cpy(ctx0, k_new, ggml_view_2d(ctx0, k_slot, n_kv, 1, k_slot->nb[1], (size_t) pos * k_slot->nb[1])));
+    ggml_build_forward_expand(gf, ggml_cpy(ctx0, v,     ggml_view_2d(ctx0, v_slot, n_kv, 1, v_slot->nb[1], (size_t) pos * v_slot->nb[1])));
+    ggml_tensor * k_all = k_new;
+    ggml_tensor * v_all = v;
+    if (pos > 0) {
+        k_all = ggml_concat(ctx0, ggml_view_2d(ctx0, k_slot, n_kv, pos, k_slot->nb[1], 0), k_new, 1);
+        v_all = ggml_concat(ctx0, ggml_view_2d(ctx0, v_slot, n_kv, pos, v_slot->nb[1], 0), v,     1);
+    }
+
+    ggml_tensor * attn = build_attn(nullptr, nullptr,
+                                    ggml_reshape_4d(ctx0, q,     d_head, n_head,    1,       1),
+                                    ggml_reshape_4d(ctx0, k_all, d_head, n_head_kv, pos + 1, 1),
+                                    ggml_reshape_4d(ctx0, v_all, d_head, n_head_kv, pos + 1, 1),
+                                    nullptr, 1.0f / sqrtf((float) d_head), il);
+    attn = ggml_mul(ctx0, attn, ggml_sigmoid(ctx0, gate));
+    cur = ggml_add(ctx0, mixed, ggml_mul_mat(ctx0, layer.o_w, attn));
+
+    ggml_tensor * ff = ggml_mul(ctx0, ggml_rms_norm(ctx0, cur, eps), layer.ln_2_w);
+    ff = ggml_swiglu_split(ctx0, ggml_mul_mat(ctx0, layer.ff_gate_w, ff), ggml_mul_mat(ctx0, layer.ff_up_w, ff));
+    cur = ggml_add(ctx0, cur, ggml_mul_mat(ctx0, layer.ff_down_w, ff));
+
+    cur = ggml_mul(ctx0, ggml_rms_norm(ctx0, cur, eps), tl.out_norm_w);
+    cb(cur, "tower_out", il);
+    return cur;
+}
+
+// one row: backbone hidden + next main code -> other codes of that frame + next backbone input
+ggml_cgraph * clip_graph_drama3_gen::build_tower() {
+    const int     n_layer   = (int) model.layers.size();
+    const int     n_heads   = (int) model.gen_code_head_w->ne[2]; // one per residual codebook
+    const int     n_shared  = n_layer - n_heads;
+    const int64_t n_cb      = model.gen_code_embd_w->ne[2];
+    const int64_t n_embd_cb = model.gen_code_embd_w->ne[0];
+
+    int pos = 0;
+    std::vector<ggml_tensor *> slots;
+    const auto slot_list = list_drama3_state_slots(hparams, model);
+    if (gen_state) {
+        float pos_f = 0.0f;
+        ggml_backend_tensor_get(gen_state->slots[0], &pos_f, 0, sizeof(float));
+        pos = (int) pos_f;
+        slots = gen_state->slots;
+    } else {
+        // buffer reserve: no gen state, use stand-in slots
+        for (const auto & slot : slot_list) {
+            ggml_tensor * t = nullptr;
+            if (slot.in_place) {
+                t = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, slot.ne0, slot.ne1);
+                ggml_set_input(t);
+            }
+            slots.push_back(t);
+        }
+    }
+    GGML_ASSERT(pos < n_pos_max);
+
+    ggml_tensor * h = ggml_new_tensor_1d(ctx0, GGML_TYPE_F32, n_mmproj_embd);
+    ggml_set_name(h, "inp_raw");
+    ggml_set_input(h);
+
+    ggml_tensor * inp_code0 = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, 1);
+    ggml_set_name(inp_code0, "inp_code0");
+    ggml_set_input(inp_code0);
+
+    // codes 2.. of the input frame, as row ids of the merged codebook table
+    ggml_tensor * inp_extra = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_cb - 2);
+    ggml_set_name(inp_extra, "inp_extra");
+    ggml_set_input(inp_extra);
+
+    // [next frame is speech, input frame is speech]
+    ggml_tensor * inp_scale = ggml_new_tensor_1d(ctx0, GGML_TYPE_F32, 2);
+    ggml_set_name(inp_scale, "inp_scale");
+    ggml_set_input(inp_scale);
+
+    ggml_tensor * inp_pos = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, 1);
+    ggml_set_name(inp_pos, "inp_pos");
+    ggml_set_input(inp_pos);
+
+    ggml_tensor * inp_rand = ggml_new_tensor_1d(ctx0, GGML_TYPE_F32, n_heads);
+    ggml_set_name(inp_rand, "inp_rand");
+    ggml_set_input(inp_rand);
+
+    ggml_tensor * mask     = ggml_view_1d(ctx0, inp_scale, 1, 0);
+    ggml_tensor * extra_on = ggml_view_1d(ctx0, inp_scale, 1, sizeof(float));
+
+    ggml_tensor * cb_all = ggml_reshape_2d(ctx0, model.gen_code_embd_w, n_embd_cb, model.gen_code_embd_w->ne[1] * n_cb);
+    ggml_tensor * extra  = ggml_cont(ctx0, ggml_transpose(ctx0, ggml_get_rows(ctx0, cb_all, inp_extra)));
+    extra = ggml_mul(ctx0, ggml_reshape_2d(ctx0, ggml_sum_rows(ctx0, extra), n_embd_cb, 1), extra_on);
+
+    clip_graph_qwen3tts_gen::code_gen cg(*this, top_k, top_p);
+    ggml_tensor * out_codes = cg.cache_set(ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, 1, n_cb), 0, inp_code0);
+
+    ggml_tensor * code = inp_code0;
+    ggml_tensor * x    = ggml_reshape_2d(ctx0, h, n_mmproj_embd, 1);
+    for (int il = 0; il < n_layer; il++) {
+        x = tower_layer(x, il, code, mask, il == 0 ? extra : nullptr, inp_pos, slots[1 + 2 * il], slots[2 + 2 * il], pos);
+        if (il < n_shared) {
+            continue;
+        }
+        const int ih = il - n_shared;
+        ggml_tensor * head   = ggml_view_2d(ctx0, model.gen_code_head_w, model.gen_code_head_w->ne[0], model.gen_code_head_w->ne[1],
+                                            model.gen_code_head_w->nb[1], (size_t) ih * model.gen_code_head_w->nb[2]);
+        ggml_tensor * logits = ggml_mul_mat(ctx0, head, x);
+        if (temp > 0.0f && temp != 1.0f) {
+            logits = ggml_scale(ctx0, logits, 1.0f / temp);
+        }
+        code = cg.do_sampling(logits, ggml_view_1d(ctx0, inp_rand, 1, (size_t) ih * sizeof(float)));
+        out_codes = cg.cache_set(out_codes, ih + 1, code);
+    }
+
+    ggml_tensor * pos_slot = slots[0];
+    ggml_build_forward_expand(gf, ggml_cpy(ctx0, ggml_scale_bias(ctx0, pos_slot, 1.0f, 1.0f), pos_slot));
+
+    out_codes = ggml_cont(ctx0, out_codes);
+    ggml_set_name(out_codes, "out_codes");
+    ggml_set_output(out_codes);
+    ggml_build_forward_expand(gf, out_codes);
+
+    // backbone input of the next row: (E_lm[main] + E_cb1[code 1]) / sqrt(2)
+    ggml_tensor * code1 = ggml_view_1d(ctx0, out_codes, 1, out_codes->nb[1]);
+    ggml_tensor * cb1   = ggml_view_2d(ctx0, model.gen_code_embd_w, n_embd_cb, model.gen_code_embd_w->ne[1],
+                                       model.gen_code_embd_w->nb[1], model.gen_code_embd_w->nb[2]);
+    ggml_tensor * out_embd = ggml_add(ctx0, ggml_get_rows(ctx0, model.gen_code_out_embd_w, inp_code0), ggml_get_rows(ctx0, cb1, code1));
+    out_embd = ggml_scale(ctx0, out_embd, 1.0f / sqrtf(2.0f));
+    cb(out_embd, "gen_audio_out", -1);
+    ggml_build_forward_expand(gf, out_embd);
+
+    return gf;
+}
+
+// frames -> PCM, qwen3tts code2wav blocks without pre_conv and tfm projections
+ggml_cgraph * clip_graph_drama3_gen::build_codec() {
+    const auto & c2w = model.c2w;
+    const int64_t n_cb = model.gen_code_embd_w->ne[2];
+
+    ggml_tensor * inp_codes = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, n_frames, n_cb);
+    ggml_set_name(inp_codes, "inp_codes");
+    ggml_set_input(inp_codes);
+
+    clip_graph_qwen3tts_gen::code2wav dec(*this);
+    dec.gelu_erf = true;
+    for (const auto & slot : list_c2w_state_slots(hparams, model)) {
+        ggml_tensor * t = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, slot.ne0, slot.ne1);
+        ggml_set_name(t, ("state_in_" + slot.name).c_str());
+        ggml_set_input(t);
+        dec.state_in[slot.name] = t;
+    }
+
+    // out_proj is folded into the codebook tables at conversion
+    auto group_ids = [&](int g) {
+        return ggml_view_1d(ctx0, inp_codes, n_frames, (size_t) g * inp_codes->nb[1]);
+    };
+    ggml_tensor * cur = ggml_get_rows(ctx0, c2w.quant_first_cb_w, group_ids(0)); // [dim, T]
+    for (int g = 1; g < n_cb; g++) {
+        ggml_tensor * table = ggml_view_2d(ctx0, c2w.quant_rest_cb_w, c2w.quant_rest_cb_w->ne[0], c2w.quant_rest_cb_w->ne[1],
+                                           c2w.quant_rest_cb_w->nb[1], (size_t) (g - 1) * c2w.quant_rest_cb_w->nb[2]);
+        cur = ggml_add(ctx0, cur, ggml_get_rows(ctx0, table, group_ids(g)));
+    }
+
+    for (int il = 0; il < hparams.wav_tfm_n_layer; il++) {
+        cur = dec.tfm_layer_forward(cur, c2w.tfm_layers[il], il);
+    }
+    cur = ggml_mul(ctx0, ggml_rms_norm(ctx0, cur, hparams.wav_tfm_eps), c2w.tfm_output_norm_w);
+    cb(cur, "wav_tfm_out", -1);
+
+    ggml_tensor * x = ggml_cont(ctx0, ggml_transpose(ctx0, cur)); // [T, dim]
+    for (size_t il = 0; il < c2w.upsample.size(); il++) {
+        const auto & up = c2w.upsample[il];
+        x = dec.causal_conv_transpose1d(x, up.conv_w, up.conv_b, 2, "");
+        x = dec.convnext_block(x, up, "up" + std::to_string(il));
+    }
+
+    static constexpr int DAC_DILATIONS[3] = { 1, 3, 9 };
+    x = dec.causal_conv1d(x, c2w.dac_entry_w, c2w.dac_entry_b, 1, "dac_entry");
+    for (size_t il = 0; il < c2w.dac.size(); il++) {
+        const auto & blk = c2w.dac[il];
+        const int stride = (int) (blk.conv_w->ne[0] / 2);
+        const std::string blk_name = "dac" + std::to_string(il);
+        x = dec.snake(x, blk.snake_alpha, blk.snake_beta);
+        x = dec.causal_conv_transpose1d(x, blk.conv_w, blk.conv_b, stride, blk_name + "_tail");
+        for (size_t ir = 0; ir < blk.res.size(); ir++) {
+            x = dec.dac_res_unit(x, blk.res[ir], DAC_DILATIONS[ir], blk_name + "_res" + std::to_string(ir));
+        }
+    }
+    x = dec.snake(x, c2w.dac_post_snake_alpha, c2w.dac_post_snake_beta);
+    x = dec.causal_conv1d(x, c2w.dac_post_conv_w, c2w.dac_post_conv_b, 1, "dac_post_conv");
+    x = ggml_tanh(ctx0, ggml_reshape_1d(ctx0, x, x->ne[0]));
+    ggml_set_name(x, "out_audio");
+    ggml_set_output(x);
+    ggml_build_forward_expand(gf, x);
+
+    for (auto & slot : dec.state_out) {
+        ggml_set_name(slot.second, ("state_out_" + slot.first).c_str());
+        ggml_set_output(slot.second);
+        ggml_build_forward_expand(gf, slot.second);
+    }
+    return gf;
+}
+
+ggml_cgraph * clip_graph_drama3_gen::build() {
+    return gen_process == CLIP_GEN_PROCESS_GEN_WAV ? build_codec() : build_tower();
+}
@@ -424,6 +424,9 @@ struct clip_graph_qwen3tts_gen : clip_graph {
        ggml_tensor * convnext_block(ggml_tensor * x, const clip_code2wav::upsample_block & blk, const std::string & state_prefix) const;
        ggml_tensor * dac_res_unit(ggml_tensor * x, const clip_code2wav::dac_res & res, int dilation, const std::string & state_name) const;

+        // drama3: erf GELU in ConvNeXt, not the tanh approximation
+        bool gelu_erf = false;
+
        // inp_codes [1, n_codes] I32 -> this frame's audio samples [n_samples] F32, clamped to [-1, 1]
        ggml_tensor * decode(ggml_tensor * inp_codes) const;
    };
@@ -487,12 +490,43 @@ struct c2w_state_slot {
    std::string name;
    int64_t     ne0;
    int64_t     ne1;
+    bool        in_place = false; // graph writes the slot directly, no state_in/state_out copy
 };
 std::vector<c2w_state_slot> list_c2w_state_slots(const clip_hparams & hparams, const clip_model & model);

+// drama3: tower KV slots, then list_c2w_state_slots()
+std::vector<c2w_state_slot> list_drama3_state_slots(const clip_hparams & hparams, const clip_model & model);
+
 // same, for the streaming mimi decoder (pocket-tts GEN_WAV)
 std::vector<c2w_state_slot> list_pockettts_state_slots(const clip_hparams & hparams, const clip_model & model);

+//
+// drama3 generation:
+// GEN_CODE = one codebook tower row, each layer attends over the whole sequence (KV in gen state)
+// GEN_WAV  = Fish DAC decoder, frames -> PCM
+//
+struct clip_graph_drama3_gen : clip_graph {
+    clip_graph_drama3_gen(clip_ctx * ctx, const clip_image_f32 & img, clip_gen_process_type gen_process, clip_gen_state * gen_state,
+                          int n_frames, int top_k, float top_p, float temp)
+        : clip_graph(ctx, img), gen_process(gen_process), gen_state(gen_state), n_frames(n_frames), top_k(top_k), top_p(top_p), temp(temp) {}
+    ggml_cgraph * build() override;
+
+    // tower KV rows, shared by the prompt and the generated frames
+    static constexpr int n_pos_max = 4096;
+
+    clip_gen_process_type gen_process;
+    clip_gen_state * gen_state; // nullptr on buffer reserve
+    int   n_frames;             // GEN_WAV only
+    int   top_k;
+    float top_p;
+    float temp;
+
+    ggml_cgraph * build_tower();
+    ggml_cgraph * build_codec();
+    ggml_tensor * tower_layer(ggml_tensor * x, int il, ggml_tensor * code, ggml_tensor * mask, ggml_tensor * extra,
+                              ggml_tensor * inp_pos, ggml_tensor * k_slot, ggml_tensor * v_slot, int pos);
+};
+
 struct clip_graph_soprano : clip_graph {
    clip_graph_soprano(clip_ctx * ctx, const clip_image_f32 & img, int n_frames) : clip_graph(ctx, img), n_frames(n_frames) {}
    ggml_cgraph * build() override;
@@ -527,7 +527,7 @@ ggml_tensor * clip_graph_qwen3tts_gen::code2wav::convnext_block(ggml_tensor * x,

    ggml_tensor * g = ggml_mul_mat(ctx0, blk.pw1_w, hc);
    g = ggml_add(ctx0, g, blk.pw1_b);
-    g = ggml_gelu(ctx0, g);
+    g = gelu_erf ? ggml_gelu_erf(ctx0, g) : ggml_gelu(ctx0, g);
    g = ggml_mul_mat(ctx0, blk.pw2_w, g);
    g = ggml_add(ctx0, g, blk.pw2_b);
    g = ggml_mul(ctx0, g, blk.gamma);
@@ -617,7 +617,7 @@ std::vector<c2w_state_slot> list_c2w_state_slots(const clip_hparams & hparams, c
    const auto & c2w = model.c2w;
    std::vector<c2w_state_slot> slots;

-    if (c2w.pre_conv_w == nullptr) {
+    if (c2w.dac_entry_w == nullptr) {
        return slots; // not a code2wav model, it keeps no state between calls
    }

@@ -632,7 +632,9 @@ std::vector<c2w_state_slot> list_c2w_state_slots(const clip_hparams & hparams, c
        slots.push_back({"tfm_v_" + std::to_string(il), kv_ch, prefix});
    }

-    slots.push_back({"pre_conv", c2w.pre_conv_w->ne[0] - 1, c2w.pre_conv_w->ne[1]});
+    if (c2w.pre_conv_w) {
+        slots.push_back({"pre_conv", c2w.pre_conv_w->ne[0] - 1, c2w.pre_conv_w->ne[1]});
+    }

    for (size_t il = 0; il < c2w.upsample.size(); il++) {
        const auto & up = c2w.upsample[il];
@@ -2077,6 +2077,345 @@ private:
    bool   wav_header_sent = false;
 };

+// Drama 3: backbone samples the main token of each frame, codebook tower samples the other codes, Fish DAC makes the PCM
+// the tower attends over the whole sequence, so it also runs on every prompt row
+class drama3_gen_audio_pipeline : public mtmd_gen_audio_pipeline {
+public:
+    using mtmd_gen_audio_pipeline::mtmd_gen_audio_pipeline;
+
+    void reset() override {
+        seq_id = 0;
+        pos = 0;
+        prompt_tokens.clear();
+        prompt_pos = 0;
+        frame.clear();
+        hist.clear();
+        codes_buf.clear();
+        n_decoded = 0;
+        gen_state.reset();
+        pcm_sent = 0;
+        wav_header_sent = false;
+        audio_pcm.clear();
+        h_state_buf.clear();
+        out_buf.clear();
+    }
+
+    int32_t set_input(const mtmd_helper_gen_audio_inp * inp) override {
+        reset();
+        seq_id   = inp->seq_id;
+        out_type = inp->out_type;
+        stream   = inp->stream;
+        if (inp->speaker_ref) {
+            LOG_WRN("mtmd_helper_gen_audio: voice cloning is not supported yet, the speaker file is ignored\n");
+        }
+
+        sem_begin = find_special_token(vocab, "<|semantic:0|>");
+        eos       = find_special_token(vocab, "<|im_end|>");
+        if (sem_begin == LLAMA_TOKEN_NULL || eos == LLAMA_TOKEN_NULL) {
+            LOG_ERR("mtmd_helper_gen_audio: missing Drama 3 tokens in vocab\n");
+            return 1;
+        }
+
+        // reference-bracket-v1 prompt without references, tokenize each segment alone as the reference does
+        std::string text(inp->prompt, inp->prompt_len);
+        if (text.rfind("<|speaker:", 0) != 0) {
+            text = "<|speaker:0|> " + text;
+        }
+        const char * segments[] = {
+            "<|im_start|>system\n", "convert the provided text to speech", "<|im_end|>\n",
+            "<|im_start|>user\n", text.c_str(), "<|im_end|>\n", "<|im_start|>assistant\n<|voice|>",
+        };
+        for (const char * seg : segments) {
+            const int32_t len = (int32_t) strlen(seg);
+            std::vector<llama_token> ids(len + 16);
+            const int32_t n = llama_tokenize(vocab, seg, len, ids.data(), (int32_t) ids.size(), false, true);
+            if (n < 0) {
+                LOG_ERR("mtmd_helper_gen_audio: tokenization failed\n");
+                return 1;
+            }
+            prompt_tokens.insert(prompt_tokens.end(), ids.begin(), ids.begin() + n);
+        }
+        if (prompt_tokens.size() >= llama_n_ctx_seq(lctx) || prompt_tokens.size() > n_tower_rows) {
+            LOG_ERR("mtmd_helper_gen_audio: prompt does not fit in the context\n");
+            return 1;
+        }
+        if (!llama_memory_seq_rm(llama_get_memory(lctx), seq_id, 0, -1)) {
+            LOG_ERR("mtmd_helper_gen_audio: cannot reset the sequence\n");
+            return 1;
+        }
+        rng.seed(inp->seed == UINT32_MAX ? std::random_device{}() : inp->seed);
+        return init_gen_state(inp->seed) ? 0 : 1;
+    }
+
+    // run the tower on every prompt row to fill its KV
+    int32_t step_prompt(int32_t n_batch) override {
+        GGML_ASSERT(n_batch > 0);
+        const int n_prompt = (int) prompt_tokens.size();
+        if (prompt_pos >= n_prompt) {
+            return 0;
+        }
+        // small batches: every row also outputs logits
+        const int n = std::min({ n_batch, n_prompt - prompt_pos, 64 });
+        llama_batch batch = llama_batch_init(n, 0, 1);
+        for (int i = 0; i < n; i++) {
+            batch.token   [i]    = prompt_tokens[(size_t) (prompt_pos + i)];
+            batch.pos     [i]    = pos + i;
+            batch.n_seq_id[i]    = 1;
+            batch.seq_id  [i][0] = seq_id;
+            batch.logits  [i]    = 1;
+        }
+        batch.n_tokens = n;
+        const int ret = llama_decode(lctx, batch);
+        llama_batch_free(batch);
+        if (ret != 0) {
+            LOG_ERR("mtmd_helper_gen_audio: prompt decode failed\n");
+            return -1;
+        }
+        for (int i = 0; i < n; i++) {
+            const float * h = llama_get_embeddings_ith(lctx, i);
+            if (prompt_pos + i == n_prompt - 1) {
+                h_state_buf.assign(h, h + n_embd); // first frame comes from the last prompt row
+                break;
+            }
+            if (!run_tower(-1, h)) {
+                return -1;
+            }
+        }
+        pos        += n;
+        prompt_pos += n;
+        return n_prompt - prompt_pos;
+    }
+
+    // sample the main token here, ignore the token of the caller
+    int32_t step_gen(llama_token sampled, const float * h_state_in, const float ** h_state_out, bool * out_stop) override {
+        GGML_UNUSED(sampled);
+        GGML_UNUSED(h_state_in);
+        *h_state_out = nullptr;
+        *out_stop    = false;
+
+        const llama_token main = sample_main(llama_get_logits_ith(lctx, -1));
+        const size_t n_rows = prompt_tokens.size() - 1 + codes_buf.size() / n_cb;
+        if (main == eos || (uint32_t) pos >= llama_n_ctx_seq(lctx) || n_rows >= n_tower_rows) {
+            *out_stop = true;
+            return 0;
+        }
+        hist.push_back(main);
+        if (hist.size() > ras_window) {
+            hist.erase(hist.begin());
+        }
+
+        const float * next_embd = run_tower(main - sem_begin, h_state_buf.data());
+        if (!next_embd) {
+            return 1;
+        }
+        codes_buf.insert(codes_buf.end(), frame.begin(), frame.end());
+
+        // text/speech stream: all mrope sections use the same position
+        const bool mrope = llama_model_rope_type(model) == LLAMA_ROPE_TYPE_MROPE ||
+                           llama_model_rope_type(model) == LLAMA_ROPE_TYPE_IMROPE;
+        decode_embd_batch batch(next_embd, 1, mrope ? 4 : 1, n_embd);
+        if (mrope) {
+            batch.set_position_mrope_1d(pos, seq_id);
+        } else {
+            batch.set_position_normal(pos, seq_id);
+        }
+        batch.logits[0] = 1;
+        if (llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, batch.render(lctx, 0, 1)) != 0) {
+            LOG_ERR("mtmd_helper_gen_audio: decode failed\n");
+            return 1;
+        }
+        pos++;
+        const float * h = llama_get_embeddings_ith(lctx, -1);
+        h_state_buf.assign(h, h + n_embd);
+        *h_state_out = h_state_buf.data();
+
+        if (stream && codes_buf.size() / n_cb - n_decoded >= stream_frames) {
+            return decode_wav() ? 0 : 1;
+        }
+        return 0;
+    }
+
+    int32_t get_output(int32_t * out_sample_rate, const char ** out_data, size_t * out_data_len, int64_t * out_n_samples) override {
+        *out_sample_rate = info.sample_rate;
+
+        if (!stream) {
+            if (flush() != 0) {
+                return 1;
+            }
+            if (out_n_samples) {
+                *out_n_samples = (int64_t) audio_pcm.size();
+            }
+            if (out_type == MTMD_HELPER_GEN_AUDIO_OUTTYPE_PCM) {
+                *out_data     = (const char *) audio_pcm.data();
+                *out_data_len = audio_pcm.size() * sizeof(float);
+                return 0;
+            }
+            out_buf.clear();
+            if (!write_wav16(out_buf, audio_pcm, info.sample_rate)) {
+                LOG_ERR("mtmd_helper_gen_audio: output too large for WAV\n");
+                return 1;
+            }
+            *out_data     = out_buf.data();
+            *out_data_len = out_buf.size();
+            return 0;
+        }
+
+        // streaming: only the audio since the previous call
+        const size_t n_new = audio_pcm.size() - pcm_sent;
+        if (out_n_samples) {
+            *out_n_samples = (int64_t) n_new;
+        }
+        if (out_type == MTMD_HELPER_GEN_AUDIO_OUTTYPE_PCM) {
+            *out_data     = (const char *) (audio_pcm.data() + pcm_sent);
+            *out_data_len = n_new * sizeof(float);
+            pcm_sent = audio_pcm.size();
+            return 0;
+        }
+        out_buf.clear();
+        if (!wav_header_sent) {
+            write_wav16_header(out_buf, UINT32_MAX, info.sample_rate);
+            wav_header_sent = true;
+        }
+        append_wav16_pcm(out_buf, audio_pcm.data() + pcm_sent, n_new);
+        pcm_sent = audio_pcm.size();
+        *out_data     = out_buf.data();
+        *out_data_len = out_buf.size();
+        return 0;
+    }
+
+    int32_t flush() override {
+        return decode_wav() ? 0 : 1;
+    }
+
+private:
+    // repetition-aware sampling: if a semantic draw repeats one of the last ras_window main tokens, draw again with fixed settings
+    static constexpr size_t ras_window = 16;
+    static constexpr size_t n_cb       = 10;
+    // codebook tower KV rows (clip_graph_drama3_gen::n_pos_max): one per prompt row but the last, one per frame
+    static constexpr size_t n_tower_rows = 4096;
+    // frames per codec call when streaming (0.33 s)
+    static constexpr size_t stream_frames = 7;
+
+    // top-p keeps the token that crosses the threshold, as the reference does
+    llama_token draw(const std::vector<float> & cand_logits, float temp, int32_t top_k, float top_p) {
+        std::vector<std::pair<float, int>> cand(cand_logits.size());
+        for (size_t i = 0; i < cand.size(); i++) {
+            cand[i] = { cand_logits[i] / temp, (int) i };
+        }
+        std::sort(cand.begin(), cand.end(), [](const auto & a, const auto & b) { return a.first > b.first; });
+        const float  lmax = cand[0].first;
+        double       sum  = 0.0;
+        std::vector<double> prob(cand.size());
+        for (size_t i = 0; i < cand.size(); i++) {
+            prob[i] = std::exp((double) (cand[i].first - lmax));
+            sum += prob[i];
+        }
+        size_t n_keep = cand.size();
+        if (top_k > 0) {
+            n_keep = std::min(n_keep, (size_t) top_k);
+        }
+        double cum = 0.0;
+        for (size_t i = 0; i < n_keep; i++) {
+            if (cum > top_p * sum) {
+                n_keep = i;
+                break;
+            }
+            cum += prob[i];
+        }
+        double kept = 0.0;
+        for (size_t i = 0; i < n_keep; i++) {
+            kept += prob[i];
+        }
+        double r = std::uniform_real_distribution<double>(0.0, 1.0)(rng) * kept;
+        for (size_t i = 0; i < n_keep; i++) {
+            r -= prob[i];
+            if (r <= 0.0) {
+                return cand[i].second;
+            }
+        }
+        return cand[n_keep - 1].second;
+    }
+
+    // candidates: eos, then semantic codes 1.. (the reference never samples code 0)
+    llama_token sample_main(const float * logits) {
+        std::vector<float> cand(n_cb_vocab);
+        cand[0] = logits[eos];
+        for (int i = 1; i < n_cb_vocab; i++) {
+            cand[i] = logits[sem_begin + i];
+        }
+        auto to_token = [&](int i) { return i == 0 ? eos : sem_begin + i; };
+        const llama_token first = to_token(draw(cand, 1.0f, 0, 0.9f));
+        if (first != eos && std::find(hist.begin(), hist.end(), first) != hist.end()) {
+            return to_token(draw(cand, 1.0f, 30, 0.95f));
+        }
+        return first;
+    }
+
+    // code0 < 0 for a prompt row, else the main code of the next frame
+    // returns the next backbone input, the codes of the next frame go to frame
+    const float * run_tower(int32_t code0, const float * h) {
+        mtmd_gen_inp inp = mtmd_gen_inp_default(mctx);
+        inp.type    = MTMD_GEN_PROCESS_TYPE_GEN_CODE;
+        inp.code0   = code0;
+        inp.embd    = const_cast<float *>(h);
+        inp.codes   = frame.empty() ? nullptr : frame.data();
+        inp.n_codes = frame.size();
+        mtmd_gen_out out{};
+        if (mtmd_gen_audio_process(mctx, gen_state.get(), &inp, &out) != 0 || (code0 >= 0 && out.n_codes != n_cb)) {
+            LOG_ERR("mtmd_helper_gen_audio: codebook tower failed\n");
+            return nullptr;
+        }
+        if (code0 >= 0) {
+            frame.assign(out.codes, out.codes + out.n_codes);
+        }
+        return out.embd;
+    }
+
+    // the codec is stateful, so frames can be decoded in pieces
+    bool decode_wav() {
+        const size_t n_frames = codes_buf.size() / n_cb;
+        while (n_decoded < n_frames) {
+            const size_t n = std::min(n_frames - n_decoded, (size_t) 64);
+            mtmd_gen_inp inp = mtmd_gen_inp_default(mctx);
+            inp.type    = MTMD_GEN_PROCESS_TYPE_GEN_WAV;
+            inp.codes   = codes_buf.data() + n_decoded * n_cb;
+            inp.n_codes = n * n_cb;
+            mtmd_gen_out out{};
+            if (mtmd_gen_audio_process(mctx, gen_state.get(), &inp, &out) != 0) {
+                LOG_ERR("mtmd_helper_gen_audio: codec decode failed\n");
+                return false;
+            }
+            audio_pcm.insert(audio_pcm.end(), out.audio, out.audio + out.n_samples);
+            n_decoded += n;
+        }
+        return true;
+    }
+
+    static constexpr int n_cb_vocab = 4096; // <|semantic:0..4095|>
+
+    llama_token sem_begin = LLAMA_TOKEN_NULL;
+    llama_token eos       = LLAMA_TOKEN_NULL;
+
+    llama_seq_id seq_id = 0;
+    llama_pos    pos    = 0;
+    std::vector<llama_token> prompt_tokens;
+    int prompt_pos = 0;
+
+    std::mt19937 rng;
+    std::vector<int32_t>     frame;     // codes of the last frame, input of the next tower row
+    std::vector<llama_token> hist;      // last main tokens, for repetition-aware sampling
+    std::vector<int32_t>     codes_buf; // frame-major
+    size_t n_decoded = 0;               // frames already through the codec
+
+    std::vector<float>   audio_pcm;
+    std::vector<float>   h_state_buf;
+    mtmd_helper_gen_audio_outtype out_type = MTMD_HELPER_GEN_AUDIO_OUTTYPE_WAV;
+    std::vector<char> out_buf;
+    bool   stream          = false;
+    size_t pcm_sent        = 0; // samples already returned by get_output()
+    bool   wav_header_sent = false;
+};
+
 static std::unique_ptr<mtmd_gen_audio_pipeline> make_pipeline(llama_context * lctx, mtmd_context * mctx) {
    switch (mtmd_gen_audio_get_info(mctx).type) {
        case MTMD_GEN_AUDIO_TYPE_QWEN3TTS:
@@ -2089,6 +2428,8 @@ static std::unique_ptr<mtmd_gen_audio_pipeline> make_pipeline(llama_context * lc
            return std::unique_ptr<mtmd_gen_audio_pipeline>(new lfm3a_gen_audio_pipeline(lctx, mctx));
        case MTMD_GEN_AUDIO_TYPE_CHATTERBOX:
            return std::unique_ptr<mtmd_gen_audio_pipeline>(new chatterbox_gen_audio_pipeline(lctx, mctx));
+        case MTMD_GEN_AUDIO_TYPE_DRAMA3:
+            return std::unique_ptr<mtmd_gen_audio_pipeline>(new drama3_gen_audio_pipeline(lctx, mctx));
        default:
            return nullptr;
    }
@@ -2373,6 +2373,10 @@ mtmd_gen_audio_info mtmd_gen_audio_get_info(const mtmd_context * ctx) {
            info.type = MTMD_GEN_AUDIO_TYPE_CHATTERBOX;
            info.sample_rate = 24000;
            break;
+        case PROJECTOR_TYPE_DRAMA3_GEN:
+            info.type = MTMD_GEN_AUDIO_TYPE_DRAMA3;
+            info.sample_rate = 44100;
+            break;
        default:
            info.type = MTMD_GEN_AUDIO_TYPE_NONE;
            break;
@@ -2776,6 +2780,12 @@ mtmd_gen_inp mtmd_gen_inp_default(const mtmd_context * ctx) {
        case PROJECTOR_TYPE_SOPRANO:
            inp.type = MTMD_GEN_PROCESS_TYPE_GEN_WAV;
            break;
+        case PROJECTOR_TYPE_DRAMA3_GEN:
+            // codebook heads, same as the reference (sglang-omni fishaudio_drama3)
+            inp.top_k = 30;
+            inp.top_p = 0.9f;
+            inp.temp  = 1.0f;
+            break;
        case PROJECTOR_TYPE_LFM3A_GEN:
            {
                // MaskGIT sampling runs in the helper, with the settings of the checkpoint
@@ -467,6 +467,7 @@ enum mtmd_gen_audio_type {
    MTMD_GEN_AUDIO_TYPE_SOPRANO,
    MTMD_GEN_AUDIO_TYPE_LFM3A,
    MTMD_GEN_AUDIO_TYPE_CHATTERBOX,
+    MTMD_GEN_AUDIO_TYPE_DRAMA3,
 };

 struct mtmd_gen_audio_info {
@@ -128,13 +128,10 @@ static std::string path_to_utf8(const fs::path & p) {
 static const std::string & home_dir() {
    static const std::string home = [] {
 #ifdef _WIN32
-        // the narrow getenv would return the profile path in the active code page
-        const wchar_t * w = _wgetenv(L"HOME");
-        if (w == nullptr) w = _wgetenv(L"USERPROFILE");
-        return w ? path_to_utf8(fs::path(w)) : std::string();
+        const std::string h = common_get_env("HOME");
+        return h.empty() ? common_get_env("USERPROFILE") : h;
 #else
-        const char * h = getenv("HOME");
-        return h ? std::string(h) : std::string();
+        return common_get_env("HOME");
 #endif
    }();
    return home;
@@ -154,3 +154,34 @@ The upstream repositories ship raw checkpoints without `config.json`. Download t
 python convert_hf_to_gguf.py path/to/chatterbox --outfile chatterbox.gguf
 python convert_hf_to_gguf.py path/to/chatterbox --mmproj --outfile mmproj-chatterbox.gguf
 ```
+
+## Drama 3
+
+[Fish Audio Drama 3](https://huggingface.co/fishaudio/drama-3): a Qwen3.5-MoE backbone (35B, 3B active) samples each frame's main token, a codebook tower samples the 9 other codes, and the Fish DAC decodes them to 44.1 kHz audio.
+
+Available params:
+- Speaker tags `<|speaker:0|>` to `<|speaker:15|>` in the text give each character its own voice, a text without a tag starts with `<|speaker:0|>`
+- The language is the one of the text, `--tts-lang` is not used
+- Note: `--tts-speaker-file` is not supported yet (voice cloning)
+- Note: `--top-k` and `--top-p` are not used, frames are sampled with the settings of the reference (main token: temperature 1.0, top-p 0.9, repetition-aware redraw over the last 16 frames; other codes: temperature 1.0, top-k 30, top-p 0.9)
+- Each frame is 2048 samples (about 46 ms), generation stops at the end-of-speech token
+- The codebook tower keeps the whole sequence in its KV cache, the prompt and the generated frames are limited to 4096 positions
+
+Example usage:
+
+```sh
+llama-tts -m drama3.gguf \
+    -mm mmproj-drama3.gguf \
+    -p "<|speaker:0|> Did you remember to buy the bread? <|speaker:1|> Of course I did, it is on the kitchen table." \
+    -n 1000 \
+    --output out.wav
+```
+
+**Note for GGUF conversion:**
+
+The routed experts of the checkpoint are FP8, they are dequantized at conversion. The mmproj holds the codebook tower and the decoder half of `codec.pth`, it is best kept in BF16:
+
+```sh
+python convert_hf_to_gguf.py path/to/drama-3 --outtype q8_0 --outfile drama3.gguf
+python convert_hf_to_gguf.py path/to/drama-3 --mmproj --outtype bf16 --outfile mmproj-drama3.gguf
+```
@@ -106,18 +106,22 @@
 		{readonly}
 		width={imageWidth}
 	/>
-{:else if isPdfFile(item.attachment, item.uploadedFile)}
-	<ChatAttachmentsListItemThumbnailFile
-		attachment={item.attachment}
-		class="flex-shrink-0 cursor-pointer {className} {scrollClasses}"
+{:else if isPdfFile(item.attachment, item.uploadedFile) && item.preview}
+	{@const pageCount =
+		item.attachment?.type === AttachmentType.PDF ? (item.attachment.images?.length ?? 0) : 0}
+
+	<ChatAttachmentsListItemThumbnailImage
+		class="{layoutClasses} cursor-pointer {className} {scrollClasses}"
+		height={imageHeight}
 		id={item.id}
+		{imageClass}
+		label={pageCount > 0 ? `PDF, ${pageCount} ${pageCount === 1 ? 'page' : 'pages'}` : 'PDF'}
 		name={item.name}
 		onRemove={onFileRemove}
 		onclick={() => onPreview?.(item)}
+		preview={item.preview}
 		{readonly}
-		size={item.size}
-		textContent={item.textContent}
-		uploadedFile={item.uploadedFile}
+		width={imageWidth}
 	/>
 {:else}
 	<ChatAttachmentsListItemThumbnailFile
@@ -7,6 +7,7 @@
 		height?: string;
 		id: string;
 		imageClass?: string;
+		label?: string;
 		onclick?: (event?: MouseEvent) => void;
 		onRemove?: (id: string) => void;
 		name: string;
@@ -20,6 +21,7 @@
 		height = 'h-16',
 		id,
 		imageClass = '',
+		label,
 		name,
 		onclick,
 		onRemove,
@@ -49,6 +51,14 @@
 		{@render image()}
 	{/if}

+	{#if label}
+		<span
+			class="pointer-events-none absolute bottom-1 left-1 rounded bg-black/60 px-1.5 py-0.5 text-[10px] font-medium text-white"
+		>
+			{label}
+		</span>
+	{/if}
+
 	{#if !readonly}
 		<div
 			class="absolute top-1 right-1 flex items-center justify-center opacity-0 transition-opacity group-focus-within:opacity-100 group-hover:opacity-100"
@@ -32,6 +32,7 @@
 		isImage: boolean;
 		isAudio: boolean;
 		isVideo: boolean;
+		isPdf: boolean;
 	}

 	interface Props {
@@ -58,6 +59,7 @@
 					...item,
 					isAudio: isAudioFile(item.attachment, item.uploadedFile),
 					isImage: isImageFile(item.attachment, item.uploadedFile),
+					isPdf: isPdfFile(item.attachment, item.uploadedFile),
 					isVideo: isVideoFile(item.attachment, item.uploadedFile)
 				})
 			)
@@ -3,11 +3,11 @@
 	import { SyntaxHighlightedCode } from '$lib/components/app';
 	import * as Alert from '$lib/components/ui/alert';
 	import { Button } from '$lib/components/ui/button';
-	import { ICON_CLASS_DEFAULT } from '$lib/constants';
-	import { PdfViewMode } from '$lib/enums';
+	import { ICON_CLASS_DEFAULT, PDF_AS_IMAGES_LEGACY_CONTENT_REGEX } from '$lib/constants';
+	import { AttachmentType, MimeTypeApplication, PdfViewMode } from '$lib/enums';
 	import type { ChatAttachmentDisplayItem } from '$lib/types';
 	import { getLanguageFromFilename } from '$lib/utils';
-	import { convertPDFToImage } from '$lib/utils/browser-only';
+	import { convertPDFToImage, convertPDFToText } from '$lib/utils/browser-only';

 	interface Props {
 		currentItem: ChatAttachmentDisplayItem | null;
@@ -23,54 +23,70 @@
 	let pdfViewMode = $state<PdfViewMode>(PdfViewMode.PAGES);
 	let pdfImages = $state<string[]>([]);
 	let pdfImagesLoading = $state(false);
+	// a failure stays on screen: the effect does not retry while an error is set
 	let pdfImagesError = $state<string | null>(null);
+	let pdfText = $state<string | null>(null);
+	let pdfTextLoading = $state(false);
+	let pdfTextError = $state<string | null>(null);

 	let language = $derived(getLanguageFromFilename(displayName));

+	// older PDFs sent as images stored a placeholder in `content`, so their text is extracted again
+	let hasStoredText = $derived(
+		currentItem?.attachment?.type === AttachmentType.PDF
+			? !(
+					currentItem.attachment.processedAsImages &&
+					PDF_AS_IMAGES_LEGACY_CONTENT_REGEX.test(currentItem.attachment.content)
+				)
+			: !!displayTextContent
+	);
+
+	let isLoading = $derived(pdfImagesLoading || pdfTextLoading);
+
+	function getPdfFile(): File | null {
+		if (currentItem?.uploadedFile?.file) {
+			return currentItem.uploadedFile.file;
+		}
+
+		if (
+			currentItem?.attachment &&
+			'base64Data' in currentItem.attachment &&
+			currentItem.attachment.base64Data
+		) {
+			const byteArray = Uint8Array.from(atob(currentItem.attachment.base64Data), (char) =>
+				char.charCodeAt(0)
+			);
+
+			return new File([byteArray], displayName, { type: MimeTypeApplication.PDF });
+		}
+
+		return null;
+	}
+
 	async function loadPdfImages() {
 		if (pdfImages.length > 0 || pdfImagesLoading || !currentItem) return;

+		if (
+			currentItem.attachment?.type === AttachmentType.PDF &&
+			currentItem.attachment.images?.length
+		) {
+			pdfImages = currentItem.attachment.images;
+
+			return;
+		}
+
+		if (pdfImagesError) return;
+
 		pdfImagesLoading = true;
-		pdfImagesError = null;

 		try {
-			let file: File | null = null;
+			const file = getPdfFile();

-			if (currentItem.uploadedFile?.file) {
-				file = currentItem.uploadedFile.file;
-			} else if (currentItem.attachment) {
-				// Check if we have pre-processed images
-				if (
-					'images' in currentItem.attachment &&
-					currentItem.attachment.images &&
-					Array.isArray(currentItem.attachment.images) &&
-					currentItem.attachment.images.length > 0
-				) {
-					pdfImages = currentItem.attachment.images;
-
-					return;
-				}
-
-				// Convert base64 back to File for processing
-				if ('base64Data' in currentItem.attachment && currentItem.attachment.base64Data) {
-					const base64Data = currentItem.attachment.base64Data;
-					const byteCharacters = atob(base64Data);
-					const byteNumbers = new Array(byteCharacters.length);
-
-					for (let i = 0; i < byteCharacters.length; i++) {
-						byteNumbers[i] = byteCharacters.charCodeAt(i);
-					}
-					const byteArray = new Uint8Array(byteNumbers);
-
-					file = new File([byteArray], displayName, { type: 'application/pdf' });
-				}
-			}
-
-			if (file) {
-				pdfImages = await convertPDFToImage(file);
-			} else {
+			if (!file) {
 				throw new Error('No PDF file available for conversion');
 			}
+
+			pdfImages = await convertPDFToImage(file);
 		} catch (error) {
 			pdfImagesError = error instanceof Error ? error.message : 'Failed to load PDF images';
 		} finally {
@@ -78,26 +94,62 @@
 		}
 	}

+	async function loadPdfText() {
+		if (pdfText !== null || pdfTextLoading || !currentItem) return;
+
+		if (hasStoredText) {
+			pdfText = displayTextContent ?? '';
+
+			return;
+		}
+
+		if (pdfTextError) return;
+
+		pdfTextLoading = true;
+
+		try {
+			const file = getPdfFile();
+
+			if (!file) {
+				throw new Error('No PDF file available for text extraction');
+			}
+
+			pdfText = await convertPDFToText(file);
+		} catch (error) {
+			pdfTextError = error instanceof Error ? error.message : 'Failed to extract PDF text';
+		} finally {
+			pdfTextLoading = false;
+		}
+	}
+
 	$effect(() => {
 		if (pdfViewMode === PdfViewMode.PAGES) {
 			loadPdfImages();
+		} else {
+			loadPdfText();
 		}
 	});
 </script>

 <div class="mb-4 flex items-center justify-end gap-2">
 	<Button
-		disabled={pdfImagesLoading}
+		disabled={isLoading}
 		onclick={() => (pdfViewMode = PdfViewMode.TEXT)}
 		size="sm"
 		variant={pdfViewMode === PdfViewMode.TEXT ? 'default' : 'outline'}
 	>
-		<FileText class="mr-1 {ICON_CLASS_DEFAULT}" />
+		{#if pdfTextLoading}
+			<div
+				class="mr-1 {ICON_CLASS_DEFAULT} animate-spin rounded-full border-2 border-current border-t-transparent"
+			></div>
+		{:else}
+			<FileText class="mr-1 {ICON_CLASS_DEFAULT}" />
+		{/if}
 		Text
 	</Button>

 	<Button
-		disabled={pdfImagesLoading}
+		disabled={isLoading}
 		onclick={() => (pdfViewMode = PdfViewMode.PAGES)}
 		size="sm"
 		variant={pdfViewMode === PdfViewMode.PAGES ? 'default' : 'outline'}
@@ -136,28 +188,24 @@
 	</Alert.Root>
 {/if}

-{#if pdfImagesLoading}
-	<div class="flex flex-1 items-center justify-center p-8">
-		<div class="text-center">
-			<div
-				class="mx-auto mb-4 h-8 w-8 animate-spin rounded-full border-4 border-white border-t-transparent"
-			></div>
-
-			<p class="text-white/70">Converting PDF to images...</p>
+{#if pdfViewMode === PdfViewMode.TEXT}
+	{#if pdfTextLoading}
+		{@render loadingStatus('Extracting PDF text...')}
+	{:else if pdfTextError}
+		{@render status('Failed to extract PDF text', pdfTextError)}
+	{:else if pdfText}
+		<div class="px-4 pb-4">
+			<SyntaxHighlightedCode class="max-w-4xl" code={pdfText} {language} maxHeight="none" />
 		</div>
-	</div>
+	{:else}
+		{@render status('No text found in this PDF')}
+	{/if}
+{:else if pdfImagesLoading}
+	{@render loadingStatus('Converting PDF to images...')}
 {:else if pdfImagesError}
-	<div class="flex flex-1 items-center justify-center p-8">
-		<div class="text-center">
-			<FileText class="mx-auto mb-4 h-16 w-16 text-white/50" />
-
-			<p class="mb-4 text-white/70">Failed to load PDF images</p>
-
-			<p class="text-sm text-white/50">{pdfImagesError}</p>
-		</div>
-	</div>
+	{@render status('Failed to load PDF images', pdfImagesError)}
 {:else if pdfImages.length > 0}
-	{#each pdfImages as image, index (image)}
+	{#each pdfImages as image, index (index)}
 		<p class="mb-2 text-sm text-white/50">Page {index + 1}</p>

 		<img alt="PDF Page {index + 1}" class="mx-auto max-w-[85vw] rounded-lg shadow-lg" src={image} />
@@ -165,22 +213,31 @@
 		<div class="h-4"></div>
 	{/each}
 {:else}
+	{@render status('No PDF pages available')}
+{/if}
+
+{#snippet loadingStatus(message: string)}
+	<div class="flex flex-1 items-center justify-center p-8">
+		<div class="text-center">
+			<div
+				class="mx-auto mb-4 h-8 w-8 animate-spin rounded-full border-4 border-white border-t-transparent"
+			></div>
+
+			<p class="text-white/70">{message}</p>
+		</div>
+	</div>
+{/snippet}
+
+{#snippet status(message: string, detail?: string)}
 	<div class="flex flex-1 items-center justify-center p-8">
 		<div class="text-center">
 			<FileText class="mx-auto mb-4 h-16 w-16 text-white/50" />

-			<p class="text-white/70">No PDF pages available</p>
+			<p class={['text-white/70', detail && 'mb-4']}>{message}</p>
+
+			{#if detail}
+				<p class="text-sm text-white/50">{detail}</p>
+			{/if}
 		</div>
 	</div>
-{/if}
-
-{#if pdfViewMode === PdfViewMode.TEXT && displayTextContent}
-	<div class="px-4 pb-4">
-		<SyntaxHighlightedCode
-			class="max-w-4xl"
-			code={displayTextContent}
-			{language}
-			maxHeight="none"
-		/>
-	</div>
-{/if}
+{/snippet}
@@ -10,6 +10,7 @@
 		isImage: boolean;
 		isAudio: boolean;
 		isVideo: boolean;
+		isPdf: boolean;
 		preview?: string;
 	}

@@ -46,7 +47,7 @@
 					]}
 					onclick={() => onNavigate(index)}
 				>
-					{#if item.isImage && item.preview}
+					{#if (item.isImage || item.isPdf) && item.preview}
 						<img alt={item.name} class="h-12 w-12 object-cover" src={item.preview} />
 					{:else}
 						<div
@@ -74,6 +74,9 @@ export const IMAGE_FILE_TYPES = {
 	}
 } as const;

+/** `content` that messages stored before the text was kept carry for a PDF sent as images. */
+export const PDF_AS_IMAGES_LEGACY_CONTENT_REGEX = /^PDF file with \d+ pages$/;
+
 export const PDF_FILE_TYPES = {
 	[FileTypePdf.PDF]: {
 		extensions: [FileExtensionPdf.PDF],
@@ -71,13 +71,22 @@ export function getAttachmentDisplayItems(
 	for (const [index, attachment] of attachments.entries()) {
 		const isImage = isImageFile(attachment);

+		let preview: string | undefined;
+
+		if (isImage && 'base64Url' in attachment) {
+			preview = attachment.base64Url;
+		} else if (attachment.type === AttachmentType.PDF && attachment.processedAsImages) {
+			// first rendered page is the thumbnail of a PDF sent as images
+			preview = attachment.images?.[0];
+		}
+
 		items.push({
 			attachment,
 			attachmentIndex: index,
 			id: `attachment-${index}`,
 			isImage,
 			name: attachment.name,
-			preview: isImage && 'base64Url' in attachment ? attachment.base64Url : undefined,
+			preview,
 			size: 'size' in attachment ? attachment.size : undefined,
 			textContent: 'content' in attachment ? attachment.content : undefined
 		});
@@ -140,6 +140,8 @@ export async function parseFilesToMessageExtras(
 					// Process PDF as images (only for vision models)
 					try {
 						const images = await convertPDFToImage(file.file);
+						// kept for the preview's Text view; the model only receives the images
+						const content = file.textContent ?? (await convertPDFToText(file.file).catch(() => ''));

 						// Show success toast for PDF image processing
 						toast.success(
@@ -151,7 +153,7 @@ export async function parseFilesToMessageExtras(

 						extras.push({
 							base64Data: base64Data,
-							content: `PDF file with ${images.length} pages`,
+							content,
 							images: images,
 							name: file.name,
 							processedAsImages: true,
@@ -0,0 +1,47 @@
+<script lang="ts" module>
+	import { defineMeta } from '@storybook/addon-svelte-csf';
+	import ChatAttachmentsPreview from '$lib/components/app/chat/ChatAttachments/ChatAttachmentsPreview/ChatAttachmentsPreview.svelte';
+	import { AttachmentType } from '$lib/enums';
+	import { expect, waitFor } from 'storybook/test';
+
+	const { Story } = defineMeta({
+		component: ChatAttachmentsPreview,
+		parameters: {
+			layout: 'fullscreen'
+		},
+		title: 'Components/ChatScreen/ChatAttachments/ChatAttachmentsPreview'
+	});
+
+	// same backdrop as DialogChatAttachmentsPreview
+	const frameClass = 'flex h-screen flex-col bg-neutral-950';
+
+	// valid base64 that is not a PDF: pdf.js throws InvalidPDFException
+	const brokenPdfAttachment: DatabaseMessageExtraPdfFile = {
+		base64Data: btoa('this is not a pdf'),
+		content: 'Extracted text of the broken PDF.',
+		name: 'broken.pdf',
+		processedAsImages: false,
+		size: 17,
+		type: AttachmentType.PDF
+	};
+</script>
+
+<!-- conversion fails once: shows the error and does not retry -->
+<Story
+	asChild
+	name="PdfConversionError"
+	play={async ({ canvas }) => {
+		await waitFor(() => expect(canvas.getByText('Failed to load PDF images')).toBeInTheDocument(), {
+			timeout: 5000
+		});
+
+		// a retry loop would flip back to the loading state
+		await new Promise((resolve) => setTimeout(resolve, 500));
+		await expect(canvas.queryByText('Converting PDF to images...')).not.toBeInTheDocument();
+		await expect(canvas.getByText('Failed to load PDF images')).toBeInTheDocument();
+	}}
+>
+	<div class={frameClass}>
+		<ChatAttachmentsPreview attachments={[brokenPdfAttachment]} class="min-h-0 flex-1" />
+	</div>
+</Story>
Author	SHA1	Message	Date
Xuan-Son NguyenandPascal	5851a8e505	model, mtmd: drama-3 tts support (#31613 ) * model, mtmd: drama-3 tts support * clean up * mtmd: drama3 helper stops at the codebook tower limit --------- Co-authored-by: Pascal <admin@serveurperso.com>	2026-11-04 15:15:09 +01:00
Aleksander Grygier	efbf6e80a5	ui: show text and page thumbnails for PDF attachments, stop endless conversion retries (#31468 ) * fix(ui): show extracted text in PDF attachment preview Assisted-by: Claude (Opus 5.5 Medium) * feat(ui): show page thumbnails for PDF attachments sent as images Assisted-by: Claude (Opus 5.5 Medium) * ui: stop retrying PDF preview conversion after it fails A failed PDF-to-image conversion or text extraction reset its loading flag, which re-ran the effect and retried forever. Remember the failed item for each view and only retry when the item changes, so the error state stays on screen. Add a PdfConversionError story that feeds a broken PDF. Assisted-by: Claude (Opus 5.5 Medium) * ui: keep the text of PDFs sent as images and tidy the PDF preview Store the extracted text of a PDF sent as images in its content, so the Text view reads it instead of parsing the PDF again; only messages stored with the old placeholder extract it. The model still receives the images only. Guard retries on the error state, as the item remounts per id, skip an empty base64 payload, and share the status markup between views. Assisted-by: Claude (Opus 5.5 Medium)	2026-11-04 15:46:51 +02:00
Masashi Yoshimura	b393d73bb0	webgpu: return maxBufferSize from get_max_size and avoid whole-buffer bindings in memset/clear (#31348 ) * Fix get_max_size and improve the memset shader * Fix buffer_memset to handle the buffer size more than max binding size.	2026-11-04 15:46:06 +02:00
Neo ZhangandSigbjørn Skjæret	00df7c22d9	sycl : fix docker issue (#31605 ) * add oneDNN in build image * Update .devops/intel.Dockerfile Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>	2026-11-04 15:45:29 +02:00
Aaron Thomas	034e2535a9	gguf-py : validate every tokenizer merge pair, not only the first (#31631 ) * gguf-py : validate every tokenizer merge pair, not only the first The transformers 4.45 pair format was detected by checking merges[0], then every element was converted unchecked: a non-pair element raised a bare TypeError, and pairs with the wrong arity or non-string parts were silently mangled into invalid merges. Validate each element (2-string list) up front and raise the existing "Unknown tokenizer merges format" error for anything else. Assisted-by: Devin * gguf-py : remove test_vocab.py per review	2026-11-04 14:42:39 +01:00
Adrien Gallouët	52e53a76df	common : delete common_get_path_from_env() (#31444 ) Signed-off-by: Adrien Gallouët <angt@huggingface.co>	2026-11-04 14:21:38 +01:00
Georgi Gerganov	cd959bae99	llama.cpp : bump version to 0.9.0 (#31646 ) * llama.cpp : bump version to 0.9.0 * authors : update	2026-11-04 14:06:26 +02:00
Georgi Gerganov	13ec8f614c	sync : ggml	2026-11-04 12:00:37 +02:00
Georgi Gerganov	cca14f64fe	ggml : bump version to 0.29.0 (ggml/1680) * ggml : bump version to 0.29.0 * authors : update	2026-11-04 12:00:37 +02:00
Aman Gupta	e853ec9a8c	llama : follow ctx_other devices and flash attention in Gemma4 assistant (#31635 ) * llama : follow ctx_other devices and flash attention in Gemma4 assistant * llama : include <algorithm> for std::any_of	2026-11-04 11:23:25 +02:00
Kirill A. Korinsky	feaf800060	ggml : Gate RVV code on the target extension (#31607 ) The RISC-V vector intrinsic specification requires __riscv_v_intrinsic to be defined when the compiler supports the API, even when the vector extension is not selected. Clang therefore defines it for rv64gc, causing ggml to enter RVV branches where vector types are rejected because the required ISA extensions are absent. Use __riscv_v for optional RVV implementations; this macro represents the selected target ISA. Retain __riscv_v_intrinsic checks in SpaceMIT sources, where they validate mandatory compiler API support rather than select optional vector code. Closes: https://github.com/ggml-org/ggml/issues/1535	2026-11-04 11:11:52 +02:00
Georgi Gerganov	4c4ca5b188	common : do not allow `_` in sidecar (#31639 )	2026-11-04 10:32:45 +02:00
EugeoSynthesisThirtyTwoandYour Name	162e2bb2e4	ggml : fix bf16 (#31609 ) Co-authored-by: Your Name <you@example.com>	2026-11-04 10:10:22 +02:00
lhez	6900fb5ac8	opencl: improve q2_0 GEMV for Adreno (#31577 ) * opencl: improve q2_0 decoding and small f32 matmuls Assisted-by: Codex * opencl: optimize small f32 matmuls on Adreno 8xx Assisted-by: Codex	2026-11-04 10:07:14 +02:00
Pascal	bbb02c93e9	ggml: bound views by the bytes their strides reach (#31584 ) * ggml: bound views by the bytes their strides reach A view was checked against its source with the size of a contiguous tensor of the same shape, computed before its strides were set. Views whose rows overlap failed that check although they stay inside their source, and views with wider strides passed it while reaching past the end. The view functions now pass their strides to the tensor constructor, which checks the bytes they actually reach, like ggml_nbytes. The conv windows of delta_net are created with all their rows at once instead of being widened after the check. * ggml: assert that a strided view starts from contiguous elements The view functions set nb[0] to the type size, so the source must have contiguous elements along dim 0 for the view to read them.	2026-11-04 09:47:11 +02:00
Georgi Gerganov	960b1d51b7	ggml : write host-buffer tensor data directly in gguf writer (#31144 ) The gguf writer staged each tensor into a temporary vector before writing it. For host-buffer tensors the data pointer is directly writable, so skip the staging copy and write it straight to the file. This avoids the per-tensor allocation, zero-init and memcpy. Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)	2026-11-04 09:29:26 +02:00
Georgi GerganovandPascal	4b29d02629	common : fix -hf tag matching inside longer quant names (#31620 ) * common : fix -hf tag matching inside longer quant names ":F16" picked a BF16 file when the repo had both, since "BF16" contains "F16". The same applied to ":Q4_XS" vs "IQ4_XS". The tag now has to start at a token boundary in the path. Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB) * common : accept underscore as -hf tag boundary Some repos separate the quant with an underscore, e.g. stepfun-ai/Step-3.5-Flash-GGUF-Q4_K_S ships step3p5_flash_Q4_K_S-00001-of-00012.gguf, which :Q4_K_S no longer matched. Add _ to the token boundary and pin the tag and default lookup cases. --------- Co-authored-by: Pascal <admin@serveurperso.com>	2026-11-04 09:22:54 +02:00
Todor BoinovskiandMax Krasnyansky	bee23322f1	hexagon: add LIGHTNING_INDEXER support (#31289 ) * hexagon: initial support for LIGHTNING_INDEXER op * hex-lght-idxr: hmx for prefill * hex-lght-idxr: q8_0 decode gather transpose, two-block decode compute * hex-lght-idxr: pipelined hmx path for decode and prefill * hex-lid: add support for K (f32) * hex-lid: add support for K (f32) * hex-lid: add support for q4_0 * hex-lid: remove scalar VTCM reads in the HMX q4_0 K prep and remove vector spills * hex-lid: host/kernel kparams alighnment and remove sw devices in kernel * hex-lit: update dma pipeline to use canonical dbl buffering in all paths * hex-lit: add GGML_HEXAGON_LI_SELECT to force hmx/hvx paths * hex-lit: add support for dims collapse * hex-lit: make kparams constistent with other ops --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>	2026-11-03 21:17:30 -08:00