[libomptarget-nvptx] Remove dead functions

[Hahnfeld]

Summary

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

[Hahnfeld]

Ping

[gtbercea]

LGTM since neither Clang nor clang-ykt use these functions.

I'd wait for someone from the XL team to confirm XL doesn't need them either.

Thanks!

[kkwli0]

The XL compilers do not use these routines. Thanks.

[jdoerfert]

+2

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/interface.h

 //===------- interface.h - NVPTX OpenMP interface definitions ---- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains debug macros to be used in the application.
 //
 //  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.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _INTERFACES_H_
 #define _INTERFACES_H_
 
 #include "option.h"

 typedef struct kmp_TaskDescr kmp_TaskDescr;
 typedef int32_t (*kmp_TaskFctPtr)(int32_t global_tid, kmp_TaskDescr *taskDescr);
 typedef struct kmp_TaskDescr {
   void *sharedPointerTable;   // ptr to a table of shared var ptrs
   kmp_TaskFctPtr sub;         // task subroutine
   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,
                                     int32_t num_threads);
 // simd
 EXTERN void __kmpc_push_simd_limit(kmp_Ident *loc, int32_t global_tid,
                                    int32_t simd_limit);
 // aee ... not supported

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/reduction.cu

 //
 //===----------------------------------------------------------------------===//
 
 #include <complex.h>
 #include <stdio.h>
 
 #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) {}
 
 EXTERN
 void __kmpc_nvptx_end_reduce_nowait(int32_t global_tid) {}
 
 EXTERN int32_t __kmpc_shuffle_int32(int32_t val, int16_t delta, int16_t size) {
   return __SHFL_DOWN_SYNC(0xFFFFFFFF, val, delta, size);