Fix build, add shuffle

Author: jdoerfert

llvm/lib/Transforms/Instrumentation/GPUSan.cpp

@@ -618,7 +618,6 @@ void GPUSanImpl::instrumentAccess(LoopInfo &LI, Instruction &I, int PtrIdx,
   if (Loop *L = LI.getLoopFor(I.getParent())) {
     auto &SE = FAM.getResult<ScalarEvolutionAnalysis>(*I.getFunction());
     const auto &LD = SE.getLoopDisposition(SE.getSCEVAtScope(PtrOp, L), L);
-    LD->
   }
 
   static int32_t ReadAccessId = -1;

offload/DeviceRTL/include/DeviceUtils.h

@@ -22,7 +22,7 @@ namespace utils {
 
 /// Return the value \p Var from thread Id \p SrcLane in the warp if the thread
 /// is identified by \p Mask.
-int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane);
+int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane, int32_t Width);
 
 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta, int32_t Width);
 

offload/DeviceRTL/src/Mapping.cpp

@@ -376,6 +376,18 @@ uint64_t ompx_ballot_sync(uint64_t mask, int pred) {
   return utils::ballotSync(mask, pred);
 }
 
+#define _TGT_KERNEL_LANGUAGE_SHFL_SYNC(TYPE, TY)                               \
+  int ompx_shfl_sync_##TY(uint64_t mask, TYPE var, int src, int width) {       \
+    return utils::shuffle(mask, var, src, width);                              \
+  }
+
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(int, i)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(float, f)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(long, l)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(double, d)
+
+#undef _TGT_KERNEL_LANGUAGE_SHFL_SYNC
+
 int ompx_shfl_down_sync_i(uint64_t mask, int var, unsigned delta, int width) {
   return utils::shuffleDown(mask, var, delta, width);
 }

offload/DeviceRTL/src/Utils.cpp

@@ -34,7 +34,7 @@ uint64_t Pack(uint32_t LowBits, uint32_t HighBits) {
   return (((uint64_t)HighBits) << 32) | (uint64_t)LowBits;
 }
 
-int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane);
+int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane, int32_t Width);
 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
                     int32_t Width);
 
@@ -45,8 +45,7 @@ uint64_t ballotSync(uint64_t Mask, int32_t Pred);
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
-int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
-  int Width = mapping::getWarpSize();
+int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane, int32_t Width) {
   int Self = mapping::getThreadIdInWarp();
   int Index = SrcLane + (Self & ~(Width - 1));
   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
@@ -82,8 +81,8 @@ bool isThreadLocalMemPtr(const void *Ptr) {
         device = {arch(nvptx, nvptx64)},                                       \
             implementation = {extension(match_any)})
 
-int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
-  return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f);
+int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane, int32_t Width) {
+  return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, Width);
 }
 
 int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta, int32_t Width) {
@@ -111,8 +110,9 @@ void utils::unpack(uint64_t Val, uint32_t &LowBits, uint32_t &HighBits) {
   impl::Unpack(Val, &LowBits, &HighBits);
 }
 
-int32_t utils::shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane) {
-  return impl::shuffle(Mask, Var, SrcLane);
+int32_t utils::shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane,
+                       int32_t Width) {
+  return impl::shuffle(Mask, Var, SrcLane, Width);
 }
 
 int32_t utils::shuffleDown(uint64_t Mask, int32_t Var, uint32_t Delta,

offload/DeviceRTL/src/Workshare.cpp

@@ -348,7 +348,7 @@ template <typename T, typename ST> struct omptarget_nvptx_LoopSupport {
     if (rank == 0) {
       warp_res = atomic::add(&Cnt, change, atomic::seq_cst);
     }
-    warp_res = utils::shuffle(active, warp_res, leader);
+    warp_res = utils::shuffle(active, warp_res, leader, mapping::getWarpSize());
     return warp_res + rank;
   }
 

openmp/runtime/src/include/ompx.h.var

@@ -139,6 +139,20 @@ _TGT_KERNEL_LANGUAGE_DECL_GRID_C(grid_dim)
 
 uint64_t ompx_ballot_sync(uint64_t mask, int pred);
 
+/// ompx_shfl_sync_{i,f,l,d}
+///{
+#define _TGT_KERNEL_LANGUAGE_SHFL_SYNC(TYPE, TY)                          \
+  TYPE ompx_shfl_sync_##TY(uint64_t mask, TYPE var, int src,                   \
+                            int width);
+
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(int, i)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(float, f)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(long, l)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(double, d)
+
+#undef _TGT_KERNEL_LANGUAGE_SHFL_SYNC
+///}
+
 /// ompx_shfl_down_sync_{i,f,l,d}
 ///{
 #define _TGT_KERNEL_LANGUAGE_SHFL_DOWN_SYNC(TYPE, TY)                          \
@@ -208,6 +222,22 @@ static inline uint64_t ballot_sync(uint64_t mask, int pred) {
   return ompx_ballot_sync(mask, pred);
 }
 
+/// shfl_sync
+///{
+#define _TGT_KERNEL_LANGUAGE_SHFL_SYNC(TYPE, TY)                          \
+  static inline TYPE shfl_sync(uint64_t mask, TYPE var, int src,               \
+                                    int width = __WARP_SIZE) {                 \
+    return ompx_shfl_down_sync_##TY(mask, var, int, width);                  \
+  }
+
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(int, i)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(float, f)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(long, l)
+_TGT_KERNEL_LANGUAGE_SHFL_SYNC(double, d)
+
+#undef _TGT_KERNEL_LANGUAGE_SHFL_DOWN_SYNC
+///}
+
 /// shfl_down_sync
 ///{
 #define _TGT_KERNEL_LANGUAGE_SHFL_DOWN_SYNC(TYPE, TY)                          \