Hacktoberfest 2026：メンテナが10月に向けて印を付けた、オープンで初心者向けの issue。 Hacktoberfest の issue を見る

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

オープン 初心者向け
  • 何か問題がありますか？
#4,589 コメント 1 件 リアクション 0 件 担当者 0 名 GitHub で見る

メンテナーはふだん 3 日以内に返信

まだ誰も着手していません。

評価

難易度
2/5
見積もり時間
1〜3時間
初心者へのやさしさ
76/100
issue の種類
バグ
明瞭さ
明確に書かれている
活発さ
活発
技術スタック
cpp
領域
backend

調査の方向性

src/kernels/cuda/iq_kernels.cuを読み、特に既存のマクロSTRATA_SHARED_ALIGN16_U32と、2571行目および2758行目にある2つのint4宣言を確認してください。ROCm上でsetup.sh --buildを使ってHIPソースをビルドし、問題を再現してから、このファイルがHIPとCUDAそれぞれで共有メモリのアラインメントをどのように処理しているかを確認してください。HIPビルドが両方の宣言を通過し、CUDAの動作も引き続きサポートされれば完了です。

索引モデルが issue の本文から書いたものです。

説明

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
  1. Clone main (46c152a1, v0.2.7.1) on a gfx1100 machine with system ROCm 7.1.1.
  2. 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/).
  3. 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).

主要言語
C++
スター
11.6k
フォーク
1k
30日以内にマージされた PR はありません

環境構築

このプロジェクトには開発コンテナ、Dockerfile、コントリビューションガイドがありません。まず README を読み、一般的な手順ははじめてのコントリビューションガイドを参照してください。

はじめの一歩

  1. issue を最後まで読み、次にプロジェクトのコントリビューションガイドを読みます。
  2. 着手することを issue にコメントします — 二人が同じ作業をするのを防げます。
  3. リポジトリをフォークし、ブランチを切って変更します。
  4. issue 番号を参照したプルリクエストを送ります。

Niko1221/Strata のほかの issue

Niko1221/Strata の issue をすべて見る

似ている issue

C++ の issue をもっと見る

新しい issue をメールで受け取る

初心者向けの GitHub issue を短くまとめたダイジェスト。