Feasibility Report: Porting ROCmFP4 to Atlas

Decision Brief

The short version: ROCmFP4 can be reimplemented inside Atlas, but the evidence does not yet justify doing it. The mandatory first step is to prove Atlas is faster than llama.cpp on AMD with standard FP8 weights.

Status: Updated 2025-06-16. Previous version contained incorrect assumptions about ROCm version, automation ceiling, and effort estimates. This version corrects those.

Executive Summary

Verdict: Technically possible, but there is no evidence it will be faster than llama.cpp + ROCmFP4. The effort is high (2-4 months) and the payoff is speculative.

The two projects are architecturally incompatible at the weight format layer:

The core problem is that Atlas does not support GGUF, Q4_0, or any 4-bit weight format other than NVFP4. To bring ROCmFP4 to Atlas, you would need to add:

1. 1. A GGUF loader (or a conversion pipeline from GGUF → Atlas's internal format)
2. 2. A new QuantFormat::RocmFp4 variant
3. 3. A new WeightDtype for ROCmFP4
4. 4. Hardware-specific HIP kernels for ROCmFP4 dequantization, matvec, matmul, FlashAttention
5. 5. Weight loading logic for every model architecture you want to support

The Critical Unknown: Does Atlas Beat llama.cpp on AMD?

This is the question that must be answered before any ROCmFP4 work. The available data suggests it might not.

Baseline: llama.cpp + ROCmFP4 (Already Works)

These are end-to-end, regression-guarded, reproducible numbers. The ROCmFP4 repo has 8 guard scripts that fail CI if these numbers drop.

Atlas on AMD: The Only Data Point

From Atlas's own AMD port README (README-HIP.md):

spark serve Qwen/Qwen3.6-27B-FP8 …; coherence check; decode tok/s via ~/bench-atlas.sh (target 20)

Atlas's own target for Qwen3.6-27B on AMD is 20 tok/s with FP8 weights.

Compare: llama.cpp gets 28-33 tok/s for the same model on the same hardware with ROCmFP4.

So Atlas is targeting to be slower than llama.cpp's current baseline, even before adding ROCmFP4.

Why Atlas's "2-3x Faster" Claim Doesn't Apply Here

Atlas's README says:

We make no compromises or generalizations. Each hardware and model combination has its own unique properties that require fine-tuning custom kernels that leverage the model for that specific hardware configuration. The end result? 2-3x faster kernels all around.

But this claim is only validated on NVIDIA GB10 (Blackwell):

There is zero AMD data. Atlas has no published AMD benchmark, no AMD comparison to llama.cpp, no AMD comparison to anything.

Why AMD is Different from NVIDIA

On NVIDIA, Atlas's model-specific kernels extract value from tensor cores and custom launch geometries. On AMD, the bottleneck is memory bandwidth, and llama.cpp's generic ggml kernels are already very well-tuned for that. The ROCmFP4-specific optimizations (unaligned dword loads, two-warp geometry, Codebook10 expanders) are already in llama.cpp.

The Honest Verdict

Bottom line: There is no business case for Atlas + ROCmFP4 on AMD based on the available data. The most likely outcome is that you spend 2-3 months and end up with something equal or slower than what already exists.

What You Should Do First: The Minimal Validation Experiment

Before any ROCmFP4 work, answer this question:

Does Atlas's model-specific kernel architecture produce meaningful speedups on AMD with standard weights?

The "Does Atlas Beat llama.cpp on AMD?" Test (1–2 weeks)

Goal: Determine if Atlas's architecture has any speed advantage on AMD at all.

What to do:

1. 1. Complete the 2–3 MMA kernels needed for Atlas AMD prefill (w4a16_gemm, dense_gemm_tc, inferspark_prefill)
2. 2. Stub the remaining 17 MMA kernels (they don't affect decode)
3. 3. Build Atlas AMD
4. 4. Serve Qwen3.6-27B-FP8 (or Qwen3.6-35B-A3B-FP8) on your Strix Halo
5. 5. Measure decode tok/s
6. 6. Run llama.cpp on the same hardware with the same model and FP8 weights
7. 7. Compare

Decision criteria:

• If Atlas FP8 > llama.cpp FP8 by ≥20%: Atlas's architecture has real value on AMD. Pursue ROCmFP4 integration.
• If Atlas FP8 ≈ llama.cpp FP8 (±20%): The architectural advantage is marginal. ROCmFP4 might not be worth the effort.
• If Atlas FP8 < llama.cpp FP8: Don't bother. The ROCmFP4 format is the speed, and llama.cpp already has it.

Why this is the right experiment:

• It answers the question before committing to the full ROCmFP4 integration
• It only requires ~2 weeks of work (the 2-3 MMA kernel ports + build wiring)
• It uses weights that already exist (FP8 from HuggingFace)
• It directly compares the two engines on the same hardware, same model, same format

Why NOT do the ROCmFP4 integration first:

• If Atlas's architecture is slower on AMD, adding ROCmFP4 won't fix it
• The ROCmFP4 integration is 2-3 months of work
• You would be optimizing a slower engine

What is ROCmFP4?

ROCmFP4 is a custom 4-bit quantization format developed for AMD Strix Halo (gfx1151). It is not MXFP4, NVFP4, or a standard Q4_0 variant. It has unique characteristics:

Block Layout

// Dual-scale (Q4_0_ROCMFP4): 18 bytes per 32-weight block, 4.50 BPW
struct block_rocmfp4 {
    uint8_t qs[16];    // 32 packed 4-bit nibbles (E2M1-derived Codebook10)
    uint8_t e[2];      // Two unsigned E4M3 scale bytes, one per 16-weight half
};

// Fast single-scale (Q4_0_ROCMFP4_FAST): 17 bytes per block, 4.25 BPW
struct block_rocmfp4_fast {
    uint8_t qs[16];    // Same 32 packed nibbles
    uint8_t e;         // One UE4M3 scale for the whole block
};

Codebook10 Value Table

static const int8_t rocmfp4_codebook[16] = {
     0,  1,  2,  3,  4,  6,  8, 10,
     0, -1, -2, -3, -4, -6, -8,-10,
};

The value table is not standard E2M1 (which would be {0, 0.5, 1, 1.5, 2, 3, 4, 6}). It is a tuned signed-integer representation where the largest magnitude is retuned from 12 to 10 to reduce outlier pull on Qwen3-family tensors.

Scale Format

ROCmFP4 uses finite unsigned E4M3 scale bytes decoded to half-scale values. This is a custom scale format, not standard FP8 E4M3. The scale decoder builds normal FP32 values directly from exponent/mantissa bits, skipping NaN/Inf handling because ROCmFP4 row validation already rejects non-finite scale bytes.

Tensor-Aware Presets

ROCmFP4 has six presets that mix layouts across tensor roles:

This tensor-aware preset system is not something you can simply "drop in" — it requires the inference engine to understand which tensor is which (token embeddings vs. attention K vs. FFN weights) and apply different quantization formats per tensor.

What is Atlas?

Atlas is a pure Rust inference engine with a modular, trait-based architecture. It is fundamentally different from llama.cpp:

Weight Loading Pipeline

safetensors file → parse header → mmap or O_DIRECT read →
  GPU allocation → WeightStore (HashMap<String, WeightTensor>) →
  ModelWeightLoader (per-model) → TransformerLayer objects

Atlas only supports safetensors files. The WeightDtype enum:

pub enum WeightDtype {
    BF16,      // 2 bytes
    FP32,      // 4 bytes
    FP8E4M3,   // 1 byte
    UInt8,     // 1 byte
}

There is no Q4_0, no Q5_K, no Q6_K, no GGUF support. Atlas's architecture assumes weights are in one of these four dtypes, loaded from HuggingFace-style safetensors checkpoints.

Quantization Format Enum

pub enum QuantFormat {
    Nvfp4,  // NVFP4 E2M1 + FP8 block scales
    Fp8,    // FP8 E4M3 block-scaled
}

The forward pass (quant_weights.rs trait) dispatches to model-specific, hardware-specific, quantization-specific kernels. The kernel directory structure:

kernels/
  gb10/
    qwen3.6-35b-a3b/
      nvfp4/          ← NVFP4 kernels for this model on GB10
    qwen3.6-27b/
      nvfp4/
  strix-hip/          ← Would be the AMD target directory
    qwen3.6-27b/
      rocmfp4/        ← Would need to exist

Kernel System

Atlas kernels are CUDA source files (.cu) compiled to PTX for NVIDIA, or compiled via hipcc with compat headers to HSACO (.co) for AMD. The build system (build.rs / build_target.rs) auto-discovers kernels from kernels/<hw>/<model>/<quant>/ directories.

Key kernel types for each (hw, model, quant) target:

• w4a16_gemm — weight-only 4-bit GEMM (prefill)
• inferspark_prefill — attention prefill kernel
• gemv / gemv_gate_up — decode matvec
• dense_gemm_tc — tensor-core dense GEMM
• moe_* — MoE routing kernels
• reshape_and_cache — KV cache management
• flash_attention — FlashAttention kernel

For ROCmFP4, you would need to write all of these for the AMD target, using the ROCmFP4 block format and Codebook10 decode in the kernels.

AMD Port Status

The port/amd-strix-halo branch has a partial AMD port:

• 72/92 CUDA kernels compile clean with hipcc using compat headers
• 20 kernels require hand-porting from NVIDIA mma.sync PTX to AMD WMMA (__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32)
• A libcuda.so → HIP shim maps 33 CUDA driver symbols to HIP equivalents
• Status: "build path validated; full-model serve pending build-system wiring"

If the Validation Experiment Passes: Integration Options

Option 1: GGUF-to-Safetensors Converter (2–3 weeks, 1 engineer)

Convert ROCmFP4 GGUF → BF16/FP8 safetensors.

• Extract and dequantize all ROCmFP4 blocks using the reference decoder
• Write standard safetensors files Atlas can load
• What you gain: You can run the models (Qwen3.6 35B A3B, etc.) in Atlas
• What you lose: 4-bit memory savings, tuned decode speed, tensor-aware precision mixing
• Risk: LOW. This is mostly a file format conversion.

Option 2: Native ROCmFP4 WeightDtype in Atlas (2–3 months, 1–2 engineers)

Add ROCmFP4 as a native packed weight format in Atlas, write decode kernels only.

• Add WeightDtype::RocmFp4 and QuantFormat::RocmFp4 to Atlas
• Write a GGUF → Atlas converter that preserves 4-bit blocks
• Write gemv and gemv_gate_up HIP kernels for decode (the hot path)
• Defer FlashAttention and prefill optimization to dequant→BF16 fallback
• What you gain: Memory savings + decode speedup
• What you lose: Full prefill speed, FlashAttention K/V cache optimization
• Risk: MEDIUM. The decode kernels are the core value.

Option 3: Full Native Integration (3–6 months, 2-3 engineers)

Full ROCmFP4 support: GGUF loader, all kernels, tensor-aware presets, regression guards.

• Implement GGUF parser in Rust
• Write all model-specific kernels (gemm, flash_attention, moe_routing, reshape_and_cache)
• Port tensor-aware preset logic to every ModelWeightLoader
• Port ROCmFP4's regression guard system
• What you gain: Full parity with ROCmFP4-llama performance
• Risk: MEDIUM-HIGH. Requires sustained Strix Halo access for validation.

Automation Reality: What an Agent Can Actually Do

The previous version of this report incorrectly stated that ~40% of the work (kernel tuning, regression validation) was "not automatable without hardware." This is wrong. The correct framing is:

• ~70% of the labor is fully automatable (format extraction, reference decoder transpilation, build system integration, test orchestration)
• ~20% is semi-automatable (kernel skeleton generation, model dispatch tables)
• ~10% requires hardware but is still automatable by an agent (compile-run-measure loops, parameter sweeps)

The kernel tuning is not a human-vs-agent problem. It is a hardware-access problem. An agent with SSH to a Strix Halo machine can do exactly what the ROCmFP4 author did: hypothesize, code, compile, measure, keep or reject. The autoresearch skill is designed for this.

The Agent's Actual Advantage

A human tuning kernels on Strix Halo would run 5-10 experiments per day (compile + benchmark takes 10-30 min each). An agent running a persistent loop can run 50-100 experiments per day (24/7, no sleep, no bias, perfect record-keeping).

The ROCmFP4 author spent months doing manual tuning. An agent could do it in 1-2 weeks of continuous search because:

1. 1. It can read the ROCmFP4 README tuning notes as a specification rather than rediscovering them empirically
2. 2. It can run the compile-measure loop continuously
3. 3. It never forgets results or gets tired

The Automation Pipeline

Phase A: Specification Automation (1 week, any machine)

• Parse ggml/rocmfp4/ source files → extract format spec, Codebook10 table, UE4M3 decoder
• Generate Rust reference decoder
• Generate HIP kernel stubs with decode helpers pre-filled

Phase B: Converter Automation (1 week, any machine)

• Read GGUF metadata → extract tensor roles, apply preset rules
• Output Atlas WeightStore + config.json + MODEL.toml

Phase C: Kernel Skeleton Automation (2 weeks, any machine)

• Generate kernels/strix-hip/<model>/rocmfp4/ directory structure
• Fill in decode logic from spec
• Leave TUNE: markers for launch geometry, tile sizes, shared memory

Phase D: Hardware-Dependent Validation (1-2 weeks, needs Strix Halo)

• Compile → benchmark → measure → decide loop
• The agent can run this autonomously: edit kernel, compile, run microbench + MTP guard, parse results, git commit or revert, update state
• Search space is guided by the ROCmFP4 README (which documents the promoted parameters)

Total effort with agent: 4-6 weeks for full integration, assuming the validation experiment proves Atlas has value on AMD.

Tools and Dependencies

Note: Atlas's README-HIP.md mentions "ROCm 6.x" but the ROCmFP4-llama repo validates on ROCm 7.2.1. Use whatever version works on your hardware.

Recommendation

Step 1 (Mandatory): Run the minimal validation experiment.

• Complete the 2-3 MMA kernel ports needed for Atlas AMD prefill
• Build Atlas, serve Qwen3.6-27B-FP8, measure decode tok/s
• Compare to llama.cpp + FP8 on the same hardware
• If Atlas FP8 is not ≥20% faster than llama.cpp FP8: Stop. The ROCmFP4 integration will not be worth it.

Step 2 (If validation passes): Build the GGUF→safetensors converter (Option 1).

• 2-3 weeks of work, immediately validates model coherence in Atlas
• Proves the format works end-to-end before committing to kernel work

Step 3 (If converter proves valuable): Add native ROCmFP4 decode kernels (Option 2).

• 2-3 months of work, gives memory savings + decode speedup
• The agent can handle the bulk of the compile-measure-iterate loop autonomously

Step 4 (Only if deeply justified): Full integration (Option 3).

• 3-6 months, only if the ROCmFP4 author collaborates or the business case is clear

Critical prerequisite: Atlas's AMD port must be completed first. The current status is 72/92 kernels compile, 20 need hand-porting, full-model serve pending. The validation experiment will also force completion of this prerequisite.

Conclusion

The ROCmFP4 format is a deeply specialized, AMD-tuned, tensor-aware quantization system built on top of llama.cpp's ggml backend. Atlas is a pure Rust, safetensors-based, hardware+model-specific kernel engine with a completely different architecture.

Bringing ROCmFP4 to Atlas is not a "port" in the conventional sense. It is a reimplementation of the ROCmFP4 quantization format, decode kernels, and tensor-aware preset system within Atlas's kernel framework.

But the most important question is not "can we do it?" — it's "should we do it?" And the answer to that depends entirely on whether Atlas's architecture produces speedups on AMD at all. The available data suggests it might not. The only way to know is the 1-2 week validation experiment.

Do not commit to the full ROCmFP4 integration until you have proven that Atlas is faster than llama.cpp on AMD with standard weights.

Report generated by autoresearch skill analysis. Updated 2025-06-16. Sources: github.com/charlie12345/rocmfp4-llama (commit 4860505ee), github.com/Avarok-Cybersecurity/atlas (branch port/amd-strix-halo).