Index: libomptarget/deviceRTLs/nvptx/src/interface.h =================================================================== --- libomptarget/deviceRTLs/nvptx/src/interface.h +++ libomptarget/deviceRTLs/nvptx/src/interface.h @@ -11,8 +11,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. // //===----------------------------------------------------------------------===// @@ -183,51 +182,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_Indent *loc); // missing -EXTERN int32_t __kmpc_bound_thread_num(kmp_Indent *loc); // missing -EXTERN int32_t __kmpc_bound_num_threads(kmp_Indent *loc); // missing -EXTERN int32_t __kmpc_in_parallel(kmp_Indent *loc); // missing - // parallel EXTERN int32_t __kmpc_global_thread_num(kmp_Indent *loc); EXTERN void __kmpc_push_num_threads(kmp_Indent *loc, int32_t global_tid, Index: libomptarget/deviceRTLs/nvptx/src/reduction.cu =================================================================== --- libomptarget/deviceRTLs/nvptx/src/reduction.cu +++ libomptarget/deviceRTLs/nvptx/src/reduction.cu @@ -17,53 +17,6 @@ #include "omptarget-nvptx.h" -// may eventually remove this -EXTERN -int32_t __gpu_block_reduce() { - int tid = GetLogicalThreadIdInBlock(); - int nt = GetNumberOfOmpThreads(tid, isSPMDMode(), isRuntimeUninitialized()); - 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_Indent *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(); - omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(threadId); - int numthread; - if (currTaskDescr->IsParallelConstruct()) { - numthread = - GetNumberOfOmpThreads(threadId, isSPMDMode(), isRuntimeUninitialized()); - } 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_Indent *loc) { - return threadIdx.x == 0 ? 2 : 0; -} - -EXTERN -int32_t __kmpc_reduce_simd(kmp_Indent *loc) { - return (threadIdx.x % 32 == 0) ? 1 : 0; -} - EXTERN void __kmpc_nvptx_end_reduce(int32_t global_tid) {}