yaminerl/escha-amd-port

Escha Qwen3.8-27B W2 on AMD (RDNA2)

This repo contains a patch and a verification script, not weights. The GGUFs live at aj9o9/Qwen3.8-27B-Escha-W2-GGUF.

I ported aj9o9's llama.cpp-escha fork to AMD. The weights are EschaLabs'. The kernel and the GGML_OP_ESCHA_MUL_MAT op are aj9o9's. This is not a requant and not a new quantization. What I changed is four edits so their CUDA kernel builds and runs under HIP on a Radeon RX 6800 XT.

The model card says CUDA and tensor cores are required. Tensor cores are required for their prefill path, not for the model. RDNA2 has none and it runs anyway, on the fp32 fallback that fork already carries.

Upstream GGUF: aj9o9/Qwen3.8-27B-Escha-W2-GGUF Upstream fork: Ajay9o9/llama.cpp-escha, branch escha-w2-dense

Stock llama.cpp cannot load these files. Neither can the unpatched fork on an AMD card. You need both the fork and the patch in this repo.

What I changed

All of it is in ggml/src/ggml-cuda/escha-moe.cu. 22 lines added, 1 removed. Every change is inside #if defined(GGML_USE_HIP), so the NVIDIA path is byte-for-byte what it was. No build-system change: ggml-hip/CMakeLists.txt already globs ../ggml-cuda/*.cu.

The codebook is written in inline PTX. Two sites, both in the hot path:

asm("lop3.b32 %0, %1, %2, %3, 0x6a;" : "=r"(x) : "r"(x), "n"(0x8fff8fffu), "n"(0x3b603b60u));

lop3 is a three-input logic op selected by an 8-bit truth table. Evaluate the expression on the canonical inputs ta=0xF0, tb=0xCC, tc=0xAA to read the table back: (0xF0 & 0xCC) ^ 0xAA == 0x6A, so immLut 0x6a is (a & b) ^ c and the portable spelling is exact, not an approximation:

x = (x & 0x8fff8fffu) ^ 0x3b603b60u;

The fork's own CPU implementation spells it that way already, at ggml/src/ggml-cpu/ops.cpp:11265. I derived it from the truth table first and found the CPU line afterwards, so the two are independent.

#include <cuda_pipeline.h> has no HIP equivalent. cp.async is only used inside the tensor-core kernel, which is compiled out on AMD, so the include is guarded and nothing else moves.

The one that would have shipped broken. The host side picked the tensor-core kernel with this:

const bool use_mma = !gen
                  && ggml_cuda_info().devices[ctx.device].cc >= GGML_CUDA_CC_TURING
                  && ...

llama.cpp defines TURING_MMA_AVAILABLE as !defined(GGML_USE_HIP), so on AMD that kernel's body compiles to NO_DEVICE_CODE. The cc test does not catch it. AMD cc values carry GGML_CUDA_CC_OFFSET_AMD, which is 0x1000000, so they compare above every NVIDIA generation and the predicate is silently true. It compiles clean, launches, and computes nothing. The fix is a host-side mirror of the device macro:

#if defined(GGML_USE_HIP)
#define ESCHA_MMA_SUPPORTED false
#else
#define ESCHA_MMA_SUPPORTED true
#endif

This is the part worth reading if you are porting some other CUDA fork. A cc comparison is not a stand-in for a feature macro once HIP is in the build.

Wave size needed no work. RDNA2 runs wave32 and llama.cpp reports Wave Size: 32, which matches the CUDA warp the kernel assumes. The hardcoded 32s in that file are all inside the tensor-core kernel, which is off here.

Correctness

The fork ships a CPU implementation of the op, so it is its own reference. I did not use generated text as the gate, because a wrong kernel can still read fluently, and I could not use the fork's own unit tests, because test-escha-mul-mat.cpp wants an escha-mul-mat-cases.gguf fixture that was never published.

