hipEngine

hipEngine is a ROCm-native local LLM inference engine designed from the ground up for AMD RDNA GPUs (starting with gfx1100, gfx1151). It pairs a small purpose-built Python host with a complete suite of custom-tuned HIP kernels developed through 100+ iterations of profiling and tuning.

hipEngine has lightweight dependencies with no PyTorch required for fully supported GPUs and models.

Core principles

• HIP-first, not CUDA-ported. Kernels directly target AMD hardware like gfx1100/RDNA3 with wave32, vec8 FMA, and the actual cache hierarchy.
• Torch-free runtime. import torch is not on the hot path. The runtime owns a thin hipengine.Tensor over raw HIP/CUDA device pointers and drives hipblasLt, hipGraph, AOTriton, and JIT builds through ctypes. Torch appears only as an optional dlpack bridge behind the hipengine[torch] extra (~125 MiB install including the vendored AOTriton subset vs ~2 GiB with torch).
• Multi-backend from day one. Kernels live under kernels/hip_gfx1100/, kernels/hip_gfx1151/, kernels/cuda_sm86/, kernels/cpu_reference/ as peer trees.
• Four-axis plugin registry. Kernels are keyed by (backend, layer, quant, variant). Models, quant schemes, and layers are plugins. No if backend == "..." or if quant == "..." branches in dispatch / engine / model code.
• Fused + unfused coexist. Every fused composite (rmsnorm+rotate, gate_combine_residual, …) has a numerically-equivalent unfused chain registered under its primitives, used as both fallback and correctness baseline.
• Evidence-backed performance. Every performance claim ships with model + quant + workload shape + hardware + exact command + correctness gate (KL ≤ 0.05, top-1 ≥ 90% vs kernels/cpu_reference/). See docs/BENCHMARK.md and benchmarks/README.md.

Status

v0.3.0 alpha. The runtime hot path is torch-free by construction, and the first two 35B-class model-loading surfaces are available on gfx1100 and gfx1151: shisa-ai/Qwen3.6-35B-A3B-PARO-packed (19.07 GiB, 4.68 bpw) in packed ParoQuant format, plus Qwen3.6 GGUF Q4_K_M / Q4_K_S files through the resident GGUF path. Older benchmark artifacts may still show the historical Qwen3.6-35B-A3B-PARO-full4096-e5-packed name or local MTP-BF16 assembly path; those rows use the same packed PARO architecture and remain the evidence for the numbers below.

• Model-aware backend="auto" / quant="auto" defaults select the registered PARO or GGUF route without environment-variable setup. Direct generation now supports exact token-id prompts, detailed outputs, logprobs, structured finish details, and backend execution telemetry.
• PARO and GGUF support ordinary sampling controls including top-k/min-p, penalties, logit bias, suppression, deterministic seeds, EOS/min-token policy, token stops, and multi-token stops. Covered PARO shapes use a native GPU sampler; unsupported shapes use an explicit host fallback.
• The OpenAI-compatible server includes capability/readiness discovery, token and context diagnostics, exact usage accounting, request batching, deadlines, cancellation, opt-in Prometheus metrics, and detailed streaming metadata.
• Local-agent support includes OpenAI-style tools, Qwen thinking controls, structured-output result validation, deterministic continuation handles, and app-local transcript sessions with fork, rollback, snapshot, and overflow policies.
• Qwen3.6 GGUF models with NextN tensors expose detailed MTP generation and a guarded explicit non-streaming server route. Dense PARO DFlash and the shared speculative proposal/verify/commit infrastructure are available as retained runtime and benchmark paths.
• PARO BF16 KV has retained W7900 evidence through 128K under the 24 GiB gate. The 256K INT8 layout also fits physically, but fails the required Qwen3.6 long-rollout quality gate and is not a supported or usable route. Current capacity, throughput, speculative-decode, and concurrency evidence is reported below with separate gfx1100/gfx1151 provenance and correctness gates.

This remains an alpha, single-GPU release. Production PARO native c>1 decode is disabled pending independent-c1 correctness, app-local sessions do not reuse resident KV, structured outputs are not grammar-constrained decoding, and the server MTP route is explicit-only. See the API limitations and concurrency status for the exact boundaries.

Hardware targets

