# 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.