Skip to content

Add gfx1100 HIP support with tuned prefill and bounded expert uploads - #121

Closed
2jztricks wants to merge 8 commits into
Niko1221:mainfrom
2jztricks:feat/hip-gfx1100-prefill
Closed

2jztricks wants to merge 8 commits into
Niko1221:mainfrom
2jztricks:feat/hip-gfx1100-prefill

Conversation

@2jztricks

@2jztricks 2jztricks commented Sep 29, 2026 •

Copy link
Copy Markdown
Contributor

RX 7900 XTX users need a native HIP path and bounded host-memory transfers to serve Flash Next reliably. This replaces #94 with the complete opt-in gfx1100 backend rebased onto current main, adds tuned prefill, and includes fresh performance evidence.

Fresh 2026-09-29 measurements on RX 7900 XTX at the unchanged 272 W cap, Orca Flash Next IQ3_XXS, MTP 4, configured 262144 context, 8192 prefill batch:

Request Prefill tokens/s Output tokens/s Total request time
4210 tokens, first use 447.5 55.5 11.729 s
8830 tokens, first use at this size 873.7 55.7 12.420 s
4210 tokens, warmed 901.1 56.7 6.937 s
8830 tokens, warmed 966.5 59.2 11.314 s

Each fresh prompt reused zero KV tokens and generated the intentional 128-token cap. All four fresh prompts and four cached follow-ups completed. A newly measured same-session existing HIP runtime control gives 1.80–2.02x prefill and 40–46% lower fresh-request wall time. This is not a comparison against unmodified upstream, and configured 262K capacity is not a full-context benchmark.

Validation: HIP and CUDA executable builds; 29/29 selected HIP CTests; three real IQ3_XXS PLE matrix graph replays; POSIX direct-file queued reads/short EOF/wake/close/reopen test. The two excluded CTests and all measurement conditions, exact source/binary identifiers and sanitized raw timings are documented in docs/AMD_HIP_PERFORMANCE.md. CUDA has compile validation here, not a new CUDA throughput run. These numerical and capped-throughput tests do not establish coding quality, full-context/vision behavior or mixed-vendor execution.

Supersedes #94. Related: #48. No model weights or personal serving/gateway configuration are included.

xyzzing added a commit to xyzzing/Strata that referenced this pull request Sep 29, 2026
…nded expert uploads (supersedes Niko1221#94)

# Conflicts:
#	CMakeLists.txt
#	include/strata/platform/direct_file.hpp
#	src/core/device.cu
#	src/kernels/cuda/native_qsa_indexer.cu
#	src/platform/direct_file.cpp
#	src/prefill/prefill.cpp
#	src/program/generate.cpp
xyzzing added a commit to xyzzing/Strata that referenced this pull request Sep 29, 2026
…end the compat layer

Upstream main added TF32/f16 tensor-core QSA scoring, cp.async prompt
attention, graph introspection and a GPU timestamp after this PR's base;
none of it compiles under HIP. The tensor-core and cp.async kernels stay
CUDA-only (the existing warp/old-kernel fallbacks serve HIP, as for pre-
sm_80 cards), the graph-node and device-attribute names join the compat
defines, math_constants.h gets a compat shim, the stage profiler uses
clock64() under HIP, and __trap becomes __builtin_trap. Verified: full
HIP build of strata and the CPU build with 8/10 ctests (the two failures
need a packed-model fixture; excluded like the PR's own validation).
Niko1221 added a commit that referenced this pull request Sep 29, 2026
@Niko1221

Copy link
Copy Markdown
Owner

Merged into 0.1.25 (released) as one squashed commit under your name, rebased onto 0.1.24. What changed on the way in:

  • Dropped duplicates: 0.1.24 already batches the QSA indexer append and the embedding gather (C-2, C-4) and has a multi-threaded PLE reader, so your variants of those were dropped and your HIP tests point at the existing functions.
  • CUDA-only paths: the new tensor-core kernels (D-1 prompt attention, QSA select) are compiled out under HIP and refused at run time, so AMD runs the portable kernels.
  • HIP compile fixes: __trap, the compute-capability attribute, the graph-node listing, and the GPU timestamp; plus NOMINMAX for Windows.
  • Installer: ./setup.sh --backend hip - KFD detection (gfx1100 only), ROCm from AMD's TheRock wheels into .venv without sudo, the engine compiled on the user's PC, and a hipBLASLt table re-tuned for hipBLASLt 100200 (the pinned ROCm) with your tuner: +42-59% prompt speed on a 7900 XTX.

Validated on an RX 7900 XTX with the Coder IQ1_M:

  • the HIP suite (30/30);
  • a live server test of every endpoint;
  • needle tests;
  • ~1,320 tokens/s prompt at 16-32K and 58-68 tokens/s output.

The CUDA side stays byte-identical to 0.1.24. Thank you @2jztricks - excellent work.

@Niko1221 Niko1221 closed this Sep 29, 2026
xyzzing added a commit to xyzzing/Strata that referenced this pull request Oct 4, 2026
…gfx1100 docs) into the 0.1.38 line — ours everywhere: the 0.1.38 line supersedes the old HIP port code; GFX1100.md and the fork CI/docs come over
Hugua700 pushed a commit to Hugua700/Strata that referenced this pull request Oct 6, 2026
Hugua700 pushed a commit to Hugua700/Strata that referenced this pull request Oct 6, 2026
…based onto 0.1.24) and its setup support
anon761 pushed a commit to anon761/Strata-upstream that referenced this pull request Oct 6, 2026
anon761 pushed a commit to anon761/Strata-upstream that referenced this pull request Oct 6, 2026
…based onto 0.1.24) and its setup support
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 539 commit(s) a0116b2..c4bcb6e on https://github.com/eddoursul/Strata, restricted to the
paths this fork still mirrors: 1 file(s) under root -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  CMakeLists.txt