backend="auto" is the public API/server default. It maps exact gfx1100 and gfx1151 detections to the matching HIP backend; unknown ROCm targets warn and select cpu_reference where a CPU implementation exists. Users on nearby targets such as gfx1101/gfx1102 can force a backend with backend="hip_gfx1100", --backend hip_gfx1100, or HIPENGINE_BACKEND=hip_gfx1100 after validating correctness/performance.

Wave32 is the default for hip_gfx1100 device code; wave64 is treated as an isolated experiment with its own gates (see docs/PLAN.md).

Memory Usage

The clean 2026-07-12 W7900 run measured hipEngine 8116c453 (rebased-equivalent reachable 8708304f; runtime and benchmark code identical) against the current Qwen3.6 packed PARO model under the 24 GiB portability gate. Both physical capacity/layout rows pass, but the required Qwen3.6 long-rollout quality gate rejects INT8 KV. Accordingly, 256K INT8 is reported as allocation capacity—not as a supported or usable route.

The INT8 layout retains 2,686,976,000 payload bytes plus 20,992,000 FP16 scale bytes across ten full-attention layers and no BF16 K/V shadow. Its matched 128K/128 BF16 comparison diverges at generated index 2, with mean/max KL 3.7646/10.0796 and 3.88% top-1 agreement. FP32 scales also reject. Memory was measured once per row; the reported timing fields are diagnostic only.

See the capacity summary, FP16-scale quality gate, and FP32-scale follow-up.

llama.cpp configuration note

The repository has no compact artifact or source revision for the former llama.cpp Q8_0 memory tables, so those numbers are not toplines. The tested configuration was:

--flash-attn on -ctk q8_0 -ctv q8_0 -c 262144 -b 128 -ub 128

A replacement capacity table must record the GGUF fingerprint, llama.cpp commit/build, GPU, full command, and whole-card sampling artifact.

Model Performance

gfx1100 (Radeon RX 7900 XTX / Radeon Pro W7900)

Status: retained. This clean 2026-07-12 refresh measured hipEngine 8116c453 (rebased-equivalent reachable 8708304f; runtime and benchmark code identical), TheRock HIP 7.15, right-sized resident sessions, production graph decode, two discarded plus five measured hipEngine runs, and five llama.cpp samples per phase. The W7900-local GGUF oracle passes external tokens and byte-exact hidden/Conv/GDN/KV state. All six rows pass clean provenance, stable finite outputs, exact Q4_K_M identity, corrected W7900 VRAM scope, and sample-variance gates. PARO remains W4 PARO/BF16 KV; the other columns use Q4_K_M with BF16/F16 KV, so bold values are descriptive rather than same-quant wins.

Prefill tok/s

Decode tok/s

Peak memory GiB

W7900 row sources: accepted summary, hipEngine PARO, hipEngine GGUF, llama.cpp HIP, llama.cpp Vulkan, and W7900 correctness oracle.

gfx1151 (AMD Ryzen AI MAX+ 395 / Radeon 8060S)

Thanks to Framework for sending a dedicated Framework Desktop Strix Halo motherboard for this profiling and tuning work.

Status: retained. GGUF and llama.cpp are the clean 2026-07-11 matched refresh. PARO 512/1K are the clean 2026-07-12 exact recovery at 9944e481; 4K and 32K-128K are the clean scoped AOTriton queue-isolation refresh at 01e2cec5, all with TheRock HIP 7.15 and TuneD accelerator-performance. hipEngine uses two discarded warmups plus five measured repetitions per right-sized resident shape; llama.cpp uses one internal warmup plus five samples per split phase. The linked artifacts pass their clean provenance, output/state, variance, model/build/device, and memory-scope gates. Bold marks the best raw value per row, but PARO is W4 PARO rather than Q4_K_M and memory scopes differ, so the emphasis is descriptive rather than a controlled same-quant/backend win.

Prefill tok/s

Decode tok/s

Peak memory GiB

The memory columns have different scopes: hipEngine reports tracked allocator high-water, while llama.cpp reports absolute whole-device amdgpu GTT used, sampled every 10 ms. Use them for within-column context growth, not small cross-column allocator comparisons. Row sources: PARO exact recovery, PARO 4K-128K AOTriton queue isolation, accepted July 11 matched summary, July 11 PARO reference, hipEngine GGUF, llama.cpp HIP, and llama.cpp Vulkan. Exact settings and gates are in the canonical benchmarks/README.md.

Current gfx1151 GGUF decode baselines

