Hacktoberfest 2026: những issue maintainer đã đánh dấu cho tháng Mười, đang mở và phù hợp người mới. Xem issue Hacktoberfest

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

Đang mở Phù hợp với người mới
  • Có gì không ổn?
#4,589 1 bình luận 0 reaction 0 người được giao Xem trên GitHub

Maintainer thường phản hồi trong vòng 3 ngày

@suvayu đang làm issue này rồi.

Từ ngày 29/10/2026.

  • #4600 của @suvayu — đang mở

Đánh giá

Độ khó
2/5
Thời gian dự kiến
1-3 giờ
Mức phù hợp với người mới
76/100
Loại issue
Lỗi
Độ rõ ràng
Đặc tả rõ ràng
Mức độ hoạt động
Sôi nổi
Công nghệ
cpp
Lĩnh vực
backend

Hướng nghiên cứu

Đọc src/kernels/cuda/iq_kernels.cu, đặc biệt là macro hiện có STRATA_SHARED_ALIGN16_U32 và hai khai báo int4 ở dòng 2571 và 2758. Tái hiện lỗi bằng cách xây dựng mã nguồn HIP với setup.sh --build trên ROCm, sau đó kiểm tra cách tệp xử lý việc căn chỉnh bộ nhớ dùng chung cho HIP và CUDA. Hoàn thành khi bản dựng HIP biên dịch qua cả hai khai báo và hành vi của CUDA vẫn được hỗ trợ.

Do mô hình lập chỉ mục viết ra từ nội dung của issue.

Mô tả

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).

Ngôn ngữ chính
C++
Star
11.6k
Fork
1k
Không có pull request nào được merge trong 30 ngày

Chuẩn bị môi trường

Dự án này không cung cấp dev container, Dockerfile hay hướng dẫn đóng góp, nên bạn cần tự thiết lập môi trường: hãy bắt đầu từ README và xem hướng dẫn đóng góp lần đầu của chúng tôi để biết các bước chung.

Bắt đầu từ đâu

  1. Đọc hết issue, rồi đọc hướng dẫn đóng góp của dự án.
  2. Bình luận trên issue rằng bạn sẽ nhận — tránh hai người làm cùng một việc.
  3. Fork repository và làm thay đổi trên một nhánh.
  4. Mở pull request có tham chiếu số hiệu của issue.

Issue khác của Niko1221/Strata

Tất cả issue của Niko1221/Strata

Issue tương tự

Thêm issue về C++

Nhận issue mới trong hộp thư của bạn

Bản tóm tắt ngắn những issue GitHub phù hợp với người mới.