addmisol wrote:
> @addmisol -- friendly heads-up: we're tracking a regression that bisects
> cleanly to this PR.
>
> In HIP code on `amdgcn` targets, FP16 wavefront reductions
> (`__reduce_{min,max}_sync<__half>`) return wrong values after this change.
> The root cause appears to be: HIP's `__half` (`struct { union { _Float16 }
> }`) is now classified by `ABIArgInfo::getDirect()` as an aggregate, while
> pre-built `rocm-device-libs` still exposes `__ockl_wfred_*_f16` with a flat
> `half` parameter. The resulting call-site / callee signature mismatch blocks
> inlining of a `convergent` wave-reduction helper, which then runs under a
> divergent-call EXEC mask and reads stale data from neighbour lanes via DPP.
>
> Reproducer is a single-file HIP program at `-O2`, ~2 s end-to-end; reverting
> this PR on top of the otherwise-identical tree makes it pass. Confirmed on
> MI200/MI210 (gfx90a).
>
> Are you aware of this, and do you have a preference on how to address it? A
> couple of options we've considered:
>
> * narrow the new classification in `clang/lib/CodeGen/Targets/AMDGPU.cpp` to
> keep the legacy `[N x i32]` coercion for small aggregates whose only data is
> `_Float16` (and `__bf16`), so the calling convention stays compatible with
> already-shipped device-libs;
> * or accept that this PR requires a coordinated `rocm-device-libs` rebuild
> and document that expectation.
>
> Happy to share more (full IR/asm diff, bisect trail, etc.) or to test a
> candidate fix. Thanks!
Hi @michaelselehov,
Sure, please revert this if needed.. i'll investigate further.
https://github.com/llvm/llvm-project/pull/185083
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