These are separate exact repeated-token SOL-G4/G5 controls. The model sweep above excludes graph capture from steady decode throughput; SOL-G5 charges one capture/instantiate and destroy to each 128-token window.

Artifacts: SOL-G4 eager audit and SOL-G5 production graph audit.

See benchmarks/README.md for the platform freshness index, exact settings, run commands, and evidence status.

Speculative decode (DFlash / MTP)

Every displayed route has its own same-protocol AR control. The exact/default and llama-compat columns are separate because only llama-compat shares the B2 natural24 structure used by the llama.cpp comparison.

GGUF MTP comparison, Radeon Pro W7900/gfx1100

The old 34.28-34.49 tok/s true-AR denominator was an eager-only benchmark path, not the fastest production no-MTP route. gfx1100 had no backend graph capability even though the state-bound implementation was already shared with gfx1151. A clean W7900 p512/d24 gate now passes all 24 hidden/GDN/KV/token transitions and moves capture-inclusive wall from 30.536 to 12.514 ms/token (2.4402x). The full natural24 suite matches every prior eager generated-token preview/tail and moves 34.28 -> 93.30 tok/s in the same MTP wrapper.

At the matched cross-engine boundary, hipEngine counts 240 complete post-prefill transitions including graph capture/instantiate/close; llama.cpp build 9648 requests 25 outputs and counts the 240 timed transitions inside predicted_ms. hipEngine is 93.30 versus 78.29 tok/s (+19.19%). BF16 versus F16 KV remains disclosed. llama.cpp stays an external diagnostic with performance_claim=false because its local instrumentation patchset is dirty but preserved.

Neither MTP route beats the corrected production AR control. Exact/default remains the semantic control; llama-compat remains explicit-only because direct partial commit is not serial-prefix-equivalent. The fixed-cycle exact and natural24 compatibility rows are different protocols and are not ranked against each other.

W7900 llama-compat full-suite gate against graph AR

All four categories and heldout lose to graph AR despite unchanged strong draft acceptance. This corrects the earlier false MTP-win conclusion without changing the compatibility semantics. Artifact: 2026-07-12-w7900-gfx1100-gguf-graph-ar-refresh.json.

GGUF MTP comparison, Radeon 8060S/gfx1151

The current exact/default B5 route no longer beats true AR after the correctness/state-lifecycle pass: 51.81 vs 54.14 tok/s (0.9571x). Its old 61.98 tok/s row is retained only as history. llama-compat remains a separate, explicit-only semantic contract and is not serial-prefix-equivalent.

The cross-engine rows use the canonical transition-matched timing contract: hipEngine uses complete cycle wall; llama.cpp requests 25 outputs and counts the 24 transitions inside predicted_ms. This removes llama.cpp's native one-untimed-token numerator advantage. hipEngine uses BF16 KV while llama.cpp uses F16 KV, which remains a model-execution difference even with matched timer boundaries. The captured llama.cpp source is dirty but fully preserved in the repository patchset; the binary hash is authoritative and performance_claim=false.

gfx1151 llama-compat full-suite gate

All four categories and the heldout split beat their true same-protocol AR controls. Train/heldout draft acceptance is 82.08% / 71.79%; the gap is kept visible rather than averaged away.

Dense PARO DFlash

Artifacts: W7900 GGUF MTP transfer, DFlash, gfx1151 exact MTP, and gfx1151 llama-compat MTP. The gfx1151 matched natural24 controls are exact/default B1-B5 and llama.cpp HIP B2. Historical hipEngine OpenAI MTP server rows are excluded. The current raw-ID route counts exact completion IDs across every choice and owns batch timing once. The corrected 2026-07-11 server matrix finds that compatibility MTP changes true-AR IDs even at c1, so it must remain explicit-only despite diagnostic c1/c2 speed gains; SOL-S1 routes automatic requests to exact/default AR while keeping the compatibility hook explicit-only. See the route-gate artifact and canonical benchmarks/README.md.

The clean gfx1151 PARO DFlash S4 profile is exact but not competitive: 9.68 versus 65.27 tok/s AR (0.148x) at B4/32 tokens. Branch-copy is faster but diverges at generated token 1, and fused target LM-head is 5.16% slower than unfused. See the compact profile and the canonical benchmark analysis.

Concurrency

