sasori TQ3P (K=3) do FLUX.2-klein-9B + text encoder ternario + kernel ggml + medicoes

.gitattributes

@@ -33,3 +33,13 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
 *.zip filter=lfs diff=lfs merge=lfs -text
 *.zst filter=lfs diff=lfs merge=lfs -text
 *tfevents* filter=lfs diff=lfs merge=lfs -text
+flux-2-klein-9b-TQ3P.gguf filter=lfs diff=lfs merge=lfs -text
+samples/fp_p0_s42.png filter=lfs diff=lfs merge=lfs -text
+samples/fp_p6_s42.png filter=lfs diff=lfs merge=lfs -text
+samples/k3_p0_s42.png filter=lfs diff=lfs merge=lfs -text
+samples/k3_p6_s42.png filter=lfs diff=lfs merge=lfs -text
+samples/k3te_p0_s42.png filter=lfs diff=lfs merge=lfs -text
+samples/k3te_p6_s42.png filter=lfs diff=lfs merge=lfs -text
+samples/q4km_p0_s42.png filter=lfs diff=lfs merge=lfs -text
+samples/q4km_p6_s42.png filter=lfs diff=lfs merge=lfs -text
+text_encoder/Qwen3-8B-TQ3P.gguf filter=lfs diff=lfs merge=lfs -text

README.md

