[Engine] HIP source build fails since v0.1.45: `shared alignas(16) int4` in iq_kernels.cu (expdec verify window)

Open Beginner friendly

#4,589 opened on Oct 29, 2026 1 comment 0 reactions 0 assignees View on GitHub

Maintainers usually reply within 3 days

@suvayu is already working on this.

Since Oct 29, 2026.

#4600 by @suvayu — open

Assessment

Difficulty: 2/5
Estimated time: 1-3 hours
Newbie friendliness: 76/100

Issue type: Bug
Clarity: Clearly specified
Activity status: Active
Tech stack: cpp
Domain: backend

Research direction

Read src/kernels/cuda/iq_kernels.cu, especially the existing STRATA_SHARED_ALIGN16_U32 macro and the two int4 declarations at lines 2571 and 2758. Reproduce with a HIP source build using setup.sh --build on ROCm, then check how the file handles shared-memory alignment for HIP and CUDA. Done when the HIP build compiles past both declarations while CUDA behavior remains supported.

Written by the indexing model from the issue text.

Description

Area

[Engine] (speed, kernels, models, the server and API)

Desktop app, Terminal or source?

Source (setup.sh --build on Linux, ROCm system install)

What happened

A source build of the AMD (HIP) engine fails. Compilation stops in src/kernels/cuda/iq_kernels.cu:

/…/src/kernels/cuda/iq_kernels.cu:2571:16: error: 'alignas' attribute cannot be applied to types
 2571 |     __shared__ alignas(16) int4 s_u[ED_NC * 2 * ED_XB];
      |                ^
/…/src/kernels/cuda/iq_kernels.cu:2758:16: error: 'alignas' attribute cannot be applied to types
 2758 |     __shared__ alignas(16) int4 s_u[ED_NC * 2 * ED_HB];
      |                ^
ninja: build stopped: subcommand failed.

I expected the engine to build as it did on v0.1.42; this is the expdec "Evaluate window" kernels.

The same file already documents this exact toolchain rule and carries a macro for it (line ~1869, #778):

// #778: `__shared__ alignas(16)` does not compile for gfx1030 (HIP); the attribute after the declarator does, but nvcc
// with MSVC as host does not take it.  Same layout either way.
#if defined(__HIPCC__)
#define STRATA_SHARED_ALIGN16_U32(name, ...) __shared__ uint32_t name[__VA_ARGS__] __attribute__((aligned(16)))
#else
#define STRATA_SHARED_ALIGN16_U32(name, ...) __shared__ alignas(16) uint32_t name[__VA_ARGS__]
#endif

The two int4 declarations added later use the bare __shared__ alignas(16) form and bypass that macro. CUDA is unaffected (nvcc accepts the form), so this is HIP-only. The declarations compile unconditionally even though STRATA_EXPDEC is a runtime opt-in, so it blocks every AMD source build.

Introduced by f7eaf2de ("Verify window: the routed experts' gate/up with the group's q8_1 rows in shared memory (STRATA_EXPDEC=1, opt-in, bitwise)…"); first tag that contains it is v0.1.45, so every release from there on. My previous clone was v0.1.42, which is why it built in October.

Smallest steps or command that reproduce it

Clone main (46c152a1, v0.2.7.1) on a gfx1100 machine with system ROCm 7.1.1.
ROCM_PATH=/opt/rocm-7.1.1 ./setup.sh --backend hip --family qwen --model IQ3_XXS --context 262144 --kv int8 --low-ram auto --vision no --draft-vocab en --port 8081 --data-dir … --build --yes --no-start
(or the equivalent cmake -DSTRATA_ENABLE_HIP=ON … && ninja strata in build-hip/).
The build fails at iq_kernels.cu, lines 2571 and 2758, as above.

Note: ./update.sh does not compile and only reports "the engine was NOT: it is still 0.1.42 and this Strata wants 0.1.46 or newer", so the compile is only reachable through setup.sh --build.

Reproduction status

Reproduced locally

Strata version

Strata 0.2.7.1 (46c152a1, v0.2.7.1) · engine 0.1.53 (built after the fix below)

Graphics card, memory and system

RX 7900 XTX 24 GB (gfx1100) · 126 GB RAM · Ubuntu 22.04.5, kernel 5.15 · system ROCm 7.1.1 (AMD clang 20.0.0git, hipBLASLt 1.1.0) · single card

Model, size and settings

Qwen3.8-Flash-Next IQ3_XXS · 262 144 context · --kv int8 · --spec 4 + MTP · --expert-cache auto · text only (no vision). (Only the build is affected; the model starts fine once the engine compiles.)

Logs

[68/164] Building HIP object CMakeFiles/strata_kernels.dir/src/kernels/cuda/kv_q4.cu.o
…
src/kernels/cuda/iq_kernels.cu:2571:16: error: 'alignas' attribute cannot be applied to types
src/kernels/cuda/iq_kernels.cu:2758:16: error: 'alignas' attribute cannot be applied to types
ninja: build stopped: subcommand failed.

Suggested fix

Give int4 the same split the file already has for uint32_t and use it at both sites:

+#if defined(__HIPCC__)
+#define STRATA_SHARED_ALIGN16_INT4(name, ...) __shared__ int4 name[__VA_ARGS__] __attribute__((aligned(16)))
+#else
+#define STRATA_SHARED_ALIGN16_INT4(name, ...) __shared__ alignas(16) int4 name[__VA_ARGS__]
+#endif
…
-    __shared__ alignas(16) int4 s_u[ED_NC * 2 * ED_XB];
+    STRATA_SHARED_ALIGN16_INT4(s_u, ED_NC * 2 * ED_XB);
…
-    __shared__ alignas(16) int4 s_u[ED_NC * 2 * ED_HB];
+    STRATA_SHARED_ALIGN16_INT4(s_u, ED_NC * 2 * ED_HB);

With that one file changed (+9/−2), the same command builds engine 0.1.53 and the model runs on this card (prefill ≈1500–1580 tok/s at 8.8 K, decode ≈80–88 tok/s, parity with v0.1.42 at the same settings).

What I did not test: whether a newer AMD clang accepts the original form. My check is ROCm 7.1.1 / clang 20 only; I have not built unpatched with ROCm 7.2.x (clang 22).

Dominant language: C++
Stars: 11.6k
Forks: 1k
PR merge metrics: 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

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

Similar issues

`elfbsp -h` missing some of the cli options
Open
- Is something wrong?
Difficulty 2/5 1-3 hours Newbie friendliness 72/100

elf-alchemist/elfbsp#113 · Opened Oct 31, 2026

Maintainers usually reply within 1 day
test(lab): banner OCR test assumes Debian font file locations
Open
- Is something wrong?
Difficulty 2/5 1-3 hours Newbie friendliness 75/100

patrick-hudson/put.gd-desktop#979 · 1 comment · Opened Oct 31, 2026

Maintainers usually reply within 1 day
# [Bug] [FA2] flash_attn_with_kvcache: the length of `cache_batch_idx` is not checked
Open
- Is something wrong?
Difficulty 2/5 1-3 hours Newbie friendliness 76/100

Dao-AILab/flash-attention#3097 · Opened Oct 31, 2026

Maintainers usually reply within 3 days
AFC backend: parse status before taking the lock
Open
- Is something wrong?
tech-debt

Difficulty 2/5 1-3 hours Newbie friendliness 82/100

prestonbrown/helixscreen#1990 · Opened Oct 31, 2026

Maintainers usually reply within 1 day
Minor doc comment oversight in CXXRTL CAPI regarding CXXRTL_ALIAS flag set on CXXRTL_ALIAS
Open
- Is something wrong?
Difficulty 1/5 Under an hour Newbie friendliness 90/100

YosysHQ/yosys#6432 · Opened Oct 31, 2026

Maintainers usually reply within 1 day

More C++ issues

[Engine] HIP source build fails since v0.1.45: `__shared__ alignas(16) int4` in iq_kernels.cu (expdec verify window)

Is something wrong?

Assessment

Research direction

Description

Area

Desktop app, Terminal or source?

What happened

Smallest steps or command that reproduce it

Reproduction status

Strata version

Graphics card, memory and system

Model, size and settings

Logs

Suggested fix

Getting set up

First steps

More from Niko1221/Strata

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Similar issues

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Is something wrong?

Get new issues in your inbox

[Engine] HIP source build fails since v0.1.45: `shared alignas(16) int4` in iq_kernels.cu (expdec verify window)