Original authors: Daehun,Daerdaal,Eddoursul,Enky

Original commits touching these paths (88 subject(s)):
  - 4-bit KV caches (--kv q4_0, --kv k8v4) in the tensor-core attention kernels
  - A fatal CUDA fault in the prompt path exits at once instead of hanging until the watchdog (Niko1221#224); CMake warns when sm_120 is built with CUDA < 13.0 (Niko1221#220, Niko1221#224); iq_pack writes native_experts.txt under a temporary name, so a stop part-way leaves no partial pack (Niko1221#172)
  - A second GPU as another expert tier (--second-gpu)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Accept sm_80+ GPUs at runtime, matching the build guard and the README
  - Add optional multi-GPU expert caches and compact transfers
  - Add reproducible Pi cache comparison and document current integration
  - Build without compiler warnings

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eddoursul/Strata
Fork-Replay-Branch: custom
Fork-Replay-Base: a0116b2
Fork-Replay-Tip: c4bcb6e
Fork-Replay-Group: root
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 539 commit(s) a0116b2..c4bcb6e on https://github.com/eddoursul/Strata, restricted to the
paths this fork still mirrors: 46 file(s) under src/kernels -- 35 clean,
11 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/kernels/cpu/native_expert.cpp
  src/kernels/cpu/pool.cpp
  src/kernels/cuda/fused_gr.cu
  src/kernels/cuda/iq_kernels.cu
  src/kernels/cuda/native_mmvq.cu
  src/kernels/cuda/native_qsa_indexer.cu
  src/kernels/cuda/native_router.cu
  src/kernels/cuda/qsa_decode_attn.cu
  src/kernels/cuda/qsa_select.cu
  src/kernels/cuda/verify_kernels.cu
  src/kernels/qsa_parity.cpp

Original authors: Andy,Eddoursul,Guillaume PUTIER,Hozák Martin

Original commits touching these paths (118 subject(s)):
  - Niko1221#152 as an opt-in (the owner's choice): STRATA_IQ_MT_MIN=1 makes CPU expert rounding independent of the drafting; the default stays 0.1.29's rule (IQ3_S decoded 1-3% slower with it on); docs; the width test runs the opt-in
  - --ple-io ram: READY only once the PLE table is in RAM (its load competed with the first requests: a first chat at ~21 tok/s instead of ~95); STRATA_PLE_RAM_NOWAIT=1 keeps the old order
  - --ple-io ram: the PLE table read into RAM once the experts are loaded (28.8 GB in ~7 s), direct reads until then; a 16K-token prompt of new text 0.5 s faster
  - --window-profile: GPU time of each stage of the verify window
  - 4-bit KV caches (--kv q4_0, --kv k8v4) in the tensor-core attention kernels
  - A one-token window commits itself
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Build without compiler warnings

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eddoursul/Strata
Fork-Replay-Branch: custom
Fork-Replay-Base: a0116b2
Fork-Replay-Tip: c4bcb6e
Fork-Replay-Group: src/kernels
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 539 commit(s) a0116b2..c4bcb6e on https://github.com/eddoursul/Strata, restricted to the
paths this fork still mirrors: 5 file(s) under src/prefill -- 1 clean,
4 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/prefill/gemm.cu
  src/prefill/kernels.cu
  src/prefill/moe_mmq.cu
  src/prefill/prefill.cpp

Original authors: Eddoursul,Konstantinos Korres,Niko1221,code-martin

Original commits touching these paths (47 subject(s)):
  - 4-bit KV caches (--kv q4_0, --kv k8v4) in the tensor-core attention kernels
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - C-1, C-2: the prompt path's QSA select grid reaches only the active blocks; the indexer appends a chunk at once
  - C-3: the prompt path's GDN short conv tiled over tokens
  - C-4: the prompt's embedding in one gather, and the n-gram table read by several threads
  - D-1: the prompt path's QSA attention on tensor cores
  - D-2: the prompt path's GDN recurrence split over value columns, the output norm in its own kernel
  - D-5: the prompt path's expert stream is issued by its own host thread

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eddoursul/Strata
Fork-Replay-Branch: custom
Fork-Replay-Base: a0116b2
Fork-Replay-Tip: c4bcb6e
Fork-Replay-Group: src/prefill
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 26 file(s) under include/strata -- 25 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  include/strata/kernels/verify_kernels.hpp

Original authors: 4EverBuilder,Andy,Anon,BlueKingMuch

Original commits touching these paths (152 subject(s)):
  - Niko1221#253: the 8 GiB arena pin cap only under WDDM; STRATA_ARENA_PIN_GIB; the split's prompt chunk steps down
  - Niko1221#284 follow-up: wait for the async commit on an event, at every point that touches the session afterwards
  - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3
  - Niko1221#340 prototype: a split's streamed ring is 96 slots when most experts are resident (STRATA_SPLIT_RING); STRATA_SPLIT_SMALL_OWN=S keeps an S-token chunk's buffers per lending stage for the session; --no-prefill-borrow on a split books the ring (it no longer starves the draft head)
  - Niko1221#465 engine: batch slots for one GPU - warn instead of refuse, windows over the active slots, slot conversation cache, prompts read beside decoding slots, BYIELD
  - Niko1221#583 default on (owner): STRATA_RING_BYTES=0 opts out; DETAILS.md with the measured quality
  - Niko1221#583: the byte-budget ring only where it gains; short prompts keep 0.1.39's ring
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: include/strata
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 1 file(s) under root -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  CMakeLists.txt

Original authors: Andy,BlueKingMuch,Chanyeong Lim,Daehun

Original commits touching these paths (152 subject(s)):
  - Niko1221#136: fused int8 tensor-core experts for the Q2_0 pack's prompt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1)
  - Niko1221#136: the fused int8 prompt experts for the native IQ packs (gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1)
  - Niko1221#290 follow-up: a BF16-row parity test (native_expert_parity --bf16-embd)
  - Niko1221#296 follow-up: only the Q5_0 MMQ instance and case stay (Q4_K / Q5_1 come with S5's STRATA_MMQ_KQUANTS)
  - Niko1221#296 follow-up: ple_q5_parity / q5_projection_parity build only with STRATA_BUILD_TESTS; the doc says what runs where
  - A fatal CUDA fault in the prompt path exits at once instead of hanging until the watchdog (Niko1221#224); CMake warns when sm_120 is built with CUDA < 13.0 (Niko1221#220, Niko1221#224); iq_pack writes native_experts.txt under a temporary name, so a stop part-way leaves no partial pack (Niko1221#172)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested.

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: root
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 21 file(s) under src/kernels -- 14 clean,
7 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/kernels/cpu/native_expert.cpp
  src/kernels/cpu/pool.cpp
  src/kernels/cuda/dequant_bf16.cu
  src/kernels/cuda/fused_gr.cu
  src/kernels/cuda/iq_kernels.cu
  src/kernels/cuda/verify_kernels.cu
  src/kernels/ngram.cpp

Original authors: 4EverBuilder,Andy,BlueKingMuch,Chanyeong Lim

Original commits touching these paths (151 subject(s)):
  - Niko1221#152 as an opt-in (the owner's choice): STRATA_IQ_MT_MIN=1 makes CPU expert rounding independent of the drafting; the default stays 0.1.29's rule (IQ3_S decoded 1-3% slower with it on); docs; the width test runs the opt-in
  - Niko1221#289 follow-up: a placeholder GPU name when the properties query fails; HIP builds compile (no device_code_error probe)
  - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3
  - Niko1221#646 on 0.1.39b: the layer split and the CPU path as before, the resident plan for 8-token windows
  - Niko1221#646: the single-matrix mmvq does not stage the IQ grids (10-20% slower on the RTX 5070 at 3-8 columns)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested.
  - C-1, C-2: the prompt path's QSA select grid reaches only the active blocks; the indexer appends a chunk at once

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: src/kernels
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 1 file(s) under src/ngram -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/ngram/ple_reader.cpp

Original authors: BlueKingMuch,Eli Elgaev,Konstantinos Korres,Niko1221

Original commits touching these paths (5 subject(s)):
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Helper threads leave the host thread's core (Linux)
  - PLE: the n-gram table in FP8 E4M3, as Qwen ships it
  - Strata: Qwen3.8-Flash-Next on a 12-24 GB NVIDIA GPU + 64 GB RAM, one-click install
  - ple: keep the SSD awake while rows are read (STRATA_SSD_KEEPALIVE)

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: src/ngram
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
…glm53f

Squashed from 709 commit(s) 61b3fb5..7a33c9c on https://github.com/lighttransport/Strata, restricted to the
paths this fork still mirrors: 49 file(s) under include/strata -- 45 clean,
4 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  include/strata/hip_compat/cublas_v2.h
  include/strata/kernels/cpu/native_expert.hpp
  include/strata/kernels/iq_kernels.hpp
  include/strata/prefill/moe_mmq.hpp

Original authors: BlueKingMuch,Guillaume PUTIER,Konstantinos Korres,Léon Avic Simmons

Original commits touching these paths (65 subject(s)):
  - Niko1221#272 follow-up: --pool-affinity all (the pool's layout so far) is the default; auto and p-cores are opt-in
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Add NUMA bandwidth benchmarks and native GLM EXL3 decode
  - Add guarded opt-in GLM prefill and bounded drafting experiments
  - Add opt-in GLM external drafting and guarded optimization records
  - Add opt-in GLM long-context prefill experiments
  - B550 decode levers: fused mHC read, direct tier DMA, MLA decode kernels, partial expert deferral
  - Coupled draft sampling for the MTP drafter (STRATA_SPEC_COUPLED=1, off by default)

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/lighttransport/Strata
Fork-Replay-Branch: glm53f
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: 7a33c9c
Fork-Replay-Group: include/strata
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 709 commit(s) 61b3fb5..7a33c9c on https://github.com/lighttransport/Strata, restricted to the
paths this fork still mirrors: 1 file(s) under root -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  CMakeLists.txt

Original authors: Andy,BlueKingMuch,Daehun,Daerdaal

Original commits touching these paths (138 subject(s)):
  - Niko1221#136: fused int8 tensor-core experts for the Q2_0 pack's prompt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1)
  - Niko1221#136: the fused int8 prompt experts for the native IQ packs (gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1)
  - Niko1221#290 follow-up: a BF16-row parity test (native_expert_parity --bf16-embd)
  - Niko1221#296 follow-up: only the Q5_0 MMQ instance and case stay (Q4_K / Q5_1 come with S5's STRATA_MMQ_KQUANTS)
  - Niko1221#296 follow-up: ple_q5_parity / q5_projection_parity build only with STRATA_BUILD_TESTS; the doc says what runs where
  - A fatal CUDA fault in the prompt path exits at once instead of hanging until the watchdog (Niko1221#224); CMake warns when sm_120 is built with CUDA < 13.0 (Niko1221#220, Niko1221#224); iq_pack writes native_experts.txt under a temporary name, so a stop part-way leaves no partial pack (Niko1221#172)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested.

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/lighttransport/Strata
Fork-Replay-Branch: glm53f
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: 7a33c9c
Fork-Replay-Group: root
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 539 commit(s) a0116b2..c4bcb6e on https://github.com/eddoursul/Strata, restricted to the
paths this fork still mirrors: 1 file(s) under root -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  CMakeLists.txt

Original authors: Daehun,Daerdaal,Eddoursul,Enky

Original commits touching these paths (88 subject(s)):
  - 4-bit KV caches (--kv q4_0, --kv k8v4) in the tensor-core attention kernels
  - A fatal CUDA fault in the prompt path exits at once instead of hanging until the watchdog (Niko1221#224); CMake warns when sm_120 is built with CUDA < 13.0 (Niko1221#220, Niko1221#224); iq_pack writes native_experts.txt under a temporary name, so a stop part-way leaves no partial pack (Niko1221#172)
  - A second GPU as another expert tier (--second-gpu)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Accept sm_80+ GPUs at runtime, matching the build guard and the README
  - Add optional multi-GPU expert caches and compact transfers
  - Add reproducible Pi cache comparison and document current integration
  - Build without compiler warnings

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eddoursul/Strata
Fork-Replay-Branch: custom
Fork-Replay-Base: a0116b2
Fork-Replay-Tip: c4bcb6e
Fork-Replay-Group: root
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 539 commit(s) a0116b2..c4bcb6e on https://github.com/eddoursul/Strata, restricted to the
paths this fork still mirrors: 46 file(s) under src/kernels -- 35 clean,
11 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/kernels/cpu/native_expert.cpp
  src/kernels/cpu/pool.cpp
  src/kernels/cuda/fused_gr.cu
  src/kernels/cuda/iq_kernels.cu
  src/kernels/cuda/native_mmvq.cu
  src/kernels/cuda/native_qsa_indexer.cu
  src/kernels/cuda/native_router.cu
  src/kernels/cuda/qsa_decode_attn.cu
  src/kernels/cuda/qsa_select.cu
  src/kernels/cuda/verify_kernels.cu
  src/kernels/qsa_parity.cpp

Original authors: Andy,Eddoursul,Guillaume PUTIER,Hozák Martin

Original commits touching these paths (118 subject(s)):
  - Niko1221#152 as an opt-in (the owner's choice): STRATA_IQ_MT_MIN=1 makes CPU expert rounding independent of the drafting; the default stays 0.1.29's rule (IQ3_S decoded 1-3% slower with it on); docs; the width test runs the opt-in
  - --ple-io ram: READY only once the PLE table is in RAM (its load competed with the first requests: a first chat at ~21 tok/s instead of ~95); STRATA_PLE_RAM_NOWAIT=1 keeps the old order
  - --ple-io ram: the PLE table read into RAM once the experts are loaded (28.8 GB in ~7 s), direct reads until then; a 16K-token prompt of new text 0.5 s faster
  - --window-profile: GPU time of each stage of the verify window
  - 4-bit KV caches (--kv q4_0, --kv k8v4) in the tensor-core attention kernels
  - A one-token window commits itself
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Build without compiler warnings

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eddoursul/Strata
Fork-Replay-Branch: custom
Fork-Replay-Base: a0116b2
Fork-Replay-Tip: c4bcb6e
Fork-Replay-Group: src/kernels
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 539 commit(s) a0116b2..c4bcb6e on https://github.com/eddoursul/Strata, restricted to the
paths this fork still mirrors: 5 file(s) under src/prefill -- 1 clean,
4 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/prefill/gemm.cu
  src/prefill/kernels.cu
  src/prefill/moe_mmq.cu
  src/prefill/prefill.cpp

Original authors: Eddoursul,Konstantinos Korres,Niko1221,code-martin

Original commits touching these paths (47 subject(s)):
  - 4-bit KV caches (--kv q4_0, --kv k8v4) in the tensor-core attention kernels
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - C-1, C-2: the prompt path's QSA select grid reaches only the active blocks; the indexer appends a chunk at once
  - C-3: the prompt path's GDN short conv tiled over tokens
  - C-4: the prompt's embedding in one gather, and the n-gram table read by several threads
  - D-1: the prompt path's QSA attention on tensor cores
  - D-2: the prompt path's GDN recurrence split over value columns, the output norm in its own kernel
  - D-5: the prompt path's expert stream is issued by its own host thread

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eddoursul/Strata
Fork-Replay-Branch: custom
Fork-Replay-Base: a0116b2
Fork-Replay-Tip: c4bcb6e
Fork-Replay-Group: src/prefill
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 26 file(s) under include/strata -- 25 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  include/strata/kernels/verify_kernels.hpp

Original authors: 4EverBuilder,Andy,Anon,BlueKingMuch

Original commits touching these paths (152 subject(s)):
  - Niko1221#253: the 8 GiB arena pin cap only under WDDM; STRATA_ARENA_PIN_GIB; the split's prompt chunk steps down
  - Niko1221#284 follow-up: wait for the async commit on an event, at every point that touches the session afterwards
  - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3
  - Niko1221#340 prototype: a split's streamed ring is 96 slots when most experts are resident (STRATA_SPLIT_RING); STRATA_SPLIT_SMALL_OWN=S keeps an S-token chunk's buffers per lending stage for the session; --no-prefill-borrow on a split books the ring (it no longer starves the draft head)
  - Niko1221#465 engine: batch slots for one GPU - warn instead of refuse, windows over the active slots, slot conversation cache, prompts read beside decoding slots, BYIELD
  - Niko1221#583 default on (owner): STRATA_RING_BYTES=0 opts out; DETAILS.md with the measured quality
  - Niko1221#583: the byte-budget ring only where it gains; short prompts keep 0.1.39's ring
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: include/strata
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 1 file(s) under root -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  CMakeLists.txt

Original authors: Andy,BlueKingMuch,Chanyeong Lim,Daehun

Original commits touching these paths (152 subject(s)):
  - Niko1221#136: fused int8 tensor-core experts for the Q2_0 pack's prompt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1)
  - Niko1221#136: the fused int8 prompt experts for the native IQ packs (gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1)
  - Niko1221#290 follow-up: a BF16-row parity test (native_expert_parity --bf16-embd)
  - Niko1221#296 follow-up: only the Q5_0 MMQ instance and case stay (Q4_K / Q5_1 come with S5's STRATA_MMQ_KQUANTS)
  - Niko1221#296 follow-up: ple_q5_parity / q5_projection_parity build only with STRATA_BUILD_TESTS; the doc says what runs where
  - A fatal CUDA fault in the prompt path exits at once instead of hanging until the watchdog (Niko1221#224); CMake warns when sm_120 is built with CUDA < 13.0 (Niko1221#220, Niko1221#224); iq_pack writes native_experts.txt under a temporary name, so a stop part-way leaves no partial pack (Niko1221#172)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested.

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: root
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 21 file(s) under src/kernels -- 14 clean,
7 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/kernels/cpu/native_expert.cpp
  src/kernels/cpu/pool.cpp
  src/kernels/cuda/dequant_bf16.cu
  src/kernels/cuda/fused_gr.cu
  src/kernels/cuda/iq_kernels.cu
  src/kernels/cuda/verify_kernels.cu
  src/kernels/ngram.cpp

Original authors: 4EverBuilder,Andy,BlueKingMuch,Chanyeong Lim

Original commits touching these paths (151 subject(s)):
  - Niko1221#152 as an opt-in (the owner's choice): STRATA_IQ_MT_MIN=1 makes CPU expert rounding independent of the drafting; the default stays 0.1.29's rule (IQ3_S decoded 1-3% slower with it on); docs; the width test runs the opt-in
  - Niko1221#289 follow-up: a placeholder GPU name when the properties query fails; HIP builds compile (no device_code_error probe)
  - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3
  - Niko1221#646 on 0.1.39b: the layer split and the CPU path as before, the resident plan for 8-token windows
  - Niko1221#646: the single-matrix mmvq does not stage the IQ grids (10-20% slower on the RTX 5070 at 3-8 columns)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested.
  - C-1, C-2: the prompt path's QSA select grid reaches only the active blocks; the indexer appends a chunk at once

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: src/kernels
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the
paths this fork still mirrors: 1 file(s) under src/ngram -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  src/ngram/ple_reader.cpp

Original authors: BlueKingMuch,Eli Elgaev,Konstantinos Korres,Niko1221

Original commits touching these paths (5 subject(s)):
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Helper threads leave the host thread's core (Linux)
  - PLE: the n-gram table in FP8 E4M3, as Qwen ships it
  - Strata: Qwen3.8-Flash-Next on a 12-24 GB NVIDIA GPU + 64 GB RAM, one-click install
  - ple: keep the SSD awake while rows are read (STRATA_SSD_KEEPALIVE)

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922
Fork-Replay-Branch: ac922
Fork-Replay-Base: fb58e0d
Fork-Replay-Tip: 1f029aa
Fork-Replay-Group: src/ngram
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
…glm53f

Squashed from 709 commit(s) 61b3fb5..7a33c9c on https://github.com/lighttransport/Strata, restricted to the
paths this fork still mirrors: 49 file(s) under include/strata -- 45 clean,
4 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  include/strata/hip_compat/cublas_v2.h
  include/strata/kernels/cpu/native_expert.hpp
  include/strata/kernels/iq_kernels.hpp
  include/strata/prefill/moe_mmq.hpp

Original authors: BlueKingMuch,Guillaume PUTIER,Konstantinos Korres,Léon Avic Simmons

Original commits touching these paths (65 subject(s)):
  - Niko1221#272 follow-up: --pool-affinity all (the pool's layout so far) is the default; auto and p-cores are opt-in
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - Add NUMA bandwidth benchmarks and native GLM EXL3 decode
  - Add guarded opt-in GLM prefill and bounded drafting experiments
  - Add opt-in GLM external drafting and guarded optimization records
  - Add opt-in GLM long-context prefill experiments
  - B550 decode levers: fused mHC read, direct tier DMA, MLA decode kernels, partial expert deferral
  - Coupled draft sampling for the MTP drafter (STRATA_SPEC_COUPLED=1, off by default)

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/lighttransport/Strata
Fork-Replay-Branch: glm53f
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: 7a33c9c
Fork-Replay-Group: include/strata
Fork-Replay-Automated: true
k0valik added a commit to k0valik/Strata that referenced this pull request Oct 15, 2026
Squashed from 709 commit(s) 61b3fb5..7a33c9c on https://github.com/lighttransport/Strata, restricted to the
paths this fork still mirrors: 1 file(s) under root -- 0 clean,
1 with conflict markers.

Conflict-marked files -- a pull request with conflicts: resolve them,
do not discard the work:
  CMakeLists.txt

Original authors: Andy,BlueKingMuch,Daehun,Daerdaal

Original commits touching these paths (138 subject(s)):
  - Niko1221#136: fused int8 tensor-core experts for the Q2_0 pack's prompt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1)
  - Niko1221#136: the fused int8 prompt experts for the native IQ packs (gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1)
  - Niko1221#290 follow-up: a BF16-row parity test (native_expert_parity --bf16-embd)
  - Niko1221#296 follow-up: only the Q5_0 MMQ instance and case stay (Q4_K / Q5_1 come with S5's STRATA_MMQ_KQUANTS)
  - Niko1221#296 follow-up: ple_q5_parity / q5_projection_parity build only with STRATA_BUILD_TESTS; the doc says what runs where
  - A fatal CUDA fault in the prompt path exits at once instead of hanging until the watchdog (Niko1221#224); CMake warns when sm_120 is built with CUDA < 13.0 (Niko1221#220, Niko1221#224); iq_pack writes native_experts.txt under a temporary name, so a stop part-way leaves no partial pack (Niko1221#172)
  - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base)
  - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested.

Review artifact produced by fork-replay.sh (a spike adapted from
sync-upstream.sh): a filtered, squashed replay for reading -- not built,
not tested, and not evidence that any of it works.

Fork-Replay-Source: https://github.com/lighttransport/Strata
Fork-Replay-Branch: glm53f
Fork-Replay-Base: 61b3fb5
Fork-Replay-Tip: 7a33c9c
Fork-Replay-Group: root
Fork-Replay-Automated: true
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants