[OpenMP][FIX] Repair accidental replacement of _shfl_sync with _shfl

[jdoerfert]

Summary

This was broken accidentally in D95752.

[ye-luo]

Fixed the problem in my test.

openmp/libomptarget/deviceRTLs/common/include/target/shuffle.h

 
 ///}
 
 /// AMDGCN implementations of the shuffle sync idiom.
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 inline int32_t __kmpc_impl_shfl_sync(uint64_t Mask, int32_t Var,
                                      int32_t SrcLane) {
   int Width = GetWarpSize();
   int Self = GetLaneId();
   int Index = SrcLane + (Self & ~(Width - 1));
   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
 }
 
 inline int32_t __kmpc_impl_shfl_down_sync(uint64_t Mask, int32_t Var,
                                           uint32_t LaneDelta, int32_t Width) {
   int Self = GetLaneId();
   int Index = Self + LaneDelta;
   Index = (int)(LaneDelta + (Self & (Width - 1))) >= Width ? Self : Index;
   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
 }
 
 #pragma omp end declare variant
 ///}
 
 /// NVPTX implementations of the shuffle and shuffle sync idiom.
 ///
 ///{
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
 inline int32_t __kmpc_impl_shfl_sync(uint64_t Mask, int32_t Var,
                                      int32_t SrcLane) {
-  return __nvvm_shfl_idx_i32(Var, SrcLane, 0x1f);
+  return __nvvm_shfl_sync_idx_i32(Mask, Var, SrcLane, 0x1f);
 }
 
 inline int32_t __kmpc_impl_shfl_down_sync(uint64_t Mask, int32_t Var,
                                           uint32_t Delta, int32_t Width) {
   int32_t T = ((GetWarpSize() - Width) << 8) | 0x1f;
-  return __nvvm_shfl_down_i32(Var, Delta, T);
+  return __nvvm_shfl_sync_down_i32(Mask, Var, Delta, T);
 }
 
 #pragma omp end declare variant
 ///}
 
 #pragma omp end declare target
 
 #endif