What I did instead: dump every graph tensor on both backends for the same prompt with llama-eval-callback, match by name and occurrence, diff the sampled values and the per-tensor sum.

The rest of the graph is the control. Every op that is not ESCHA_MUL_MAT is stock llama.cpp going through the same HIP backend on the same card, so whatever CPU-vs-GPU deviation those show is this build's numerical envelope: reduction order, fp16 rounding in flash attention, and so on. That turns "is 0.0005 close enough" into a comparison instead of a guess.

	tensors	max abs Δ value	max rel Δ sum
ESCHA_MUL_MAT, ported	325	0.0005	0.014
stock ops, worst of 19	1919	0.0046	0.103

The ported op deviates about 9x less than ops that are already known good. 325 tensors is every escha projection across all 64 layers. No non-finite values. compare-dumps.py in this repo does the matching and prints that table.

llama-eval-callback -m model.gguf -ngl 99 -p "Hello world" > gpu.log
llama-eval-callback -m model.gguf -ngl 0  -p "Hello world" > cpu.log
python3 compare-dumps.py gpu.log cpu.log

Use a prompt, not a long generation. The CPU op shards only over batch rows, so at batch 1 it is effectively single-threaded and a 27B decode will not finish in useful time.

What this establishes is that the HIP kernel computes what the CPU kernel computes. It says nothing about quality against EschaLabs' runtime. aj9o9 did that work on the CUDA side, at 99.4% top-1 and JS divergence 0.0000 bits, and none of my four edits touch the math on a path that agreement covered. For what quality looks like on this build, on plain text, see Quality below.

Speed

One RX 6800 XT, stock clocks, single stream, batch 1, full offload. Ryzen 5 5600X, CachyOS, ROCm 7.2.4. llama-bench, -r 2:

	RX 6800 XT	RTX 3090, 250 W, upstream
pp128	120.3 ± 17.6	—
pp512	132.4 ± 3.9	700.4
tg64 / tg128	18.9 ± 0.4	24.03

Decode lands at 79% of the 3090. That tracks: decode is memory-sensitive and this card has 512 GB/s against the 3090's 936.

Prefill is 19% of it, and that is the tensor cores, not the port. RDNA2 has no matrix units at all. WMMA starts at RDNA3, gfx1100. Upstream measured 212.4 tok/s before they added the tensor-core path and 700.4 after, so 132 on the fp32 fallback with a slower card is roughly where it should sit.

There is no ESCHA_NO_MMA=1 decision to make here. On AMD the fallback is forced at compile time by edit three, because there is nothing to fall back from.

Quality

Escha is a quantization method on the same base model as every other quant here, Qwen/Qwen3.8-27B, so a KL-divergence audit against a shared reference is a fair comparison, not an apples-to-oranges one.

Methodology follows a benchmark of the same model family: llama-perplexity --kl-divergence, wikitext-2, 32 chunks, 16,384 tokens, against a shared reference model's logits. First pass used Qwen3.8-27B-UD-Q5_K_XL, the same reference that piece uses, so I could check the pipeline against its published numbers before trusting it on Escha: that piece measured Q3_K_XL at 92.62%/0.2955 (top-1/99% KLD) and IQ4_XS at 94.41%/0.1993 against Q5_K_XL, on different hardware. Mine came out 93.08%/0.2855 and 94.19%/0.2006 — inside noise of theirs.

Q5_K_XL is still one quantization step away from the real thing, though, so the table below uses Qwen3.8-27B-Q8_0 as reference instead — 8.5 bpw, the closest to lossless llama.cpp ships for this model short of raw BF16. Q8_0 and Q5_K_XL agree with each other on perplexity to four significant figures (5.5969 vs 5.5958), so this is about as close to a genuine BF16 baseline as this card can produce without downloading 55 GB of safetensors and running them uncompressed.

