[libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify

[JonChesterfield]

Summary

[libomptarget][devicertl][nfc] Remove some cuda intrinsics, simplify

Replace popc, ffs with clang intrinsics. Move kmpc_impl_min to only file
that uses it and replace template with explictly typed.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

68

reordered to closer match nvptx

openmp/libomptarget/deviceRTLs/common/src/reduction.cu

187

Also considered 'min', but expect it to collide with one of the included headers, and 'min_uint32_t' but couldn't see a pretty way to write that inTheNamingConvention.

INLINE static is a slightly strange construct, chosen to match the functions above.

266

Three call sites, all uint32_t. Don't need a template.

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

95–96

These might be better changing typed as __kmpc_impl_lanemask_t, implemented in a common header that picks builtin_ffs or builtin_ffs based on sizeof(__kmpc_impl_lanemask_t). This keeps the current name mangling dispatch and separate implementations.

[tianshilei1992]

openmp/libomptarget/deviceRTLs/common/src/reduction.cu

187

The semantics of min(x, y) is same as x < y ? x : y but the implementation could be very different. For some data type (especially integers) we don't even need a branch here. I don't think it's good for performance to just use branch. If we know the type is uint32_t, then we can use bitwise operation here.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/common/src/reduction.cu

187

This will codegen as selec (https://godbolt.org/z/8fxPo7)

define dso_local i64 @min(i64 %0, i64 %1) local_unnamed_addr #0 {
  %3 = icmp ult i64 %0, %1
  %4 = select i1 %3, i64 %0, i64 %1
  ret i64 %4
}

which the backend will have an easier time with than decoding the equivalent bit shuffling.

[tianshilei1992]

openmp/libomptarget/deviceRTLs/common/src/reduction.cu

187

Okay. If this can solve the problem, __kmpc_impl_min can be removed as well.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/common/src/reduction.cu

187

Yep, __kmpc_impl_min is gone.

I've kept kmpcMin as a local function in preference to open coding x<y?x:y at each call site because it's shorter and fractionally easier to read than checking whether the argument order is min or max

[JonChesterfield]

• clang-format

[jdoerfert]

LGTM

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

 //===------- target_impl.h - AMDGCN OpenMP GPU implementation ----- HIP -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Declarations and definitions of target specific functions and constants
 //
 //===----------------------------------------------------------------------===//
 #ifndef OMPTARGET_AMDGCN_TARGET_IMPL_H
 #define OMPTARGET_AMDGCN_TARGET_IMPL_H
 
 #ifndef __AMDGCN__
 #error "amdgcn target_impl.h expects to be compiled under __AMDGCN__"
 #endif
 
 #include "amdgcn_interface.h"
 
 #include <assert.h>
 #include <inttypes.h>
 #include <stddef.h>
 #include <stdint.h>
enum DATA_SHARING_SIZES {
   // The size reserved for data in a shared memory slot.
   DS_Slot_Size = 256,
   // The slot size that should be reserved for a working warp.
   DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,
   // The maximum number of warps in use
   DS_Max_Warp_Number = 16,
 };
 
+enum : __kmpc_impl_lanemask_t {
+  __kmpc_impl_all_lanes = ~(__kmpc_impl_lanemask_t)0
+};
+
 INLINE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   lo = (uint32_t)(val & UINT64_C(0x00000000FFFFFFFF));
   hi = (uint32_t)((val & UINT64_C(0xFFFFFFFF00000000)) >> 32);
 }
 
 INLINE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   return (((uint64_t)hi) << 32) | (uint64_t)lo;
 }
 
-enum : __kmpc_impl_lanemask_t {
-  __kmpc_impl_all_lanes = ~(__kmpc_impl_lanemask_t)0
-};
-
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt();
-
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt();
-
 DEVICE uint32_t __kmpc_impl_smid();
-
 DEVICE double __kmpc_impl_get_wtick();
-
 DEVICE double __kmpc_impl_get_wtime();
 
 INLINE uint64_t __kmpc_impl_ffs(uint64_t x) { return __builtin_ffsl(x); }
-
 INLINE uint64_t __kmpc_impl_popc(uint64_t x) { return __builtin_popcountl(x); }
 
-template <typename T> INLINE T __kmpc_impl_min(T x, T y) {
-  return x < y ? x : y;
-}
-
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask();
 
 DEVICE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
                                      int32_t SrcLane);
 
 DEVICE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
                                           uint32_t Delta, int32_t Width);
 
 INLINE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
 
 INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t) {
   // AMDGCN doesn't need to sync threads in a warp
 }
 
 // AMDGCN specific kernel initialization
 DEVICE void __kmpc_impl_target_init();
 
 // Equivalent to ptx bar.sync 1. Barrier until num_threads arrive.
 DEVICE void __kmpc_impl_named_sync(uint32_t num_threads);
 
 INLINE void __kmpc_impl_threadfence() {
   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
 }
 
 INLINE void __kmpc_impl_threadfence_block() {
   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
 }
 
 INLINE void __kmpc_impl_threadfence_system() {
   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
 }
 
 // Calls to the AMDGCN layer (assuming 1D layout)
 INLINE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 INLINE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
 DEVICE int GetNumberOfBlocksInKernel();
 DEVICE int GetNumberOfThreadsInBlock();
 DEVICE unsigned GetWarpId();
 DEVICE unsigned GetLaneId();
 
 // Locks
 DEVICE void __kmpc_impl_init_lock(omp_lock_t *lock);
 DEVICE void __kmpc_impl_destroy_lock(omp_lock_t *lock);

openmp/libomptarget/deviceRTLs/common/src/reduction.cu

 }
 
 INLINE static uint32_t roundToWarpsize(uint32_t s) {
   if (s < WARPSIZE)
     return 1;
   return (s & ~(unsigned)(WARPSIZE - 1));
 }
 
+INLINE static uint32_t kmpcMin(uint32_t x, uint32_t y) { return x < y ? x : y; }
+
 DEVICE static volatile uint32_t IterCnt = 0;
 DEVICE static volatile uint32_t Cnt = 0;
 EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
     kmp_Ident *loc, int32_t global_tid, void *global_buffer,
     int32_t num_of_records, void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
     kmp_InterWarpCopyFctPtr cpyFct, kmp_ListGlobalFctPtr lgcpyFct,
     kmp_ListGlobalFctPtr lgredFct, kmp_ListGlobalFctPtr glcpyFct,
     kmp_ListGlobalFctPtr glredFct) {
EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
   // Each thread will have a local struct containing the values to be
   // reduced:
   //      1. do reduction within each warp.
   //      2. do reduction across warps.
   //      3. write the final result to the main reduction variable
   //         by returning 1 in the thread holding the reduction result.
 
   // Check if this is the very last team.
-  unsigned NumRecs = __kmpc_impl_min(NumTeams, uint32_t(num_of_records));
+  unsigned NumRecs = kmpcMin(NumTeams, uint32_t(num_of_records));
   if (ChunkTeamCount == NumTeams - Bound - 1) {
     //
     // Last team processing.
     //
     if (ThreadId >= NumRecs)
       return 0;
-    NumThreads = roundToWarpsize(__kmpc_impl_min(NumThreads, NumRecs));
+    NumThreads = roundToWarpsize(kmpcMin(NumThreads, NumRecs));
     if (ThreadId >= NumThreads)
       return 0;
 
     // Load from buffer and reduce.
     glcpyFct(global_buffer, ThreadId, reduce_data);
     for (uint32_t i = NumThreads + ThreadId; i < NumRecs; i += NumThreads)
       glredFct(global_buffer, i, reduce_data);
 
     // Reduce across warps to the warp master.
     if (NumThreads > 1) {
       gpu_regular_warp_reduce(reduce_data, shflFct);
 
       // When we have more than [warpsize] number of threads
       // a block reduction is performed here.
-      uint32_t ActiveThreads = __kmpc_impl_min(NumRecs, NumThreads);
+      uint32_t ActiveThreads = kmpcMin(NumRecs, NumThreads);
       if (ActiveThreads > WARPSIZE) {
         uint32_t WarpsNeeded = (ActiveThreads + WARPSIZE - 1) / WARPSIZE;
         // Gather all the reduced values from each warp
         // to the first warp.
         cpyFct(reduce_data, WarpsNeeded);
 
         uint32_t WarpId = ThreadId / WARPSIZE;
         if (WarpId == 0)

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

 //===------------ target_impl.h - NVPTX OpenMP GPU options ------- CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
 #ifndef _TARGET_IMPL_H_
 #define _TARGET_IMPL_H_
 
 #include <assert.h>
 #include <cuda.h>
 #include <inttypes.h>
 #include <stdio.h>
 #include <stdlib.h>
 
 #include "nvptx_interface.h"
 
 #define DEVICE __device__
 #define INLINE __forceinline__ DEVICE

 
 DEVICE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi);
 DEVICE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi);
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt();
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt();
 DEVICE uint32_t __kmpc_impl_smid();
 DEVICE double __kmpc_impl_get_wtick();
 DEVICE double __kmpc_impl_get_wtime();
 
-INLINE uint32_t __kmpc_impl_ffs(uint32_t x) { return __ffs(x); }
-
-INLINE uint32_t __kmpc_impl_popc(uint32_t x) { return __popc(x); }
-
-template <typename T> INLINE T __kmpc_impl_min(T x, T y) {
-  return min(x, y);
-}
+INLINE uint32_t __kmpc_impl_ffs(uint32_t x) { return __builtin_ffs(x); }
+INLINE uint32_t __kmpc_impl_popc(uint32_t x) { return __builtin_popcount(x); }
 
 #ifndef CUDA_VERSION
 #error CUDA_VERSION macro is undefined, something wrong with cuda.
 #endif
 
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask();
 
 DEVICE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t Mask, int32_t Var,