[Openmp-commits] [llvm] [openmp] [OpenMP][OMPX] Add ballot_sync (PR #91297)

Joseph Huber via Openmp-commits openmp-commits at lists.llvm.org
Thu May 23 07:57:57 PDT 2024


================
@@ -57,6 +59,12 @@ int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
 }
 
+uint64_t ballotSync(uint64_t Mask, int32_t Pred) {
+  return Mask &
+         (__AMDGCN_WAVEFRONT_SIZE == 32 ? __builtin_amdgcn_ballot_w32(Pred)
+                                        : __builtin_amdgcn_ballot_w64(Pred));
----------------
jhuber6 wrote:

```suggestion
  return Mask & __builtin_amdgcn_ballot_w64(Pred));
```
The `w64` intrinsic works on both targets, the backend lowers it correctly. See https://godbolt.org/z/Ph6sn3Mh6.

https://github.com/llvm/llvm-project/pull/91297


More information about the Openmp-commits mailing list