Pull requests / #3429
#3429 [Engine] HIP: the expdec shared int4 arrays use the attribute-after-declarator form
open · @chenxyz129 · 0 Kommentare · Auf GitHub
Setup & installAMD / HIPNVIDIA / CUDA
Beschreibung
## Problem and result **Before:** on `main` (9ae93f7b, Control 0.1.46.1 / engine 0.1.46), a HIP build for gfx1032 stops in `src/kernels/cuda/iq_kernels.cu`: ```text src/kernels/cuda/iq_kernels.cu:2263:16: error: 'alignas' attribute cannot be applied to types src/kernels/cuda/iq_kernels.cu:2450:16: error: 'alignas' attribute cannot be applied to types ``` Those are the two expdec kernels added in v0.1.45 (gate/up and down), which declare `__shared__ alignas(16) int4 s_u[...]`. The same file already carries the working form right next to them: `STRATA_SHARED_ALIGN16_U32`, with the #778 comment saying that HIP takes the attribute after the declarator. **After:** a `STRATA_SHARED_ALIGN16_I4` macro mirrors that one (`__shared__ int4 name[...] __attribute__((aligned(16)))` under `__HIPCC__`, `alignas` elsewhere), and the two kernels use it. Array type, size and alignment are unchanged. ## Related discussion none ## Approach One macro beside the existing `STRATA_SHARED_ALIGN16_U32`, and the two declarations switched to it. No kernel logic, no launch configuration and no numerics change: this only makes the expdec source compile for HIP again. ## Validation **Passed:** - After the change, at commit `f1e9d33a` (main `9ae93f7b` + this patch), `cmake --build build-hip -j10` for gfx1032 -> all targets built and `strata` linked; `./build-hip/strata --version` -> `strata 0.1.46.1`. Environment: Fedora, ROCm 7.1.1 (hipcc 7.1.52802), AMD RX 6650 XT (gfx1032, 8 GB). - `./build-hip/strata-device --selftest` at the same commit -> `strata-device selftest OK` (arena allocation and the over-allocation refusal). - The same build also passes with this patch on top of the gfx1032 setup branch (`3eff80b1`): all targets built, `strata-device --selftest` -> `strata-device selftest OK`. **Failed:** - Before the change, at `9ae93f7b`, the same `cmake --build build-hip -j10` for gfx1032 fails with the two `'alignas' attribute cannot be applied to types` errors quoted above. That is the problem this fixes. **Not run:** - gfx1030 hardware: same shared-memory declaration and same HIP compiler path, but no such card here. - CUDA, SYCL and Metal builds. - A full model run, llama.cpp parity and any timing: this is a compile fix, and no number here is a speed claim. ## Limits - Tested with one HIP compiler version (7.1.52802) on one gfx1032 card; other HIP targets compile the same macro but were not built here. - Nothing about the expdec kernels' speed or results is measured in this pull request.
Mehr auf der Site
Links zu Install, Modellen, Releases.