[libomptarget-nvptx] Remove dead functions

These entry points are never called by Clang trunk nor clang-ykt. If
XL doesn't use them either, they can finally go away.

Differential Revision: https://reviews.llvm.org/D52700

git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@365817 91177308-0d34-0410-b5e6-96231b3b80d8

2 files changed, 1 insertions, 85 deletions

diff --git a/libomptarget/deviceRTLs/nvptx/src/interface.h b/libomptarget/deviceRTLs/nvptx/src/interface.h
index 534ed62..b2a13a4 100644
--- a/libomptarget/deviceRTLs/nvptx/src/interface.h
+++ b/libomptarget/deviceRTLs/nvptx/src/interface.h
@@ -10,8 +10,7 @@
 //
 //  This file contains all the definitions that are relevant to
 //  the interface. The first section contains the interface as
-//  declared by OpenMP.  A second section includes library private calls
-//  (mostly debug, temporary?) The third section includes the compiler
+//  declared by OpenMP.  The second section includes the compiler
 //  specific interfaces.
 //
 //===----------------------------------------------------------------------===//
@@ -211,51 +210,14 @@ typedef struct kmp_TaskDescr {
   int32_t partId;             // unused
   kmp_TaskFctPtr destructors; // destructor of c++ first private
 } kmp_TaskDescr;
-// task dep defs
-#define KMP_TASKDEP_IN 0x1u
-#define KMP_TASKDEP_OUT 0x2u
-typedef struct kmp_TaskDep_Public {
-  void *addr;
-  size_t len;
-  uint8_t flags; // bit 0: in, bit 1: out
-} kmp_TaskDep_Public;
-
-// flags that interpret the interface part of tasking flags
-#define KMP_TASK_IS_TIED 0x1
-#define KMP_TASK_FINAL 0x2
-#define KMP_TASK_MERGED_IF0 0x4 /* unused */
-#define KMP_TASK_DESTRUCTOR_THUNK 0x8
-
-// flags for task setup return
-#define KMP_CURRENT_TASK_NOT_SUSPENDED 0
-#define KMP_CURRENT_TASK_SUSPENDED 1
 
 // sync defs
 typedef int32_t kmp_CriticalName[8];
 
 ////////////////////////////////////////////////////////////////////////////////
-// flags for kstate (all bits initially off)
-////////////////////////////////////////////////////////////////////////////////
-
-// first 2 bits used by kmp_Reduction (defined in kmp_reduction.cpp)
-#define KMP_REDUCTION_MASK 0x3
-#define KMP_SKIP_NEXT_CALL 0x4
-#define KMP_SKIP_NEXT_CANCEL_BARRIER 0x8
-
-////////////////////////////////////////////////////////////////////////////////
-// data
-////////////////////////////////////////////////////////////////////////////////
-
-////////////////////////////////////////////////////////////////////////////////
 // external interface
 ////////////////////////////////////////////////////////////////////////////////
 
-// query
-EXTERN int32_t __kmpc_global_num_threads(kmp_Ident *loc); // missing
-EXTERN int32_t __kmpc_bound_thread_num(kmp_Ident *loc);   // missing
-EXTERN int32_t __kmpc_bound_num_threads(kmp_Ident *loc);  // missing
-EXTERN int32_t __kmpc_in_parallel(kmp_Ident *loc);        // missing
-
 // parallel
 EXTERN int32_t __kmpc_global_thread_num(kmp_Ident *loc);
 EXTERN void __kmpc_push_num_threads(kmp_Ident *loc, int32_t global_tid,
diff --git a/libomptarget/deviceRTLs/nvptx/src/reduction.cu b/libomptarget/deviceRTLs/nvptx/src/reduction.cu
index 8d60381..c925638 100644
--- a/libomptarget/deviceRTLs/nvptx/src/reduction.cu
+++ b/libomptarget/deviceRTLs/nvptx/src/reduction.cu
@@ -16,52 +16,6 @@
 
 #include "omptarget-nvptx.h"
 
-// may eventually remove this
-EXTERN
-int32_t __gpu_block_reduce() {
-  bool isSPMDExecutionMode = isSPMDMode();
-  int nt = GetNumberOfOmpThreads(isSPMDExecutionMode);
-  if (nt != blockDim.x)
-    return 0;
-  unsigned tnum = __ACTIVEMASK();
-  if (tnum != (~0x0)) // assume swapSize is 32
-    return 0;
-  return 1;
-}
-
-EXTERN
-int32_t __kmpc_reduce_gpu(kmp_Ident *loc, int32_t global_tid, int32_t num_vars,
-                          size_t reduce_size, void *reduce_data,
-                          void *reduce_array_size, kmp_ReductFctPtr *reductFct,
-                          kmp_CriticalName *lck) {
-  int threadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
-  omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(threadId);
-  int numthread;
-  if (currTaskDescr->IsParallelConstruct()) {
-    numthread = GetNumberOfOmpThreads(checkSPMDMode(loc));
-  } else {
-    numthread = GetNumberOfOmpTeams();
-  }
-
-  if (numthread == 1)
-    return 1;
-  if (!__gpu_block_reduce())
-    return 2;
-  if (threadIdx.x == 0)
-    return 1;
-  return 0;
-}
-
-EXTERN
-int32_t __kmpc_reduce_combined(kmp_Ident *loc) {
-  return threadIdx.x == 0 ? 2 : 0;
-}
-
-EXTERN
-int32_t __kmpc_reduce_simd(kmp_Ident *loc) {
-  return (threadIdx.x % 32 == 0) ? 1 : 0;
-}
-
 EXTERN
 void __kmpc_nvptx_end_reduce(int32_t global_tid) {}