Hacktoberfest 2026: le issue che i maintainer hanno segnato per ottobre, aperte e adatte ai principianti. Sfoglia le issue Hacktoberfest

gfx1011 (Navi 12) on real hardware: 96/105 ctest pass; `hip_q2_zero` fails, fixed by adding `__gfx1011__` to the Q2_0 signed-zero workaround

Aperta Adatta ai principianti
  • Qualcosa non va?
#2,374 0 commenti 0 reazioni 0 assegnatari Vedi su GitHub

I maintainer di solito rispondono entro 1 giorno

Nessuno ha ancora preso questa issue.

Valutazione

Difficoltà
2/5
Tempo stimato
1-3 ore
Idoneità per principianti
66/100
Tipo di issue
Bug
Chiarezza
Specificata chiaramente
Stato di attività
Attiva
Stack tecnologico
cpp
Ambito
backend

Direzione di ricerca

Inizia dalla guardia #if in dq_q2_0 in src/kernels/cuda/iq_kernels.cu, che elenca le architetture che unificano lo zero con segno. Il test che fallisce è hip_q2_zero, e la modifica di una riga aggiunge __gfx1011__ a quell'elenco. Correggi anche la formulazione su Navi 12 in docs/OLDER_GPUS.md. Il lavoro è completato quando hip_q2_zero passa su hardware gfx1011, il che richiede una scheda gfx1011 per la conferma. Se anche gfx1012 necessiti della guardia resta una questione aperta.

Scritto dal modello di indicizzazione a partire dal testo della issue.

Descrizione

docs/OLDER_GPUS.md lists gfx1011 as "untested on hardware". Here is a hardware run.

Hardware / software

  • 3× AMD BC-160 (Navi 12, gfx1011, 8 GB HBM2 each); host i5-6500 (4 cores, no AVX-512), 16 GB RAM
  • Ubuntu, kernel 7.0.0-38-generic (in-tree amdgpu)
  • Strata v0.1.43.1, commit a0116b2
  • ROCm from TheRock wheels: rocm[libraries,devel]==7.14.0a20260608, index https://rocm.nightlies.amd.com/v2/gfx101X-dgpu/
  • configure: -DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF -DSTRATA_BUILD_TESTS=ON -DSTRATA_PREFILL_MMQ=ON -DCMAKE_HIP_ARCHITECTURES=gfx1011

Results

  • Build: 0 errors.
  • strata-device: HIP arch gfx1011 wave32 (compiled for gfx1011), 18 CUs, 7.984 GiB VRAM.
  • ctest, run separately on each card (HIP_VISIBLE_DEVICES=0/1/2), with ple_parity and platform_memory_test excluded: identical on all three cards — 96 passed, 8 skipped, 1 failed out of 105, ~170–180 s per card. No amdgpu errors in the kernel log.
  • Skipped (expected for this hardware): hip_prefill_wmma_gemm_parity, hip_prefill_hcd_exact_parity, hip_prompt_attn_wmma (no WMMA on RDNA1), hip_prefill_hipblaslt_gemm, hip_prefill_fused_moe, hip_prefill_fused_iq, expert_cache_segmented_test, expert_multi_test.
  • Failed: hip_q2_zero: 
    index=65 got=0x0000 expected=0x8000
...
hip_q2_zero: 1024 values, 128 failures

Cause and fix

This is the same signed-zero folding the guard in dq_q2_0 (src/kernels/cuda/iq_kernels.cu) already works around for gfx1012 on HIP < 7, gfx1151 and gfx1200/gfx1201. gfx1011 is not in that list, and with ROCm 7.14 it shows the same behaviour. Adding it fixes the test on all three cards (1024 values, 0 failures):

-#if defined(__HIPCC__) && ((defined(__gfx1012__) && HIP_VERSION_MAJOR < 7) || defined(__gfx1151__) || \
+#if defined(__HIPCC__) && ((defined(__gfx1012__) && HIP_VERSION_MAJOR < 7) || defined(__gfx1011__) || defined(__gfx1151__) || \
     defined(__gfx1200__) || defined(__gfx1201__))

Since gfx1011 still folds on HIP 7.14, gfx1012 on HIP ≥ 7 may need the guard too. I have no gfx1012 card to check.

Not tested: running a model. This box has 16 GB of RAM, below the documented minimum. The memory planner (--max-context 20480) reports a 4105-slot expert cache (5.68 GB) plus 0.27 GB KV, which fits a single 8 GB card.

Small doc note: OLDER_GPUS.md says gfx1011 is "the same RDNA1 ISA as gfx1010 (Navi 12)". Navi 12 is gfx1011 itself; gfx1010 is Navi 10.

Lingua principale
C++
Stelle
11.6k
Fork
1k
Nessuna PR unita negli ultimi 30g

Preparare l'ambiente

Questo progetto non fornisce container di sviluppo, Dockerfile né guida per i contributori, quindi l'ambiente è a tuo carico: parti dal suo README e consulta la nostra guida al primo contributo per i passaggi generali.

Come iniziare

  1. Leggi tutta la issue e poi la guida ai contributi del progetto.
  2. Commenta sulla issue per dire che te ne occupi tu — evita che due persone facciano lo stesso lavoro.
  3. Fai un fork del repository e lavora su un branch.
  4. Apri una pull request che faccia riferimento al numero della issue.

Altre issue di Niko1221/Strata

Tutte le issue di Niko1221/Strata

Issue simili

Altre issue su C++

Ricevi le nuove issue nella tua casella

Un breve riepilogo di issue GitHub adatte ai principianti.