The current publishable gfx1151 table is the exact PARO production-routing catalog below. c1 has a retained timing; c2-c8 use width-1 sessions because every native candidate fails the independent-c1 oracle. P2 proves that serial route through ragged c8-to-c1 EOS/cancel transitions and front/middle/tail sparse slots. Native batching remains closed until a general c>N algorithm passes the same token/state/KV gates. See docs/CONCURRENCY.md for the design history.

The linked records keep gfx1100 and gfx1151 separate because the model files, ROCm stacks, and comparison backends differ. Aggregate is total tok/s across the batch; per-sequence is tok/s seen by one request. See docs/VLLM_RDNA3.md for vLLM RDNA3 setup notes.

gfx1100 / W7900 decode tok/s vs concurrency (Qwen3.6 35B-A3B, 512/128)

Status: stale diagnostic. This is a median-of-3 scaling snapshot, not an apples-to-apples engine ranking. hipEngine uses PARO W4/BF16 KV, llama.cpp uses Vulkan Q4_K_M/f16 KV, and vLLM uses GPTQ Int4. hipEngine and llama.cpp report backend decode timing; vLLM reports OpenAI client wall throughput.

No eligible concurrency row; the mixed-quant, mixed-timing sweep remains linked below pending rerun.

Protocol: prompt 512, decode 128, 8 warmup decode tokens, median of 3. hipEngine c=1 uses the single-sequence graph-replay benchmark and c>1 uses the native batch benchmark. llama.cpp restarts its server for each concurrency and repetition with -np c -c 1024*c.

Source artifacts: hipEngine W7900, llama.cpp Vulkan W7900, vLLM local build W7900, full W7900 refresh summary, and vLLM RDNA3 notes.

gfx1151 / Radeon 8060S PARO exact shape catalog (2026-07-11, Qwen3.6 35B-A3B, 512/128)

Status: retained c1 performance, c1-c8 routing correctness, and production lifecycle safety. Clean a18ff7bc uses the same exact 512-token fixture at every width. c1 graph replay is retained; every c2-c8 native candidate fails independent-c1 equality at generated index 2 and is explicitly routed through width-1 sessions. Clean 6f1910c9 then passes ragged c8-to-c1 EOS/cancel lifecycle coverage without compacting physical slots.

Protocol: W4 PARO/BF16 KV, 40 layers, exact prompt-ID SHA-256 b162b2d0...2388, 8 warmup decode steps, 128 measured decode steps, and greedy sampling. c1 is a clean median of three (66.948/66.754/66.910 tok/s). Native c2-c8 diagnostic rates are withheld from the topline because all rows fail the 137-token oracle. The P2 gate uses prompt lengths 449 through 512 and matches all eight generated sequences, 30 linear-state families, and 10 live K/V families through EOS plus front/middle/tail sparse cancellation. Ragged prefill uses the correctness-first per_segment_ragged_exact fallback and makes no throughput claim. See the P1 compact catalog, P2 lifecycle artifact, and canonical run record.

gfx1151 / Radeon 8060S historical cross-engine concurrency (2026-06-15)

Status: stale diagnostic. hipEngine uses PARO W4/BF16 KV; llama.cpp uses Vulkan Q4_K_S/f16 KV. vLLM did not produce a healthy server. The summary lacks the measured hipEngine commit, and the then-used per-run device properties could report gfx1100 even though the run forced HIPENGINE_HIP_ARCH=gfx1151.

No eligible concurrency row; the performance_claim=false snapshot remains linked below pending rerun.

Protocol: prompt 512, decode 128, 8 warmup decode tokens, median of 3. Primitive c>1 attention/KV checks passed. The generated-token field used the older batch-shaped reference and is not independent-c1 evidence. Profiler, scaling, and provenance gates also did not pass.

Source artifacts: gfx1151 summary, hipEngine PARO, llama.cpp Vulkan, and vLLM blocked.

A 2026-06-13 RX 7900 XTX rerun reached c1/c2/c4 but c8 blocked with HIP OOM; see XTX partial. Replicate the W7900 hipEngine, llama.cpp Vulkan, and vLLM concurrency rows with:

scripts/run_w7900_readme_refresh.sh concurrency
scripts/run_w7900_readme_refresh.sh vllm

The exact settings and gfx1151 runner gap are recorded in benchmarks/README.md.

GGUF Support

As of v0.2.0, hipEngine includes resident Qwen3.6 GGUF support for Q4_K_M and Q4_K_S model files (with more formats planned). This is a major runtime path, not just a loader shim: GGUF has its own quant readers, bulk-prefill path, decode-repacked T16 layouts, and fast-path controls.

Current caveats:

• PARO models take ~22s to load on the W7900 test host in the current refresh; GGUF Q4_K_M currently takes about 74s because decode-repack happens on load. On-disk caching could reduce startup time later, but would require additional storage for repacked layouts.

• GGUF has higher base weight residency than packed PARO before KV cache is the deciding factor. The full-attention KV slope is the same 10-layer Qwen3.6 shape; the 24 GiB long-context gap is mostly the loaded-weight baseline. Packed PARO is ~19.07 GiB on disk, while the local GGUF tensor payloads are:

  GGUF tensor family	Q4_K_M GiB	Q4_K_M mix	Q4_K_S GiB	Q4_K_S mix	Q4_K_S - Q4_K_M
  Q4_K	11.531	54.7%	16.875	84.8%	+5.344
  Q5_K	6.531	31.0%	0.000	0.0%	-6.531
  Q8_0	1.932	9.2%	1.932	9.7%	+0.000
  Q6_K	1.004	4.8%	1.004	5.0%	+0.000
  F32/BF16 metadata	0.098	0.5%	0.098	0.5%	+0.000
  Total tensor payload	21.097	100.0%	19.909	100.0%	-1.188

  In other words, Q4_K_S saves ~1.19 GiB versus Q4_K_M by replacing the selected-MoE Q5_K expert-down payload with Q4_K; it still starts above packed PARO, and hipEngine's resident T16/pack8 decode layouts add their own allocator shape. On 24 GiB cards, current Q4_K_M BF16-KV support is a mid-context path unless a lower-memory KV/weight policy is explicitly enabled. The current clean gfx1151 p512/d128 census is 21.478 GiB owned/tracked: 20.461 GiB replacement weights, 0.503 GiB required raw token embedding, 0.097 GiB dense weights/metadata, and 0.417 GiB scratch/session buffers. It confirms the default T16 path replaces source layouts rather than retaining raw+packed copies; this short-context margin is not a 128K capacity claim.

• GGUF is close enough to PARO to share some high-level scheduling ideas, but in practice it needs substantial GGUF-only kernels and dispatch. The goal for future releases is to keep closing the remaining PARO/GGUF speed gap.

Architecture at a glance

┌─────────────────────────────────────────────────────────────────┐
│  USER API                                                       │
│  hipengine.LLM.generate()           library API                 │
│  hipengine serve                    OpenAI-compatible server    │
├─────────────────────────────────────────────────────────────────┤
│  LOADING (torch-free)                                           │
│  safetensors mmap + hipMemcpyAsync / HF config / jinja2 chat    │
│  templates / HF tokenizers (Rust)                               │
├─────────────────────────────────────────────────────────────────┤
│  DISPATCH                                                       │
│  Scheduler / Block Manager (KVPolicy) / Prefix Cache            │
│  Fusion Planner (chain → kernel plan, fused preferred)          │
│  Model / Quant / Layer plugins / Engine loop (hipGraph replay)  │
├─────────────────────────────────────────────────────────────────┤
│  CORE (torch-free primitives)                                   │
│  hipengine.Tensor / device / memory / stream / graph / blas     │
│  build (hipcc subprocess + ctypes.CDLL + .so cache)             │
├─────────────────────────────────────────────────────────────────┤
│  KERNELS (backend-keyed custom HIP implementations)             │
│  kernels/hip_gfx1100/  attention / linear_attn / moe / quant    │
│                        wmma / norm / rotary / fused             │
│  kernels/hip_gfx1151/  native target-arch peer backend          │
│  kernels/cuda_sm86/    (future)                                 │
│  kernels/cpu_reference/ correctness oracle, no GPU required     │
└─────────────────────────────────────────────────────────────────┘

Full layer diagram, plugin axes, KV cache ABI, and roadmap are in docs/PLAN.md.

Installation

# PyPI wheel: runtime, JIT kernel sources, vendored AOTriton, and server
pip install hipengine

# Source checkout: fetch Git LFS payloads before an editable install
git lfs install
git lfs pull
pip install -e .

# with the optional dlpack torch bridge for user-boundary interop
pip install "hipengine[torch]"

# dev / test
pip install -e ".[dev]"

Python 3.10+. A working ROCm install with libamdhip64.so on the loader path is required for any GPU run; CPU-reference correctness tests run without a GPU.

ROCm / TheRock setup for retained benchmark rows