quant	bpw	top-1 vs Q8_0	perplexity	PPL penalty	99% KLD
Escha-W2	2.47	84.60 ± 0.40%	6.1041	+9.10%	1.3325
UD-Q3_K_XL	3.94	93.36 ± 0.28%	5.6851	+1.61%	0.2575
IQ4_XS	4.60	94.29 ± 0.26%	5.6148	+0.35%	0.1827
UD-Q5_K_XL	5.92	97.70 ± 0.17%	5.5958	+0.01%	0.0329
Q8_0	8.50	reference	5.5969	—	—

Swapping the reference barely moved Escha's number: 84.60% here against 84.78% against Q5_K_XL. If the earlier result had been an artifact of a weak reference, closing most of the remaining bpw gap to BF16 would have closed some of that gap too. It didn't, which means the divergence is coming from Escha, not from what I measured it against.

Escha's own card claims 99.4% top-1 against BF16. I still can't close that gap from here — different corpus, different runtime (SGLang vs llama.cpp), and Q8_0 is not BF16 — and I am not saying their number is wrong. What I can say: on plain natural-language text, on this build, against the closest thing to BF16 this card can hold, Escha disagrees with the reference on roughly one token in six. Q3_K_XL and IQ4_XS disagree on roughly one in fifteen and one in seventeen, in files 30% and 52% larger than Escha's.

llama-perplexity -m Qwen3.8-27B-Q8_0.gguf -f wikitext-2-raw/wiki.test.raw \
  --kl-divergence-base kld-ref-q8.dat --chunks 32 -ngl 28 -fa on

llama-perplexity -m Escha-Qwen3.8-27B-W2-Q8E.gguf -f wikitext-2-raw/wiki.test.raw \
  --kl-divergence --kl-divergence-base kld-ref-q8.dat --chunks 32 -ngl 99 -fa on

Get wiki.test.raw from scripts/get-wikitext-2.sh in either llama.cpp tree. Qwen3.8-27B-Q8_0.gguf doesn't fit in 16 GB either; -ngl 28 was what fit here with room left for the compute buffer. Generating the reference logits is a one-off and doesn't need to be fast — this pass, more than half on 6 CPU cores, still finished in under 10 minutes.

Practical read: Escha buys VRAM — 10.3 GB resident against Q3_K_XL's 13.4 GB file — at a real quality cost, not the near-lossless one the card advertises. If VRAM headroom isn't what's binding for you, Q3_K_XL or IQ4_XS are the better model at this precision, on this card.

Is the gap in the conversion, not the method? EschaLabs' safetensors are already the quantized checkpoint — escha_code / escha_rin / escha_rout sidecars per tensor, not BF16 — so turning that into a GGUF isn't quantization, it's a container repack, and that's checkable directly against aj9o9's GGUF with no inference involved: pull the same tensor out of both files and diff it.

I did that for every escha-quantized tensor in the model, not a sample. diff-escha-tensors.py in this repo does the matching:

python3 diff-escha-tensors.py /path/to/EschaLabs-Qwen3.8-27B-Escha-W2 Escha-Qwen3.8-27B-W2-Q8E.gguf

family	layers	tensors	code identical	scale-fold verified
mlp.down/gate/up_proj	64	192	192/192	192/192
self_attn.q/k/v/o_proj	16	64	64/64	64/64
linear_attn.in_proj_qkv/z, out_proj	48	144	0/144	0/144 (positional)

256 of 400 tensors — the MLP and standard-attention weights, and the largest tensors in the model by parameter count — are bit-identical on escha_code between EschaLabs' original and aj9o9's GGUF, with rin/rout matching original × scale to fp16 rounding, every time. The remaining 144, the linear-attention (SSM) projections, never match position for position, but all 144 have an exact codebook-histogram match: same 65,536 values, same counts, just reordered. A conversion bug does not reproduce a perfect histogram by chance 144 times running; that pattern is a deliberate memory-layout repack for that codec variant, not corruption, though I did not decode the permutation itself.

So the conversion checks out, on every tensor family I could check it on. The gap above isn't the AMD port (Correctness, above) and isn't aj9o9's GGUF (this). What's left: Escha costing more on plain text than the card states, or EschaLabs measuring on a distribution wikitext-2 doesn't resemble. Both stay open.

