[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
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 @@
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) {}