@@ -0,0 +1,194 @@
+---
+license: other
+license_name: flux-2-klein-non-commercial-license
+license_link: https://huggingface.co/black-forest-labs/FLUX.2-klein-9B/blob/main/LICENSE.md
+base_model:
+  - black-forest-labs/FLUX.2-klein-9B
+  - Qwen/Qwen3-8B
+pipeline_tag: text-to-image
+tags:
+  - text-to-image
+  - flux
+  - flux.2
+  - gguf
+  - quantization
+  - ternary
+  - sasori
+  - stable-diffusion-cpp
+library_name: sasori
+---
+
+# sasori TQ3P — FLUX.2 [klein] 9B ternarizado (K=3 trit-planes)
+
+**O que é:** o denoiser (DiT) do FLUX.2 [klein] 9B e o seu text encoder (Qwen3-8B) convertidos para
+**TQ3P**, o formato ternário multi-plano do [sasori](https://github.com/Adriel007/PhD-propose): cada peso
+é a soma de **3 planos de trits** `{-1,0,+1}` com escala f16 por grupo de 256 — **6,1875 bits por peso**.
+Sem retreino, sem dados de calibração (data-free, layer-local).
+
+> ### ⚠️ Precisa de um runtime patchado — não roda no stable-diffusion.cpp oficial
+> `TQ3P` é um tipo de quantização **customizado do ggml** (type id 43) que o upstream não conhece: um
+> binário sem patch **recusa a carregar** estes arquivos. Os patches e o script de build estão em
+> [`kernel/`](./kernel) — são 25 hunks, aplicam limpo, e o build leva ~5 min. Veja
+> [Como rodar](#como-rodar).
+
+## Arquivos
+
+| arquivo | o que é | tamanho |
+|---|---|---|
+| `flux-2-klein-9b-TQ3P.gguf` | denoiser DiT, K=3 (112 tensores ternários / 89 em FP16) | **6,94 GiB** |
+| `text_encoder/Qwen3-8B-TQ3P.gguf` | text encoder, K=3 | **7,33 GiB** |
+| `kernel/sdcpp-cpu.patch` | tipo TQ{K}P no ggml-CPU (bloco, `vec_dot` AVX2 `maddubs`) | 17 hunks |
+| `kernel/sdcpp-cuda.patch` | gate CUDA (`dequant→cuBLAS`) para o denoiser rodar na GPU | 8 hunks |
+| `kernel/build.sh` · `kernel/BUILD.md` | build reproduzível do runtime | — |
+| `bench/` | os JSONs de todas as medições abaixo | — |
+| `samples/` | amostras pareadas FP16 vs TQ3P (mesmo prompt, mesma seed) | — |
+
+**Você também precisa** (não redistribuídos aqui):
+- **VAE**: `full_encoder_small_decoder.safetensors` de
+  [`black-forest-labs/FLUX.2-small-decoder`](https://huggingface.co/black-forest-labs/FLUX.2-small-decoder)
+  (o que foi usado em todas as medições; sem gate de licença). A VAE oficial `flux2_ae.safetensors`
+  vive no repo gated [`FLUX.2-dev`](https://huggingface.co/black-forest-labs/FLUX.2-dev) e também serve.
+
+## Números medidos
+
+Tudo medido **nesta conversão**, num pod A100-80GB PCIe, `sd.cpp master-802-e92e86f` + os patches
+deste repo, CUDA sm_80, 1024×1024, 4 steps (o klein-9B é step-distilled), euler, `--diffusion-fa`.
+
+### VRAM e footprint (medidos)
+
+| braço | DiT (arquivo) | text encoder (arquivo) | **pacote** | VRAM medida (1024²) |
+|---|---|---|---|---|
+| FP16 / BF16 (16 bpw) | 16,91 GiB | 15,26 GiB | **32,17 GiB** | 33 058 MB |
+| DiT TQ3P + TE BF16 | 6,94 GiB | 15,26 GiB | 22,20 GiB | 22 853 MB |
+| **DiT TQ3P + TE TQ3P** (6,1875 bpw) | **6,94 GiB** | **7,33 GiB** | **14,27 GiB** | **14 729 MB** |
+| Q4_K_M + Q8_0 (baseline) | 5,50 GiB | 8,11 GiB | 13,61 GiB | 13 425 MB |
+
+### Velocidade de denoise — s/it, mediana de 16 gerações por braço
+
+| braço | s/it (mediana) | faixa | vs FP16 |
+|---|---|---|---|
+| FP16 | 1,660 | 1,577–1,768 | — |
+| **DiT TQ3P** | **1,200** | 1,177–1,230 | **1,38× mais rápido** |
+| **DiT TQ3P + TE TQ3P** | **1,221** | 1,195–1,278 | **1,36× mais rápido** |
+| Q4_K_M | 1,619 | 1,603–1,657 | 1,03× mais rápido |
+
+`sampling completed` do sd-cli mede só o laço de denoise (exclui load do modelo, encode do prompt e
+decode da VAE), então este contraste isola o DiT. O text encoder roda uma vez por imagem, fora do
+laço — daí `TQ3P` e `TQ3P + TE TQ3P` empatarem (1,8 % = ruído).
+
+### Fidelidade — CLIP-score pareado (ViT-L/14, N=16 pares: 8 prompts × 2 seeds)
+
+| braço | CLIP-score | retenção vs FP16 | Wilcoxon p | placar par-a-par (FP/braço) |
+|---|---|---|---|---|
+| FP16 | 0,7303 ± 0,0589 | — | — | — |
+| **DiT TQ3P** | 0,7285 ± 0,0702 | **99,8 %** | 0,744 (n.s.) | **8/8 — empate** |
+| **DiT TQ3P + TE TQ3P** | 0,7276 ± 0,0631 | 99,6 % | 0,464 (n.s.) | 8/8 — empate |
+| Q4_K_M + Q8_0 | 0,7227 ± 0,0608 | 99,0 % | 0,231 (n.s.) | 9/7 |
+
+Inspeção visual das amostras pareadas (`samples/`): nos 4 braços a contagem do prompt "three yellow
+rubber ducks" sai correta e o fotorrealismo se mantém; o que muda é a trajetória de denoise
+(composição levemente diferente), não a qualidade.
+
+### Leitura honesta destes números
+
+- **O TQ3P não é dominado pelo Q4_K_M aqui — é um trade-off real.** O DiT ternário é **26 % maior**
+  que o `Q4_K_M` (6,94 vs 5,50 GiB) e ao mesmo tempo **35 % mais rápido** no denoise (1,20 vs
+  1,62 s/it). Como *pacote* (DiT + text encoder) a diferença de tamanho quase desaparece — 14,27 vs
+  13,61 GiB — porque o Qwen3-8B em TQ3P (7,33 GiB) fica **menor** que em Q8_0 (8,11 GiB).
+- **Isto contraria a expectativa que trouxemos do run anterior** (SD3.5-Medium, onde TQ3P 0,343 s/it
+  ≈ FP 0,347 ≈ Q4_K_M 0,329 — empate nos três). Aqui o ternário ganha. **Hipótese** (não isolada
+  neste run): o gate CUDA roteia TQ{K}P por `dequant → cuBLAS`, e a 1024² o denoise processa ~4096
+  tokens de imagem por passo — regime de GEMM grande, onde o cuBLAS bate o kernel `mmq` que serve o
+  `Q4_K`; contra o FP16, o ternário lê 2,4× menos bytes pelo mesmo caminho de GEMM. Falsificar isso
+  exigiria variar resolução e comparar os caminhos de kernel isoladamente — **não foi feito**.
+- **Nenhuma diferença de fidelidade é detectável a N=16** — em nenhum braço, inclusive o baseline.
+  Isso **não** é o mesmo que "TQ3P = FP16": N=16 é subpotente para um Δ de ~0,2 pp. O que se pode
+  afirmar é mais fraco e mais honesto: **o placar par-a-par do TQ3P contra o FP16 é 8/8**, sem
+  tendência direcional (no run do SD3.5 o FP ganhava 11/16), e a inspeção visual não mostra
+  degradação.
+- **CLIP-score não é FID.** Mede alinhamento imagem-texto pareado, não fidelidade perceptual da
+  distribuição. Nenhum FID de Fréchet foi computado.
+- **Escopo da medição:** 1 modelo, 1 GPU (A100-80 PCIe), 1024², 4 steps, euler, 2 seeds, 8 prompts,
+  1 CLIP (ViT-L/14), VAE `small-decoder`. Não medido: outras resoluções, CPU, ARM, edição
+  (`-r` image-to-image), outros samplers, FID.
+
+## Como rodar
+
+```bash
+# 1) runtime com o kernel ternário (~5 min)
+git clone https://huggingface.co/FardoX/sasori-flux2-klein-9b-tq3p && cd sasori-flux2-klein-9b-tq3p
+CUDA_ARCH=86 bash kernel/build.sh          # 80=A100 86=3090 89=4090 90=H100; BUILD_CUDA=OFF p/ CPU
+strings stable-diffusion.cpp/build/bin/sd-cli | grep TQ3P   # sanity: o tipo tem que existir
+
+# 2) pesos (este repo + a VAE)
+hf download FardoX/sasori-flux2-klein-9b-tq3p --local-dir .
+hf download black-forest-labs/FLUX.2-small-decoder full_encoder_small_decoder.safetensors --local-dir vae
+
+# 3) gerar
+stable-diffusion.cpp/build/bin/sd-cli \
+  --diffusion-model flux-2-klein-9b-TQ3P.gguf \
+  --llm text_encoder/Qwen3-8B-TQ3P.gguf \
+  --vae vae/full_encoder_small_decoder.safetensors \
+  -p "a lovely cat sitting on a wooden table, soft window light" \
+  --cfg-scale 1.0 --steps 4 -W 1024 -H 1024 --sampling-method euler \
+  --diffusion-fa -o cat.png -v
+```
+
+### Paralelismo — o que realmente acelera
+
+- **`--diffusion-fa`** (flash-attention no denoiser) e **4 steps**: o klein-9B é step-distilled, então
+  4 steps é a operação normal, não um atalho.
+- **NÃO passe `--offload-to-cpu` se você tem VRAM.** A doc oficial do sd.cpp usa essa flag em todos os
+  exemplos, e ela põe **todos os pesos na RAM** (`VRAM 0.00MB` no log) — medido aqui: com `--offload-to-cpu` o log mostra `VRAM 0.00MB` e o sampling de 4 steps a 512² leva **27,8 s**; sem a flag, a 1024² (4× mais pixels) leva **4,8 s**.
+  Use a flag só se o modelo não couber na sua GPU.
+- **VRAM necessária** (TQ3P, 1024², sem offload): **14,7 GB** com o text encoder também em TQ3P (7,1 GB o DiT + 7,5 GB o TE + 0,1 GB a VAE), ou 22,9 GB mantendo o text encoder em BF16.
+- **CPU**: o kernel TQ{K}P usa `_mm256_maddubs_epi16` (AVX2) e é um mpGEMM direto — o trit multiplica
+  sem dequantizar. Passe `-t <núcleos>`. Em CPU, medições anteriores do projeto colocam o TQ2P
+  empatado com o Q4_K_M; o **TQ3P perde de ambos** (mais bytes para ler por peso).
+
+### Ternarizar você mesmo (e o que aprendi sobre paralelizar isso)
+
+```bash
+python3 sasori/scripts/inject_tqkp_diffusion.py IN_F16.gguf OUT_TQ3P.gguf \
+        --K 3 --group 256 --device cuda --dry-run   # audita a deny-list ANTES de gastar GPU
+```
+
+O ajuste é **layer-local e data-free** (cada matriz é independente), então a tentação é paralelizar em
+processos. **Medido neste run:** rodar a conversão do DiT (GPU) e a do text encoder (CPU, 32 workers) ao
+mesmo tempo **piorou** o caminho crítico de **3,7 s → 25 s por tensor (6,7× mais lento)**. Motivo: o
+solver joint-ridge na GPU fica em ~7 % de utilização; o gargalo real é o **`pack` dos trits em numpy na
+CPU** — os dois jobs disputavam o mesmo recurso, não recursos disjuntos. **Serializar os dois jobs na
+GPU, com toda a CPU para o pack de um só, foi mais rápido:** DiT 8,72 G params em 15,3 min, TE em
+**2,7 min** (6,94 G params).
+
+## O que fica em FP16 (e por quê)
+
+Dos 201 tensores do DiT, **89 (354 M de 9,07 G params = 3,9 %)** continuam em FP16 — custa ~0,7 GB e
+protege os pontos onde um erro de trit **não** é local:
+
+| mantido FP16 | por quê |
+|---|---|
+| `*modulation*` (adaLN, 285 M) | emite shift/scale/gate por bloco a partir do timestep: um erro aqui reescala **toda** a ativação do bloco. A literatura de PTQ para DiT (Q-DiT, PTQ4DiT, ViDiT-Q) aponta adaLN/modulação como o sítio sensível. |
+| `txt_in` (50 M) | única porta do prompt para dentro do denoiser; ternarizar poupa ~0,1 GB e arrisca todo o condicionamento. |
+| `time_in` (17,8 M) | MLP de condicionamento do timestep. |
+| `final_layer.linear` (0,5 M) | projeção de volta ao espaço latente — mesma razão pela qual o caminho LLM do sasori mantém `output.weight` em FP. |
+| `img_in`, normas, `key/query_norm` | `img_in` tem in_dim 128 (não múltiplo de 256); normas são 1-D. |
+
+**Nota de rigor:** essas escolhas são **hipóteses fundamentadas** (mecanismo + literatura de PTQ de
+DiT), **não uma ablação medida neste run**. Ternarizar a modulação e medir o custo é o experimento
+seguinte — não foi feito aqui.
+
+## Licença
+
+O denoiser é um **derivado do FLUX.2 [klein] 9B** e herda a **FLUX Non-Commercial License** da Black
+Forest Labs: **uso não-comercial apenas**, e as mesmas restrições se aplicam a qualquer derivado deste.
+Leia a licença antes de usar. O text encoder é derivado do **Qwen3-8B (Apache-2.0)**. Os patches do
+kernel e os scripts são do projeto sasori.
+
+## Citação / proveniência
+
+- Pesos-fonte: `unsloth/FLUX.2-klein-9B-GGUF` (`flux-2-klein-9b-F16.gguf`) e `unsloth/Qwen3-8B-GGUF`
+  (`Qwen3-8B-BF16.gguf`) — GGUFs full-precision, sem quantização em cascata.
+- Runtime: `leejet/stable-diffusion.cpp` `master-802-e92e86f`, ggml `eced84c8` + os patches deste repo.
+- Conversão: sasori `inject_tqkp_diffusion.py --K 3 --group 256 --device cuda`, niter=25.
+- Método: PTQ ternário multi-plano (joint-ridge alternado), data-free, layer-local.

bench/clipscore.json

@@ -0,0 +1,173 @@
+{
+ "n_pairs": 16,
+ "ref": "fp",
+ "clip": "openai/clip-vit-large-patch14",
+ "pairs": [
+  {
+   "prompt_i": 0,
+   "seed": 42
+  },
+  {
+   "prompt_i": 0,
+   "seed": 1337
+  },
+  {
+   "prompt_i": 1,
+   "seed": 42
+  },
+  {
+   "prompt_i": 1,
+   "seed": 1337
+  },
+  {
+   "prompt_i": 2,
+   "seed": 42
+  },
+  {
+   "prompt_i": 2,
+   "seed": 1337
+  },
+  {
+   "prompt_i": 3,
+   "seed": 42
+  },
+  {
+   "prompt_i": 3,
+   "seed": 1337
+  },
+  {
+   "prompt_i": 4,
+   "seed": 42
+  },
+  {
+   "prompt_i": 4,
+   "seed": 1337
+  },
+  {
+   "prompt_i": 5,
+   "seed": 42
+  },
+  {
+   "prompt_i": 5,
+   "seed": 1337
+  },
+  {
+   "prompt_i": 6,
+   "seed": 42
+  },
+  {
+   "prompt_i": 6,
+   "seed": 1337
+  },
+  {
+   "prompt_i": 7,
+   "seed": 42
+  },
+  {
+   "prompt_i": 7,
+   "seed": 1337
+  }
+ ],
+ "arms": {
+  "fp": {
+   "mean": 0.730270504951477,
+   "std": 0.05890607833862305,
+   "per_pair": [
+    0.7898228168487549,
+    0.7958123087882996,
+    0.69093257188797,
+    0.6075867414474487,
+    0.8519022464752197,
+    0.7870793342590332,
+    0.7254836559295654,
+    0.7475548982620239,
+    0.7404217720031738,
+    0.6994374990463257,
+    0.733026921749115,
+    0.7118973731994629,
+    0.6827402710914612,
+    0.7448596954345703,
+    0.7178919911384583,
+    0.6578782796859741
+   ]
+  },
+  "k3": {
+   "mean": 0.7284521460533142,
+   "std": 0.0701667070388794,
+   "per_pair": [
+    0.8078278303146362,
+    0.7883470058441162,
+    0.6257125735282898,
+    0.5858814716339111,
+    0.8263368606567383,
+    0.8231972455978394,
+    0.7344119548797607,
+    0.740487813949585,
+    0.7451135516166687,
+    0.682620108127594,
+    0.6912571787834167,
+    0.7204871773719788,
+    0.7584888339042664,
+    0.7658802270889282,
+    0.7210639119148254,
+    0.638121485710144
+   ],
+   "retention_pct": 99.7509994506836,
+   "wilcoxon_p": 0.743560791015625,
+   "ref_wins": 8,
+   "arm_wins": 8
+  },
+  "k3te": {
+   "mean": 0.7276306748390198,
+   "std": 0.06310403347015381,
+   "per_pair": [
+    0.7939058542251587,
+    0.8016656637191772,
+    0.6442081332206726,
+    0.5878970623016357,
+    0.8186874389648438,
+    0.7942061424255371,
+    0.7080155611038208,
+    0.7588667869567871,
+    0.7321103811264038,
+    0.7036865949630737,
+    0.6975305080413818,
+    0.7260567545890808,
+    0.7676522731781006,
+    0.7142478823661804,
+    0.7418999075889587,
+    0.6514531970024109
+   ],
+   "retention_pct": 99.63851928710938,
+   "wilcoxon_p": 0.4637451171875,
+   "ref_wins": 8,
+   "arm_wins": 8
+  },
+  "q4km": {
+   "mean": 0.722662627696991,
+   "std": 0.060831647366285324,
+   "per_pair": [
+    0.7809169292449951,
+    0.7693759202957153,
+    0.6436231732368469,
+    0.6148441433906555,
+    0.8311357498168945,
+    0.7981688976287842,
+    0.7179467678070068,
+    0.7483753561973572,
+    0.7414851188659668,
+    0.6667605638504028,
+    0.7295082807540894,
+    0.7041403651237488,
+    0.7126471996307373,
+    0.7543122172355652,
+    0.7212250232696533,
+    0.628136157989502
+   ],
+   "retention_pct": 98.95821380615234,
+   "wilcoxon_p": 0.231201171875,
+   "ref_wins": 9,
+   "arm_wins": 7
+  }
+ }
+}

bench/plan_dit_k3.json

@@ -0,0 +1,1426 @@
+{
+ "ternarize": [
+  [
+   "double_blocks.0.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.0.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.0.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.0.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.0.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.0.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.0.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.0.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.1.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.1.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.1.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.1.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.1.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.1.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.1.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.1.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.2.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.2.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.2.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.2.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.2.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.2.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.2.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.2.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.3.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.3.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.3.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.3.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.3.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.3.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.3.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.3.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.4.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.4.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.4.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.4.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.4.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.4.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.4.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.4.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.5.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.5.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.5.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.5.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.5.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.5.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.5.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.5.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.6.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.6.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.6.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.6.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.6.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.6.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.6.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.6.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.7.img_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.7.img_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.7.img_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.7.img_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.7.txt_attn.proj.weight",
+   [
+    4096,
+    4096
+   ]
+  ],
+  [
+   "double_blocks.7.txt_attn.qkv.weight",
+   [
+    4096,
+    12288
+   ]
+  ],
+  [
+   "double_blocks.7.txt_mlp.0.weight",
+   [
+    4096,
+    24576
+   ]
+  ],
+  [
+   "double_blocks.7.txt_mlp.2.weight",
+   [
+    12288,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.0.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.0.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.1.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.1.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.10.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.10.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.11.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.11.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.12.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.12.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.13.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.13.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.14.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.14.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.15.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.15.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.16.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.16.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.17.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.17.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.18.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.18.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.19.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.19.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.2.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.2.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.20.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.20.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.21.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.21.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.22.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.22.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.23.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.23.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.3.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.3.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.4.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.4.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.5.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.5.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.6.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.6.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.7.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.7.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.8.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.8.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ],
+  [
+   "single_blocks.9.linear1.weight",
+   [
+    4096,
+    36864
+   ]
+  ],
+  [
+   "single_blocks.9.linear2.weight",
+   [
+    16384,
+    4096
+   ]
+  ]
+ ],
+ "keep_fp16": [
+  [
+   "double_blocks.0.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.0.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.0.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.0.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.1.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.1.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.1.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.1.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.2.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.2.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.2.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.2.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.3.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.3.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.3.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.3.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.4.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.4.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.4.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.4.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.5.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.5.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.5.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.5.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.6.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.6.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.6.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.6.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.7.img_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.7.img_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.7.txt_attn.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_blocks.7.txt_attn.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "double_stream_modulation_img.lin.weight",
+   [
+    4096,
+    24576
+   ],
+   "deny:modulation"
+  ],
+  [
+   "double_stream_modulation_txt.lin.weight",
+   [
+    4096,
+    24576
+   ],
+   "deny:modulation"
+  ],
+  [
+   "final_layer.adaLN_modulation.1.weight",
+   [
+    4096,
+    8192
+   ],
+   "deny:modulation"
+  ],
+  [
+   "final_layer.linear.weight",
+   [
+    4096,
+    128
+   ],
+   "deny:final_layer.linear"
+  ],
+  [
+   "img_in.weight",
+   [
+    128,
+    4096
+   ],
+   "deny:img_in"
+  ],
+  [
+   "single_blocks.0.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.0.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.1.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.1.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.10.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.10.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.11.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.11.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.12.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.12.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.13.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.13.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.14.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.14.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.15.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.15.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.16.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.16.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.17.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.17.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.18.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.18.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.19.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.19.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.2.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.2.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.20.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.20.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.21.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.21.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.22.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.22.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.23.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.23.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.3.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.3.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.4.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.4.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.5.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.5.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.6.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.6.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.7.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.7.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.8.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.8.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.9.norm.key_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_blocks.9.norm.query_norm.scale",
+   [
+    128
+   ],
+   "not-2D"
+  ],
+  [
+   "single_stream_modulation.lin.weight",
+   [
+    4096,
+    12288
+   ],
+   "deny:modulation"
+  ],
+  [
+   "time_in.in_layer.weight",
+   [
+    256,
+    4096
+   ],
+   "deny:time_in"
+  ],
+  [
+   "time_in.out_layer.weight",
+   [
+    4096,
+    4096
+   ],
+   "deny:time_in"
+  ],
+  [
+   "txt_in.weight",
+   [
+    12288,
+    4096
+   ],
+   "deny:txt_in"
+  ]
+ ],
+ "n_ternarize": 112,
+ "n_keep_fp16": 89,
+ "params_ternarize_M": 8724.15232,
+ "params_keep_M": 354.428928
+}

bench/speed.json

@@ -0,0 +1,54 @@
+{
+ "fp": {
+  "n_gen": 16,
+  "s_per_it_median": 1.66,
+  "s_per_it_min": 1.5775,
+  "s_per_it_max": 1.7675,
+  "sampling_total_median_s": 6.64,
+  "steps": 4,
+  "total_params_MB": 33058.29,
+  "vram_MB": 33058.29,
+  "ram_MB": 0.0,
+  "text_encoder_MB": 15623.18,
+  "diffusion_model_MB": 17316.04
+ },
+ "k3": {
+  "n_gen": 16,
+  "s_per_it_median": 1.2,
+  "s_per_it_min": 1.1775,
+  "s_per_it_max": 1.23,
+  "sampling_total_median_s": 4.8,
+  "steps": 4,
+  "total_params_MB": 22853.29,
+  "vram_MB": 22853.29,
+  "ram_MB": 0.0,
+  "text_encoder_MB": 15623.18,
+  "diffusion_model_MB": 7111.04
+ },
+ "k3te": {
+  "n_gen": 16,
+  "s_per_it_median": 1.2212,
+  "s_per_it_min": 1.195,
+  "s_per_it_max": 1.2775,
+  "sampling_total_median_s": 4.885,
+  "steps": 4,
+  "total_params_MB": 14728.54,
+  "vram_MB": 14728.54,
+  "ram_MB": 0.0,
+  "text_encoder_MB": 7498.43,
+  "diffusion_model_MB": 7111.04
+ },
+ "q4km": {
+  "n_gen": 16,
+  "s_per_it_median": 1.6187,
+  "s_per_it_min": 1.6025,
+  "s_per_it_max": 1.6575,
+  "sampling_total_median_s": 6.475,
+  "steps": 4,
+  "total_params_MB": 13424.88,
+  "vram_MB": 13424.88,
+  "ram_MB": 0.0,
+  "text_encoder_MB": 7669.77,
+  "diffusion_model_MB": 5636.04
+ }
+}

flux-2-klein-9b-TQ3P.gguf

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:5fba79be29f1f36beed03ed47a0e5fb183d43869275f864c2beb4e0b61c472f8
+size 7456479520

kernel/BUILD.md

@@ -0,0 +1,74 @@
+# O kernel TQ{K}P — o que os patches fazem e como buildar
+
+Os GGUFs deste repo usam um tipo de quantização **customizado do ggml**. Sem os patches abaixo, o
+`stable-diffusion.cpp` oficial **não carrega** os arquivos (ele nem sabe que o type id 43 existe).
+Isto não é empacotamento opcional: é o runtime.
+
+## Build rápido
+
+```bash
+CUDA_ARCH=86 bash build.sh        # 80=A100  86=3090/A6000  89=4090  90=H100
+BUILD_CUDA=OFF bash build.sh      # CPU-only (AVX2)
+SDCPP_REF=<tag|sha> bash build.sh # fixar a revisão do sd.cpp
+```
+
+Saída: `stable-diffusion.cpp/build/bin/sd-cli` (e `sd-server`). Sanity check:
+
+```bash
+strings stable-diffusion.cpp/build/bin/sd-cli | grep TQ3P   # tem que imprimir algo
+```
+
+Requisitos: `cmake >= 3.24` (o do apt costuma ser 3.22 → `pip install -U cmake`), CUDA toolkit no PATH
+para o build de GPU (`export PATH=/usr/local/cuda/bin:$PATH`), e um compilador com AVX2 para a CPU.
+
+## O que cada patch faz
+
+### `sdcpp-cpu.patch` (17 hunks, 9 arquivos do ggml)
+
+Adiciona a família de tipos TQ{K}P ao ggml-CPU end to end:
+
+- **`include/ggml.h`** — os enums. `TQ2P=42`, `TQ3P=43`, `TQ1P=44`, `TQ4P=45`, variantes por
+  group-size (`_G32/_G64/_G128`) e por formato de escala (`_Q4K6`), e `GGML_TYPE_COUNT` vai de 42 → 62.
+  Os ids **não podem colidir** com tipos novos do upstream; `build.sh` aborta se `COUNT != 42` na
+  revisão escolhida.
+- **`src/ggml-common.h`** — o layout do bloco. `block_tq3p` = 3 planos × 256 trits empacotados a 2 bits
+  (`qs1/qs2/qs3`, 64 bytes cada) + 3 escalas `ggml_half`, com `static_assert` no tamanho. Um peso é
+  `w ≈ α₁t₁ + α₂t₂ + α₃t₃`, `tᵏ ∈ {-1,0,+1}`.
+- **`src/ggml-cpu/arch/x86/quants.c`** — o `vec_dot` AVX2. É um **mpGEMM direto**: `_mm256_maddubs_epi16`
+  multiplica trit (2 bits) por ativação int8 e acumula em int16 **sem dequantizar**. Para ternário a
+  multiplicação é trivial, então este caminho já é próximo do ótimo (uma implementação LUT-GEMM estilo
+  T-MAC foi medida no projeto e ficou **1,3× mais lenta**).
+- **`src/ggml-cpu/{ggml-cpu.c,quants.c,quants.h}`, `src/ggml-quants.{c,h}`, `src/ggml.c`** — registro do
+  tipo (traits, block size, `to_float`, dispatch) para o resto do grafo funcionar sem casos especiais.
+
+### `sdcpp-cuda.patch` (8 hunks, 4 arquivos)
+
+Faz o denoiser rodar na **GPU** em vez de abortar ou cair para a CPU:
+
+- **`src/ggml-cuda/mmvq.cu`** — `should_use_mmvq → false` para os tipos TQ{K}P. Sem isso o
+  `mul_mat_vec_q` bate num `GGML_ABORT` (não existe kernel mmvq para o tipo).
+- **`src/ggml-cuda/ggml-cuda.cu`** — `supports_op`, `should_fuse_mul_mat_vec_q → false` e o dispatch do
+  decode fundido, roteando TQ{K}P pelo caminho **`dequant → cuBLAS`**.
+- **`src/ggml-cuda/convert.{cu,cuh}`** — o dequant CUDA do bloco de 3 planos.
+
+Por que `dequant→cuBLAS` e não um kernel ternário dedicado: no denoiser o GEMM domina e o dequant é
+barato, então o custo fica igual ao FP16. Medido no SD3.5-Medium (run anterior do projeto): TQ3P
+0,343 s/it vs FP 0,347 vs Q4_K_M 0,329 — e ~200× mais rápido que o caminho CPU (~54 s/it).
+
+## Compatibilidade verificada
+
+| item | valor |
+|---|---|
+| sd.cpp | `master-802-e92e86f` |
+| ggml (submódulo) | `eced84c8` (`GGML_TYPE_Q1_0=41`, `COUNT=42`) |
+| aplicação | 25 hunks, **0 rejeitos** (`git apply --check` passou nos dois patches) |
+| build | CUDA sm_80, `Release`, RC=0 |
+
+Se `git apply` falhar numa revisão futura, é drift de contexto: os hunks precisam ser re-portados
+(o bloco de enums é o único ponto que exige decisão — os ids têm que continuar livres).
+
+## Portabilidade
+
+O mesmo par de patches (com paths ajustados) vem do fork de `llama.cpp` do sasori — o op TQ{K}P é
+compartilhado entre os dois runtimes ggml. O kernel também compila para **wasm32** (gate de 12/12 no
+projeto), mas o runtime completo de difusão no browser não está integrado.

kernel/build.sh

@@ -0,0 +1,48 @@
+#!/usr/bin/env bash
+# Build stable-diffusion.cpp with the sasori TQ{K}P ternary kernel.
+#
+# The GGUFs in this repo use a CUSTOM ggml quantization type (TQ3P, ggml type id 43) that upstream
+# stable-diffusion.cpp does NOT know: an unpatched binary will refuse to load them. These two
+# patches add the type end to end (block layout, CPU vec_dot, CUDA dispatch), so this script is
+# not optional packaging -- it is how you get a runtime that can read the weights.
+#
+# Verified 2026-07-29 against sd.cpp master-802-e92e86f / ggml eced84c8 (25 hunks, 0 rejects).
+set -euo pipefail
+HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+SDCPP_REF="${SDCPP_REF:-master}"          # pin to a tag/sha for a byte-reproducible build
+CUDA_ARCH="${CUDA_ARCH:-}"                # e.g. 80 (A100) 86 (3090) 89 (4090) 90 (H100); empty = native
+BUILD_CUDA="${BUILD_CUDA:-ON}"
+
+echo "== clone sd.cpp ($SDCPP_REF) =="
+[ -d stable-diffusion.cpp ] || git clone --recursive https://github.com/leejet/stable-diffusion.cpp
+cd stable-diffusion.cpp
+git fetch --all --tags -q
+git checkout -q "$SDCPP_REF"
+git submodule update --init --recursive -q
+
+echo "== apply the TQ{K}P kernel (into the ggml submodule) =="
+cd ggml
+# ggml type ids 42..60 must be FREE for the sasori types; if upstream ever claims them the check
+# below fails loudly instead of silently producing a binary that misreads every tensor.
+grep -q "GGML_TYPE_COUNT   = 42," include/ggml.h || {
+  echo "!! this ggml revision does not have COUNT=42: the TQ{K}P type ids would collide."
+  echo "!! pin SDCPP_REF to a revision where they are free, or re-port the enum block by hand."
+  exit 1; }
+for p in cpu cuda; do
+  [ "$p" = cuda ] && [ "$BUILD_CUDA" != ON ] && continue
+  git apply --check "$HERE/sdcpp-$p.patch" || { echo "!! sdcpp-$p.patch does not apply to $SDCPP_REF"; exit 1; }
+  git apply "$HERE/sdcpp-$p.patch"
+  echo "   applied sdcpp-$p.patch"
+done
+cd ..
+
+echo "== build =="
+CMAKE_ARGS=(-DCMAKE_BUILD_TYPE=Release -DSD_CUDA="$BUILD_CUDA")
+[ -n "$CUDA_ARCH" ] && CMAKE_ARGS+=(-DCMAKE_CUDA_ARCHITECTURES="$CUDA_ARCH")
+cmake -B build "${CMAKE_ARGS[@]}"
+cmake --build build -j"$(nproc)"
+
+echo
+echo "OK -> $(pwd)/build/bin/sd-cli"
+echo "Sanity check (the type must be known to the binary):"
+echo "  strings build/bin/sd-cli | grep TQ3P"

kernel/sdcpp-cuda.patch

@@ -0,0 +1,465 @@
+diff --git a/src/ggml-cuda/convert.cu b/src/ggml-cuda/convert.cu
+index 61630a3..03ca5fe 100644
+--- a/src/ggml-cuda/convert.cu
++++ b/src/ggml-cuda/convert.cu
+@@ -709,6 +709,262 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
+     }
+ }
+ 
++// ---- sasori ternary trit-plane types (TQ{K}P / TQ{K}P_G{g} / TQ{K}P_Q4K6) ----
++// K trit-planes summed. The per-plane 2-bit packing mirrors block_tq2_0's strided layout, so the
++// byte index equals the thread id within a plane; only the scale decode differs per family. This
++// dequant mirrors dequantize_row_tq{K}p* / q4k6_dequantize in ggml-quants.c. MVP path: dequant ->
++// cuBLAS fp16 GEMM (no mmvq/mmq). Activation goes through fp16 here, not Q8_K as on CPU, so expect
++// tiny numeric drift vs the CPU vec_dot (same class of caveat as PR ggml-org/llama.cpp#11183).
++
++// f16-scale K-plane dequant: fixed TQ{K}P == G=256, plus TQ{K}P_G{32,64,128}. nsub = 256/G scales/plane.
++template<typename dst_t, int K, int G>
++static __global__ void dequantize_block_tqkp_gv(const void * __restrict__ vx, dst_t * __restrict__ yy) {
++    const int64_t i   = blockIdx.x;                 // super-block (256 weights)
++    const int64_t tid = threadIdx.x;                // 0..63: byte within a plane
++    const int nsub = 256 / G;
++    const int qb   = QK_TQKP_SUPER / 4;             // 64 bytes per plane
++    const int bpb  = K*qb + K*nsub*2;               // bytes per super-block
++    const uint8_t * __restrict__ blk = (const uint8_t *) vx + i*bpb;
++    const half    * __restrict__ d   = (const half *) (blk + K*qb);
++    const int n = tid / 32;
++    const int l = tid - 32*n;
++    dst_t * y = yy + i*QK_TQKP_SUPER + 128*n;
++#pragma unroll
++    for (int bp = 0; bp < 4; ++bp) {
++        const int j   = 128*n + bp*32 + l;          // logical position in the super-block
++        const int sub = j / G;
++        float acc = 0.0f;
++#pragma unroll
++        for (int kk = 0; kk < K; ++kk) {
++            const int t = ((blk[kk*qb + tid] >> (bp*2)) & 3) - 1;   // trit in {-1,0,1}
++            acc += __half2float(d[kk*nsub + sub]) * (float) t;
++        }
++        y[l + bp*32] = (dst_t) acc;
++    }
++}
++
++// q4k6-scale K-plane dequant (g=32 fixed): per-plane fp16 super-scale, sign-flipped by the sub bit7,
++// times a 6-bit unsigned sub-scale. Mirrors q4k6_dequantize() in ggml-quants.c bit-for-bit.
++template<typename dst_t, int K>
++static __global__ void dequantize_block_tqkp_q4k6(const void * __restrict__ vx, dst_t * __restrict__ yy) {
++    const int64_t i   = blockIdx.x;
++    const int64_t tid = threadIdx.x;                // 0..63
++    const int nsub = QK_TQKP_SUPER / Q4K6_G;        // 8
++    const int qb   = QK_TQKP_SUPER / 4;             // 64
++    const int bpb  = K*qb + K*nsub + K*2;
++    const uint8_t * __restrict__ blk = (const uint8_t *) vx + i*bpb;
++    const uint8_t * __restrict__ sb  = blk + K*qb;                          // K*nsub sub-scale bytes
++    const half    * __restrict__ sup = (const half *)(blk + K*qb + K*nsub); // K fp16 super-scales
++    const int n = tid / 32;
++    const int l = tid - 32*n;
++    dst_t * y = yy + i*QK_TQKP_SUPER + 128*n;
++#pragma unroll
++    for (int bp = 0; bp < 4; ++bp) {
++        const int j = 128*n + bp*32 + l;
++        const int c = j / Q4K6_G;                   // sub-group index 0..7
++        float acc = 0.0f;
++#pragma unroll
++        for (int kk = 0; kk < K; ++kk) {
++            const uint8_t s = sb[kk*nsub + c];
++            uint32_t bits = __float_as_uint(__half2float(sup[kk]));
++            bits ^= ((uint32_t)(s & 0x80u)) << 24;  // flip sign iff sub bit7 set
++            const float sc = __uint_as_float(bits) * (float)(s & 0x3F);
++            const int t = ((blk[kk*qb + tid] >> (bp*2)) & 3) - 1;
++            acc += sc * (float) t;
++        }
++        y[l + bp*32] = (dst_t) acc;
++    }
++}
++
++#define DEQUANT_TQKP_GV_CUDA(NAME, K, G)                                                 \
++    template<typename dst_t>                                                             \
++    static void NAME(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) { \
++        const int nb = k / QK_TQKP_SUPER;                                                \
++        dequantize_block_tqkp_gv<dst_t, K, G><<<nb, 64, 0, stream>>>(vx, y);             \
++    }
++#define DEQUANT_TQKP_Q4K6_CUDA(NAME, K)                                                  \
++    template<typename dst_t>                                                             \
++    static void NAME(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) { \
++        const int nb = k / QK_TQKP_SUPER;                                                \
++        dequantize_block_tqkp_q4k6<dst_t, K><<<nb, 64, 0, stream>>>(vx, y);              \
++    }
++
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq1p_cuda,      1, 256)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq2p_cuda,      2, 256)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq3p_cuda,      3, 256)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq4p_cuda,      4, 256)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq1p_g32_cuda,  1, 32)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq1p_g64_cuda,  1, 64)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq1p_g128_cuda, 1, 128)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq2p_g32_cuda,  2, 32)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq2p_g64_cuda,  2, 64)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq2p_g128_cuda, 2, 128)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq3p_g32_cuda,  3, 32)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq3p_g64_cuda,  3, 64)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq3p_g128_cuda, 3, 128)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq4p_g32_cuda,  4, 32)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq4p_g64_cuda,  4, 64)
++DEQUANT_TQKP_GV_CUDA(dequantize_row_tq4p_g128_cuda, 4, 128)
++DEQUANT_TQKP_Q4K6_CUDA(dequantize_row_tq1p_q4k6_cuda, 1)
++DEQUANT_TQKP_Q4K6_CUDA(dequantize_row_tq2p_q4k6_cuda, 2)
++DEQUANT_TQKP_Q4K6_CUDA(dequantize_row_tq3p_q4k6_cuda, 3)
++
++
++// ---- sasori fused ternary mat-vec (decode path) ----
++// dst[row] = sum_col W[row,col] * x[col], reading PACKED TQ{K}P weights directly (no fp16 spill) ->
++// ~K*2/8 bytes/weight instead of 2 (fp16), i.e. much less weight bandwidth on the bandwidth-bound
++// decode. Activation x is fp32 (src1). Strided-256 K-plane layout mirrors dequantize_block_tqkp_gv.
++// One warp per output row; correctness-first (not yet memory-coalesced). MVP: f16-scale families only
++// (TQ{K}P == g256 + _G{32,64,128}); q4k6 still uses the dequant->cuBLAS path.
++template<int K, int G>
++static __global__ void mul_mat_vec_tqkp_kernel(
++        const void * __restrict__ vx, const float * __restrict__ x, float * __restrict__ dst,
++        const int ncols, const int nrows) {
++    const int row = blockIdx.x;
++    if (row >= nrows) return;
++    const int tid  = threadIdx.x;
++    const int warp = tid >> 5;
++    const int lane = tid & 31;
++    const int nwarps = blockDim.x >> 5;
++    const int nsub = 256 / G;
++    const int qb   = 64;
++    const int bpb  = K*qb + K*nsub*2;
++    const int nsuper = ncols / 256;
++    const uint8_t * __restrict__ row_base = (const uint8_t *) vx + (size_t) row * nsuper * bpb;
++    float acc = 0.0f;
++    for (int sb = warp; sb < nsuper; sb += nwarps) {      // warps split super-blocks
++        const uint8_t * __restrict__ blk = row_base + (size_t) sb * bpb;
++        const half    * __restrict__ d   = (const half *) (blk + K*qb);
++        const float   * __restrict__ xs  = x + sb*256;
++        #pragma unroll
++        for (int tb = lane; tb < 64; tb += 32) {          // coalesced byte reads within a super-block
++            const int n = tb >> 5, l = tb & 31;
++            #pragma unroll
++            for (int bp = 0; bp < 4; ++bp) {
++                const int j   = 128*n + bp*32 + l;
++                const int sub = j / G;
++                float w = 0.0f;
++                #pragma unroll
++                for (int kk = 0; kk < K; ++kk) {
++                    const int code = (blk[kk*qb + tb] >> (bp*2)) & 3;
++                    w += __half2float(d[kk*nsub + sub]) * (float) (code - 1);
++                }
++                acc += w * xs[j];
++            }
++        }
++    }
++    acc = warp_reduce_sum(acc);
++    __shared__ float sh[32];
++    if (lane == 0) sh[warp] = acc;
++    __syncthreads();
++    if (warp == 0) {
++        float v = (lane < nwarps) ? sh[lane] : 0.0f;
++        v = warp_reduce_sum(v);
++        if (lane == 0) dst[row] = v;
++    }
++}
++
++
++// q4k6-scale fused mat-vec (g=32): per-plane fp16 super-scale sign-flipped by sub bit7, times a 6-bit
++// unsigned sub-scale. Mirrors dequantize_block_tqkp_q4k6's scale decode; same coalesced structure.
++template<int K>
++static __global__ void mul_mat_vec_tqkp_q4k6_kernel(
++        const void * __restrict__ vx, const float * __restrict__ x, float * __restrict__ dst,
++        const int ncols, const int nrows) {
++    const int row = blockIdx.x;
++    if (row >= nrows) return;
++    const int tid  = threadIdx.x;
++    const int warp = tid >> 5;
++    const int lane = tid & 31;
++    const int nwarps = blockDim.x >> 5;
++    const int nsub = 256 / Q4K6_G;
++    const int qb   = 64;
++    const int bpb  = K*qb + K*nsub + K*2;
++    const int nsuper = ncols / 256;
++    const uint8_t * __restrict__ row_base = (const uint8_t *) vx + (size_t) row * nsuper * bpb;
++    float acc = 0.0f;
++    for (int sb = warp; sb < nsuper; sb += nwarps) {
++        const uint8_t * __restrict__ blk  = row_base + (size_t) sb * bpb;
++        const uint8_t * __restrict__ sbsc = blk + K*qb;
++        const half    * __restrict__ sup  = (const half *)(blk + K*qb + K*nsub);
++        const float   * __restrict__ xs   = x + sb*256;
++        #pragma unroll
++        for (int tb = lane; tb < 64; tb += 32) {
++            const int n = tb >> 5, l = tb & 31;
++            #pragma unroll
++            for (int bp = 0; bp < 4; ++bp) {
++                const int j = 128*n + bp*32 + l;
++                const int c = j / Q4K6_G;
++                float w = 0.0f;
++                #pragma unroll
++                for (int kk = 0; kk < K; ++kk) {
++                    const uint8_t s = sbsc[kk*nsub + c];
++                    uint32_t bits = __float_as_uint(__half2float(sup[kk]));
++                    bits ^= ((uint32_t)(s & 0x80u)) << 24;
++                    const float sc = __uint_as_float(bits) * (float)(s & 0x3F);
++                    const int t = ((blk[kk*qb + tb] >> (bp*2)) & 3) - 1;
++                    w += sc * (float) t;
++                }
++                acc += w * xs[j];
++            }
++        }
++    }
++    acc = warp_reduce_sum(acc);
++    __shared__ float sh[32];
++    if (lane == 0) sh[warp] = acc;
++    __syncthreads();
++    if (warp == 0) {
++        float v = (lane < nwarps) ? sh[lane] : 0.0f;
++        v = warp_reduce_sum(v);
++        if (lane == 0) dst[row] = v;
++    }
++}
++
++void ggml_cuda_mul_mat_vec_tqkp(ggml_backend_cuda_context & ctx,
++        const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
++    const int ncols = src0->ne[0];
++    const int nrows = src0->ne[1];
++    const float * x = (const float *) src1->data;
++    float * d = (float *) dst->data;
++    cudaStream_t stream = ctx.stream();
++    const dim3 grid(nrows, 1, 1);
++    switch (src0->type) {
++        case GGML_TYPE_TQ1P:      mul_mat_vec_tqkp_kernel<1,256><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ2P:      mul_mat_vec_tqkp_kernel<2,256><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ3P:      mul_mat_vec_tqkp_kernel<3,256><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ4P:      mul_mat_vec_tqkp_kernel<4,256><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ1P_G32:  mul_mat_vec_tqkp_kernel<1,32 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ1P_G64:  mul_mat_vec_tqkp_kernel<1,64 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ1P_G128: mul_mat_vec_tqkp_kernel<1,128><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ2P_G32:  mul_mat_vec_tqkp_kernel<2,32 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ2P_G64:  mul_mat_vec_tqkp_kernel<2,64 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ2P_G128: mul_mat_vec_tqkp_kernel<2,128><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ3P_G32:  mul_mat_vec_tqkp_kernel<3,32 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ3P_G64:  mul_mat_vec_tqkp_kernel<3,64 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ3P_G128: mul_mat_vec_tqkp_kernel<3,128><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ4P_G32:  mul_mat_vec_tqkp_kernel<4,32 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ4P_G64:  mul_mat_vec_tqkp_kernel<4,64 ><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ4P_G128: mul_mat_vec_tqkp_kernel<4,128><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ1P_Q4K6: mul_mat_vec_tqkp_q4k6_kernel<1><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ2P_Q4K6: mul_mat_vec_tqkp_q4k6_kernel<2><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        case GGML_TYPE_TQ3P_Q4K6: mul_mat_vec_tqkp_q4k6_kernel<3><<<grid,128,0,stream>>>(src0->data,x,d,ncols,nrows); break;
++        default: GGML_ABORT("mul_mat_vec_tqkp: unsupported type");
++    }
++}
++
++bool ggml_cuda_is_tqkp_fused(enum ggml_type t) {
++    switch (t) {
++        case GGML_TYPE_TQ1P: case GGML_TYPE_TQ2P: case GGML_TYPE_TQ3P: case GGML_TYPE_TQ4P:
++        case GGML_TYPE_TQ1P_G32: case GGML_TYPE_TQ1P_G64: case GGML_TYPE_TQ1P_G128:
++        case GGML_TYPE_TQ2P_G32: case GGML_TYPE_TQ2P_G64: case GGML_TYPE_TQ2P_G128:
++        case GGML_TYPE_TQ3P_G32: case GGML_TYPE_TQ3P_G64: case GGML_TYPE_TQ3P_G128:
++        case GGML_TYPE_TQ4P_G32: case GGML_TYPE_TQ4P_G64: case GGML_TYPE_TQ4P_G128:
++        case GGML_TYPE_TQ1P_Q4K6: case GGML_TYPE_TQ2P_Q4K6: case GGML_TYPE_TQ3P_Q4K6:
++            return true;
++        default: return false;
++    }
++}
++
+ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
+     switch (type) {
+         case GGML_TYPE_Q1_0:
+@@ -736,6 +992,44 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
+             return dequantize_row_q5_K_cuda;
+         case GGML_TYPE_Q6_K:
+             return dequantize_row_q6_K_cuda;
++        case GGML_TYPE_TQ1P:
++            return dequantize_row_tq1p_cuda;
++        case GGML_TYPE_TQ2P:
++            return dequantize_row_tq2p_cuda;
++        case GGML_TYPE_TQ3P:
++            return dequantize_row_tq3p_cuda;
++        case GGML_TYPE_TQ4P:
++            return dequantize_row_tq4p_cuda;
++        case GGML_TYPE_TQ1P_G32:
++            return dequantize_row_tq1p_g32_cuda;
++        case GGML_TYPE_TQ1P_G64:
++            return dequantize_row_tq1p_g64_cuda;
++        case GGML_TYPE_TQ1P_G128:
++            return dequantize_row_tq1p_g128_cuda;
++        case GGML_TYPE_TQ2P_G32:
++            return dequantize_row_tq2p_g32_cuda;
++        case GGML_TYPE_TQ2P_G64:
++            return dequantize_row_tq2p_g64_cuda;
++        case GGML_TYPE_TQ2P_G128:
++            return dequantize_row_tq2p_g128_cuda;
++        case GGML_TYPE_TQ3P_G32:
++            return dequantize_row_tq3p_g32_cuda;
++        case GGML_TYPE_TQ3P_G64:
++            return dequantize_row_tq3p_g64_cuda;
++        case GGML_TYPE_TQ3P_G128:
++            return dequantize_row_tq3p_g128_cuda;
++        case GGML_TYPE_TQ4P_G32:
++            return dequantize_row_tq4p_g32_cuda;
++        case GGML_TYPE_TQ4P_G64:
++            return dequantize_row_tq4p_g64_cuda;
++        case GGML_TYPE_TQ4P_G128:
++            return dequantize_row_tq4p_g128_cuda;
++        case GGML_TYPE_TQ1P_Q4K6:
++            return dequantize_row_tq1p_q4k6_cuda;
++        case GGML_TYPE_TQ2P_Q4K6:
++            return dequantize_row_tq2p_q4k6_cuda;
++        case GGML_TYPE_TQ3P_Q4K6:
++            return dequantize_row_tq3p_q4k6_cuda;
+         case GGML_TYPE_IQ2_XXS:
+             return dequantize_row_iq2_xxs_cuda;
+         case GGML_TYPE_IQ2_XS:
+@@ -791,6 +1085,44 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
+             return dequantize_row_q5_K_cuda;
+         case GGML_TYPE_Q6_K:
+             return dequantize_row_q6_K_cuda;
++        case GGML_TYPE_TQ1P:
++            return dequantize_row_tq1p_cuda;
++        case GGML_TYPE_TQ2P:
++            return dequantize_row_tq2p_cuda;
++        case GGML_TYPE_TQ3P:
++            return dequantize_row_tq3p_cuda;
++        case GGML_TYPE_TQ4P:
++            return dequantize_row_tq4p_cuda;
++        case GGML_TYPE_TQ1P_G32:
++            return dequantize_row_tq1p_g32_cuda;
++        case GGML_TYPE_TQ1P_G64:
++            return dequantize_row_tq1p_g64_cuda;
++        case GGML_TYPE_TQ1P_G128:
++            return dequantize_row_tq1p_g128_cuda;
++        case GGML_TYPE_TQ2P_G32:
++            return dequantize_row_tq2p_g32_cuda;
++        case GGML_TYPE_TQ2P_G64:
++            return dequantize_row_tq2p_g64_cuda;
++        case GGML_TYPE_TQ2P_G128:
++            return dequantize_row_tq2p_g128_cuda;
++        case GGML_TYPE_TQ3P_G32:
++            return dequantize_row_tq3p_g32_cuda;
++        case GGML_TYPE_TQ3P_G64:
++            return dequantize_row_tq3p_g64_cuda;
++        case GGML_TYPE_TQ3P_G128:
++            return dequantize_row_tq3p_g128_cuda;
++        case GGML_TYPE_TQ4P_G32:
++            return dequantize_row_tq4p_g32_cuda;
++        case GGML_TYPE_TQ4P_G64:
++            return dequantize_row_tq4p_g64_cuda;
++        case GGML_TYPE_TQ4P_G128:
++            return dequantize_row_tq4p_g128_cuda;
++        case GGML_TYPE_TQ1P_Q4K6:
++            return dequantize_row_tq1p_q4k6_cuda;
++        case GGML_TYPE_TQ2P_Q4K6:
++            return dequantize_row_tq2p_q4k6_cuda;
++        case GGML_TYPE_TQ3P_Q4K6:
++            return dequantize_row_tq3p_q4k6_cuda;
+         case GGML_TYPE_IQ2_XXS:
+             return dequantize_row_iq2_xxs_cuda;
+         case GGML_TYPE_IQ2_XS:
+diff --git a/src/ggml-cuda/convert.cuh b/src/ggml-cuda/convert.cuh
+index f5d37c7..707fcc8 100644
+--- a/src/ggml-cuda/convert.cuh
++++ b/src/ggml-cuda/convert.cuh
+@@ -64,3 +64,8 @@ template<typename dst_t, typename src_t>
+         return float(x);
+     }
+ }
++
++// sasori fused ternary mat-vec (decode): reads packed TQ{K}P weights directly (f16-scale families).
++void ggml_cuda_mul_mat_vec_tqkp(ggml_backend_cuda_context & ctx,
++        const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);
++bool ggml_cuda_is_tqkp_fused(enum ggml_type t);
+diff --git a/src/ggml-cuda/ggml-cuda.cu b/src/ggml-cuda/ggml-cuda.cu
+index f5293ad..0509752 100644
+--- a/src/ggml-cuda/ggml-cuda.cu
++++ b/src/ggml-cuda/ggml-cuda.cu
+@@ -2500,6 +2500,20 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_f(const ggml_tensor * tensor) {
+ }
+ 
+ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) {
++    // sasori: ternary trit-plane types have no fused mmvq kernel -> never fuse (would GGML_ABORT);
++    // they take the dequant -> cuBLAS path via ggml_get_to_fp16_cuda.
++    switch (tensor->src[0]->type) {
++        case GGML_TYPE_TQ1P:  case GGML_TYPE_TQ2P:  case GGML_TYPE_TQ3P:
++        case GGML_TYPE_TQ4P:  case GGML_TYPE_TQ1P_G32:  case GGML_TYPE_TQ1P_G64:
++        case GGML_TYPE_TQ1P_G128:  case GGML_TYPE_TQ2P_G32:  case GGML_TYPE_TQ2P_G64:
++        case GGML_TYPE_TQ2P_G128:  case GGML_TYPE_TQ3P_G32:  case GGML_TYPE_TQ3P_G64:
++        case GGML_TYPE_TQ3P_G128:  case GGML_TYPE_TQ4P_G32:  case GGML_TYPE_TQ4P_G64:
++        case GGML_TYPE_TQ4P_G128:  case GGML_TYPE_TQ1P_Q4K6:  case GGML_TYPE_TQ2P_Q4K6:
++        case GGML_TYPE_TQ3P_Q4K6:
++            return false;
++        default:
++            break;
++    }
+     ggml_tensor *       src0 = tensor->src[0];
+     ggml_tensor *       src1 = tensor->src[1];
+     const ggml_tensor * dst  = tensor;
+@@ -2604,6 +2618,17 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
+         return;
+     }
+ 
++    // sasori: fused ternary mat-vec for the DECODE case (ne11==1) — reads packed TQ weights directly
++    // instead of the dequant->cuBLAS path (much less weight bandwidth). f16-scale families only.
++    // Toggle off with SASORI_NO_FUSED_MMVEC=1 (A/B correctness vs the dequant->cuBLAS reference).
++    static const bool sasori_fused_off = getenv("SASORI_NO_FUSED_MMVEC") != nullptr;
++    if (!sasori_fused_off && !split && ggml_cuda_is_tqkp_fused(src0->type) &&
++        src1->type == GGML_TYPE_F32 && src1->ne[1] == 1 && src1->ne[2] == 1 && src1->ne[3] == 1 &&
++        ggml_is_contiguous(src1) && ggml_is_contiguous(src0)) {
++        ggml_cuda_mul_mat_vec_tqkp(ctx, src0, src1, dst);
++        return;
++    }
++
+     if (!split && use_mul_mat_vec_f) {
+         // the custom F16 vector kernel can be used over batched cuBLAS GEMM
+         // but this is only faster for GPUs without tensor cores or with a thin src0 matrix (particularly KQV in attention)
+@@ -5166,6 +5191,25 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
+                     case GGML_TYPE_Q5_K:
+                     case GGML_TYPE_Q6_K:
+                     case GGML_TYPE_Q8_K:
++                    case GGML_TYPE_TQ1P:
++                    case GGML_TYPE_TQ2P:
++                    case GGML_TYPE_TQ3P:
++                    case GGML_TYPE_TQ4P:
++                    case GGML_TYPE_TQ1P_G32:
++                    case GGML_TYPE_TQ1P_G64:
++                    case GGML_TYPE_TQ1P_G128:
++                    case GGML_TYPE_TQ2P_G32:
++                    case GGML_TYPE_TQ2P_G64:
++                    case GGML_TYPE_TQ2P_G128:
++                    case GGML_TYPE_TQ3P_G32:
++                    case GGML_TYPE_TQ3P_G64:
++                    case GGML_TYPE_TQ3P_G128:
++                    case GGML_TYPE_TQ4P_G32:
++                    case GGML_TYPE_TQ4P_G64:
++                    case GGML_TYPE_TQ4P_G128:
++                    case GGML_TYPE_TQ1P_Q4K6:
++                    case GGML_TYPE_TQ2P_Q4K6:
++                    case GGML_TYPE_TQ3P_Q4K6:
+                     case GGML_TYPE_IQ1_M:
+                     case GGML_TYPE_IQ1_S:
+                     case GGML_TYPE_IQ2_S:
+diff --git a/src/ggml-cuda/mmvq.cu b/src/ggml-cuda/mmvq.cu
+index 4b04265..555d74b 100644
+--- a/src/ggml-cuda/mmvq.cu
++++ b/src/ggml-cuda/mmvq.cu
+@@ -278,6 +278,21 @@ int get_mmvq_mmid_max_batch(ggml_type type, int cc) {
+ }
+ 
+ bool ggml_cuda_should_use_mmvq(enum ggml_type type, int cc, int64_t ne11) {
++    // sasori ternary trit-plane types have no mmvq (or mmq) kernel yet: force the
++    // dequant -> cuBLAS GEMM path (they are wired into ggml_get_to_fp16_cuda). Without
++    // this, use_mul_mat_vec_q would be true and mul_mat_vec_q_switch_type would GGML_ABORT.
++    switch (type) {
++        case GGML_TYPE_TQ1P:  case GGML_TYPE_TQ2P:  case GGML_TYPE_TQ3P:
++        case GGML_TYPE_TQ4P:  case GGML_TYPE_TQ1P_G32:  case GGML_TYPE_TQ1P_G64:
++        case GGML_TYPE_TQ1P_G128:  case GGML_TYPE_TQ2P_G32:  case GGML_TYPE_TQ2P_G64:
++        case GGML_TYPE_TQ2P_G128:  case GGML_TYPE_TQ3P_G32:  case GGML_TYPE_TQ3P_G64:
++        case GGML_TYPE_TQ3P_G128:  case GGML_TYPE_TQ4P_G32:  case GGML_TYPE_TQ4P_G64:
++        case GGML_TYPE_TQ4P_G128:  case GGML_TYPE_TQ1P_Q4K6:  case GGML_TYPE_TQ2P_Q4K6:
++        case GGML_TYPE_TQ3P_Q4K6:
++            return false;
++        default:
++            break;
++    }
+     if (GGML_CUDA_CC_IS_CDNA(cc)) {
+         if (GGML_CUDA_CC_IS_CDNA1(cc)) {
+             switch (type) {

text_encoder/Qwen3-8B-TQ3P.gguf

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9cb46cba22a6cf01989ddf8f8d1c67eabb3e7c32d9c0a872fc810e883c371274
+size 7868626528