For retained gfx1100 benchmark rows, use the pinned AMD TheRock environment in docs/THEROCK.md, not an ad-hoc mixed /opt/rocm runtime. Current retained rows use TheRock ROCm 7.13.0a20260423 with:

HIP version: 7.13.26162-1140233ffe

On this host (Linux 7.0.10-1-cachyos, W7900 VBIOS 113-D7070100-138, RX 7900 XTX VBIOS 113-EXT89622-001), ROCm 7.14 nightly diagnostics showed GGUF prefill and MTP wall-time regressions, so 7.13 remains the canonical stack until a newer ROCm release beats the same gates. See docs/THEROCK.md for the exact pip install/repair commands, clean process wrapper, and the upstream TheRock RELEASES.md reference.

The installed app exposes a small command group:

hipengine --help
hipengine serve --help
hipengine bench list

Quickstart

Model loading does not start network downloads. Populate the Hugging Face cache before using a repository ID:

hf download shisa-ai/Qwen3.6-35B-A3B-PARO-packed

Then construct LLM with the same repository ID:

from hipengine import LLM, SamplingParams

llm = LLM("shisa-ai/Qwen3.6-35B-A3B-PARO-packed")
outputs = llm.generate(
    ["Hello, hipEngine."],
    SamplingParams(max_tokens=64, temperature=0.0),
)
print(outputs[0])

LLM(model) auto-detects gfx1100 or gfx1151 and selects the model plugin's quantization. The Qwen3.6 GGUF path also selects T16 decode-repack plus the retained WMMA-prefill/GEMV-decode session profile. Explicit backend= and quant= arguments are overrides; supported PARO and GGUF models do not require hipEngine environment variables. Unsupported registry combinations fail instead of falling back to a torch path.

OpenAI-compatible server

The OpenAI-compatible FastAPI layer is installed by default:

pip install hipengine
hipengine serve \
  --model shisa-ai/Qwen3.6-35B-A3B-PARO-packed \
  --served-model-name qwen-paro

--model accepts either a local filesystem path or a Hugging Face model ID already present in the local HF cache; hipEngine resolves IDs locally and does not download weights during startup.

Core endpoints are GET /v1/models, POST /v1/completions, and POST /v1/chat/completions, with token-level SSE streaming, logprobs, OpenAI-style tools, structured-output validation, and Qwen thinking controls. hipEngine extensions provide readiness/capability discovery, token and context diagnostics, and app-local session management. Chat responses separate <think> reasoning into reasoning_content. The server eagerly warms the model on startup, caps omitted chat max_tokens with --chat-default-max-tokens (default 4096), and has an explicit --debug mode for full request/response payload logging. See the complete docs/API.md endpoint table for bearer-token auth, request examples, feature contracts, diagnostics, and current limitations.

Documentation

Development

# narrowest test suite (CPU-only paths run without a GPU)
pytest -q

# kernel source-lineage drift check before any port
python3 scripts/check_lineage.py --kind kernel --diff stat

See AGENTS.md for the full workflow: when to run the CPU-reference correctness gate, when to add a rocprofv3 --kernel-trace smoke, and what a retained benchmark row requires.

References & lineage

hipEngine is not a fork of any project; it is a brand new codebase with from-scratch code and kernels. Of course it builds on the work of many others:

• ROCm - of course this all sits on AMD's open-source compute stack, notably on HIP.
• Nano-vLLM - most of the original kernel tuning iteration loops used this as a host-layer. Some of the performance limitations of the architecture motivated the hipEngine rewrite, but we remain grateful and deeply appreciative of nano-vllm as a great research platform.
• ParoQuant - after reviewing the current SOTA on model quantization, we chose ParoQuant as the first target due to both its excellent accuracy and its efficiency (QTIP/YAQA is very cool but proved challenging to implement performant RDNA3 kernels)
• FastDMS - our KVCache ABI is shaped by the lessons learned from building our DMS reference implementation.

Greetz: hipfire, Lucebox, DS4, ExLlamaV3 and ofc the og llama.cpp

See also: Marlin, kernel-anvil, wmma_ops, tilelang, fsr4-rdna3-optimization, ROCm examples

License

hipEngine source code is licensed under AGPL-3.0-or-later. It is built and distributed for anyone who has an AMD card that hasn't been living up to its compute potential.

Model weights, checkpoints, and external datasets remain under their own licenses.