Hacktoberfest 2026: the issues maintainers tagged for October, open and beginner-friendly. Browse Hacktoberfest issues

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

Open Beginner friendly
  • Is something wrong?
#2,374 0 comments 0 reactions 0 assignees View on GitHub

Maintainers usually reply within 1 day

Nobody has claimed this yet.

Assessment

Difficulty
2/5
Estimated time
1-3 hours
Newbie friendliness
66/100
Issue type
Bug
Clarity
Clearly specified
Activity status
Active
Tech stack
cpp
Domain
backend

Research direction

Start with the #if guard in dq_q2_0 in src/kernels/cuda/iq_kernels.cu, which lists the architectures that fold signed zero. The failing test is hip_q2_zero, and the one-line change adds __gfx1011__ to that list. Also fix the Navi 12 wording in docs/OLDER_GPUS.md. Done means hip_q2_zero passes on gfx1011 hardware, which needs a gfx1011 card to confirm; whether gfx1012 also needs the guard is still open.

Written by the indexing model from the issue text.

Description

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.

Dominant language
C++
Stars
11.6k
Forks
1k
No merged PRs in 30d

Getting set up

This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

More from Niko1221/Strata

All issues in Niko1221/Strata

Similar issues

More C++ issues

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.