Build

git clone -b escha-w2-dense https://github.com/Ajay9o9/llama.cpp-escha.git
cd llama.cpp-escha
git apply /path/to/escha-hip-gfx1030.patch

HIPCXX=/opt/rocm/llvm/bin/clang++ HIP_PATH=/opt/rocm \
cmake -S . -B build-hip -DGGML_HIP=ON -DAMDGPU_TARGETS=gfx1030 \
      -DCMAKE_BUILD_TYPE=Release -DLLAMA_CURL=OFF
cmake --build build-hip -j 6

gfx1030 is a 6800 XT / 6900 XT. Change it for your card. Builds clean, one -Wcuda-compat warning about a #pragma unroll in parentheses, which is upstream's and harmless.

gfx1030 is natively supported by current ROCm. You do not need HSA_OVERRIDE_GFX_VERSION.

Download

hf download aj9o9/Qwen3.8-27B-Escha-W2-GGUF Escha-Qwen3.8-27B-W2-Q8E.gguf --local-dir .

Take the Q8_0 build, for upstream's reasons. I only tested that one.

Run it

./build-hip/bin/llama-server \
  -m Escha-Qwen3.8-27B-W2-Q8E.gguf \
  --host 127.0.0.1 --port 8080 \
  -ngl 99 -fa on --jinja \
  -np 1 -t 10 \
  -c 65536 \
  --cache-type-k q8_0 --cache-type-v q8_0

It is a reasoning model. The chat template uses <think> tags and takes a reasoning_effort kwarg, so --jinja --reasoning-format deepseek gets you a separate reasoning_content field. Give it room to think or you get an empty answer.

VRAM

Measured with rocm-smi on a 16 GiB card, weights plus q8_0 KV plus compute buffer, model resident:

context	VRAM
8192	10.20 GiB
65536	11.96 GiB
65536, -nkvo	9.83 GiB

Those three rows are measured. I did not test past 65536. Upstream measured 18.69 GiB at 262144 on a 3090, which does not fit here, so somewhere above 64k this card runs out. The KV cache is cheap for this architecture, about 34 KB per token at q8_0, because only 16 of the 64 layers are full attention.

-nkvo pushes the KV cache to host RAM. This model does not need it at 64k, and it costs you speed, but if you are sharing the card it buys back 2.1 GiB.

What this is not

• Not a requant. The 2-bit payload is decoded in the kernel. I did not touch the format.
• Not my kernel. aj9o9 wrote it. I made it compile and dispatch correctly under HIP.
• Not a full quality evaluation. Quality above is one axis, KL-divergence against a Q8_0 reference on wikitext-2. No MMLU, no long-context eval, no comparison against SGLang or BF16 directly. Correctness and Quality are two different claims, kernel agreement and output quality — read both, not just one.
• Not tested on any card but a 6800 XT. gfx1030 only. Other RDNA2 parts should work unchanged. RDNA3 and RDNA4 have WMMA and AMD_WMMA_AVAILABLE, so a real prefill path is plausible there, but it would mean mapping the fork's mma.cuh tile ops onto it and I have not tried. CDNA has MFMA and the same applies.
• Not the MoE 35B. Dense 27B only, on the escha-w2-dense branch. I did not touch the escha-w2 branch or GGML_OP_ESCHA_MOE, beyond keeping it compiling.
• Not the F16 build. I only downloaded and ran the Q8_0 file.
• Not upstream. Neither the fork nor this patch.
• Not fast at prefill. See above. If prefill throughput is what you need, this is the wrong card, not the wrong patch.

License

Apache-2.0 for the weights, MIT for the llama.cpp patch, following the projects they belong to.

Weights: EschaLabs/Qwen3.8-27B-Escha-W2 GGUF and kernel: aj9o9/Qwen3.8-27B-Escha-W2-GGUF