[libc-commits] [clang] [libc] [Clang] Add width handling for <gpuintrin.h> shuffle helper (PR #125896)

[Joseph Huber via libc-commits]

================
@@ -145,17 +145,21 @@ _DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_lane(uint64_t __lane_mask) {
 
 // Shuffles the the lanes inside the wavefront according to the given index.
 _DEFAULT_FN_ATTRS static __inline__ uint32_t
-__gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x) {
-  return __builtin_amdgcn_ds_bpermute(__idx << 2, __x);
+__gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x,
+                      uint32_t __width) {
+  uint32_t __lane = __idx + (__gpu_lane_id() & ~(__width - 1));
----------------
jhuber6 wrote:

I was thinking about how to document this, but I couldn't find anywhere in `clang` that did usage documentation for things like this. I suppose at a minimum I could put usage headers on everything, but for now I figured it'd best to just be consistent with `__shfl` and figured that was enough.

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