Pull requests / #3429

#3429 [Engine] HIP: the expdec shared int4 arrays use the attribute-after-declarator form

open · @chenxyz129 · 0 comentarios · En GitHub

Setup & installAMD / HIPNVIDIA / CUDA

Descripción

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

En el sitio

Enlaces a install, modelos, releases.