[libomptarget][NVPTX] Drop dead code and data structures, NFCI.

[Hahnfeld]

Summary

• cg and HasCancel in WorkDescr were never read and can be removed.
• This eliminates the last use of priv in ThreadPrivateContext.
• CounterGroup is unused afterwards.
• Remove duplicate external declares in omptarget-nvptx.cu that are already in the header omptarget-nvptx.h.

[ABataev]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

24–42 ↗

(On Diff #163791)

Are you sure that you want to delete these variables?

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

24–42 ↗

(On Diff #163791)

Yes, they are also declared in omptarget-nvptx.h (I added an inline comment). I've been criticizing these duplicated declarations in a few reviews, time to get rid of them. Note that omptarget_nvptx_device_State and omptarget_nvptx_device_simpleState are not in the header because they are only used in this file.

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

431–446 ↗

(On Diff #163791)

Here.

[ABataev]

LG

[grokos]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptxi.h

182 ↗

(On Diff #163791)

Also remove this dead code?

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptxi.h

182 ↗

(On Diff #163791)

Yes, will do on commit.

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

-//===------ counter_group.h - NVPTX OpenMP loop scheduling ------- CUDA -*-===//
-//
-//                     The LLVM Compiler Infrastructure
-//
-// This file is dual licensed under the MIT and the University of Illinois Open
-// Source Licenses. See LICENSE.txt for details.
-//
-//===----------------------------------------------------------------------===//
-//
-// Interface to implement OpenMP loop scheduling
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef _OMPTARGET_NVPTX_COUNTER_GROUP_H_
-#define _OMPTARGET_NVPTX_COUNTER_GROUP_H_
-
-#include "option.h"
-
-// counter group type for synchronizations
-class omptarget_nvptx_CounterGroup {
-public:
-  // getters and setters
-  INLINE Counter &Event() { return v_event; }
-  INLINE volatile Counter &Start() { return v_start; }
-  INLINE Counter &Init() { return v_init; }
-
-  // Synchronization Interface
-
-  INLINE void Clear();             // first time start=event
-  INLINE void Reset();             // init = first
-  INLINE void Init(Counter &priv); // priv = init
-  INLINE Counter Next();           // just counts number of events
-
-  // set priv to n, to be used in later waitOrRelease
-  INLINE void Complete(Counter &priv, Counter n);
-
-  // check priv and decide if we have to wait or can free the other warps
-  INLINE void Release(Counter priv, Counter current_event_value);
-  INLINE void WaitOrRelease(Counter priv, Counter current_event_value);
-
-private:
-  Counter v_event; // counter of events (atomic)
-
-  // volatile is needed to force loads to read from global
-  // memory or L2 cache and see the write by the last master
-  volatile Counter v_start; // signal when events registered are finished
-
-  Counter v_init; // used to initialize local thread variables
-};
-
-#endif /* SRC_COUNTER_GROUP_H_ */

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

-//===----- counter_groupi.h - NVPTX OpenMP loop scheduling ------- CUDA -*-===//
-//
-//                     The LLVM Compiler Infrastructure
-//
-// This file is dual licensed under the MIT and the University of Illinois Open
-// Source Licenses. See LICENSE.txt for details.
-//
-//===----------------------------------------------------------------------===//
-//
-// Interface implementation for OpenMP loop scheduling
-//
-//===----------------------------------------------------------------------===//
-
-#include "option.h"
-
-INLINE void omptarget_nvptx_CounterGroup::Clear() {
-  PRINT0(LD_SYNCD, "clear counters\n")
-  v_event = 0;
-  v_start = 0;
-  // v_init does not need to be reset (its value is dead)
-}
-
-INLINE void omptarget_nvptx_CounterGroup::Reset() {
-  // done by master before entering parallel
-  ASSERT(LT_FUSSY, v_event == v_start,
-         "error, entry %lld !=start %lld at reset\n", P64(v_event),
-         P64(v_start));
-  v_init = v_start;
-}
-
-INLINE void omptarget_nvptx_CounterGroup::Init(Counter &priv) {
-  PRINT(LD_SYNCD, "init priv counter 0x%llx with val %lld\n", P64(&priv),
-        P64(v_start));
-  priv = v_start;
-}
-
-// just counts number of events
-INLINE Counter omptarget_nvptx_CounterGroup::Next() {
-  Counter oldVal = atomicAdd(&v_event, (Counter)1);
-  PRINT(LD_SYNCD, "next event counter 0x%llx with val %lld->%lld\n",
-        P64(&v_event), P64(oldVal), P64(oldVal + 1));
-
-  return oldVal;
-}
-
-// set priv to n, to be used in later waitOrRelease
-INLINE void omptarget_nvptx_CounterGroup::Complete(Counter &priv, Counter n) {
-  PRINT(LD_SYNCD, "complete priv counter 0x%llx with val %llu->%llu (+%llu)\n",
-        P64(&priv), P64(priv), P64(priv + n), n);
-  priv += n;
-}
-
-INLINE void omptarget_nvptx_CounterGroup::Release(Counter priv,
-                                                  Counter current_event_value) {
-  if (priv - 1 == current_event_value) {
-    PRINT(LD_SYNCD, "Release start counter 0x%llx with val %lld->%lld\n",
-          P64(&v_start), P64(v_start), P64(priv));
-    v_start = priv;
-  }
-}
-
-// check priv and decide if we have to wait or can free the other warps
-INLINE void
-omptarget_nvptx_CounterGroup::WaitOrRelease(Counter priv,
-                                            Counter current_event_value) {
-  if (priv - 1 == current_event_value) {
-    PRINT(LD_SYNCD, "Release start counter 0x%llx with val %lld->%lld\n",
-          P64(&v_start), P64(v_start), P64(priv));
-    v_start = priv;
-  } else {
-    PRINT(LD_SYNCD,
-          "Start waiting while start counter 0x%llx with val %lld < %lld\n",
-          P64(&v_start), P64(v_start), P64(priv));
-    while (priv > v_start) {
-      // IDLE LOOP
-      // start is volatile: it will be re-loaded at each while loop
-    }
-    PRINT(LD_SYNCD,
-          "Done waiting as start counter 0x%llx with val %lld >= %lld\n",
-          P64(&v_start), P64(v_start), P64(priv));
-  }
-}

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

 //  (mostly debug, temporary?) The third section includes the compiler
 //  specific interfaces.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _INTERFACES_H_
 #define _INTERFACES_H_
 
+#include "option.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // OpenMP interface
 ////////////////////////////////////////////////////////////////////////////////
 
 typedef uint32_t omp_lock_t;      /* arbitrary type of the right length */
 typedef uint64_t omp_nest_lock_t; /* arbitrary type of the right length */
 
 typedef enum omp_sched_t {

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 
 #include <inttypes.h>
 
 // cuda includes
 #include <cuda.h>
 #include <math.h>
 
 // local includes
-#include "counter_group.h"
 #include "debug.h"     // debug
 #include "interface.h" // interfaces with omp, compiler, and user
 #include "option.h"    // choices we have
 #include "state-queue.h"
 #include "support.h"
 
 #define OMPTARGET_NVPTX_VERSION 1.1
 

 
 ////////////////////////////////////////////////////////////////////////////////
 // Descriptor of a parallel region (worksharing in general)
 
 class omptarget_nvptx_WorkDescr {
 
 public:
   // access to data
-  INLINE omptarget_nvptx_CounterGroup &CounterGroup() { return cg; }
   INLINE omptarget_nvptx_TaskDescr *WorkTaskDescr() { return &masterTaskICV; }
-  // init
-  INLINE void InitWorkDescr();
 
 private:
-  omptarget_nvptx_CounterGroup cg; // for barrier (no other needed)
   omptarget_nvptx_TaskDescr masterTaskICV;
-  bool hasCancel;
 };
 
 ////////////////////////////////////////////////////////////////////////////////
 
 class omptarget_nvptx_TeamDescr {
 public:
   // access to data
   INLINE omptarget_nvptx_TaskDescr *LevelZeroTaskDescr() {
public:
   // parallel
   INLINE uint16_t &NumThreadsForNextParallel(int tid) {
     return nextRegion.tnum[tid];
   }
   // simd
   INLINE uint16_t &SimdLimitForNextSimd(int tid) {
     return nextRegion.slim[tid];
   }
-  // sync
-  INLINE Counter &Priv(int tid) { return priv[tid]; }
-  INLINE void IncrementPriv(int tid, Counter val) { priv[tid] += val; }
   // schedule (for dispatch)
   INLINE kmp_sched_t &ScheduleType(int tid) { return schedule[tid]; }
   INLINE int64_t &Chunk(int tid) { return chunk[tid]; }
   INLINE int64_t &LoopUpperBound(int tid) { return loopUpperBound[tid]; }
   INLINE int64_t &NextLowerBound(int tid) { return nextLowerBound[tid]; }
   INLINE int64_t &Stride(int tid) { return stride[tid]; }
 
   INLINE omptarget_nvptx_TeamDescr &TeamContext() { return teamContext; }
private:
   omptarget_nvptx_TaskDescr *topTaskDescr[MAX_THREADS_PER_TEAM];
   union {
     // Only one of the two is live at the same time.
     // parallel
     uint16_t tnum[MAX_THREADS_PER_TEAM];
     // simd limit
     uint16_t slim[MAX_THREADS_PER_TEAM];
   } nextRegion;
-  // sync
-  Counter priv[MAX_THREADS_PER_TEAM];
   // schedule (for dispatch)
   kmp_sched_t schedule[MAX_THREADS_PER_TEAM]; // remember schedule type for #for
   int64_t chunk[MAX_THREADS_PER_TEAM];
   int64_t loopUpperBound[MAX_THREADS_PER_TEAM];
   // state for dispatch with dyn/guided OR static (never use both at a time)
   int64_t nextLowerBound[MAX_THREADS_PER_TEAM];
   int64_t stride[MAX_THREADS_PER_TEAM];
   // Queue to which this object must be returned.

 INLINE omptarget_nvptx_WorkDescr &getMyWorkDescriptor();
 INLINE omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor();
 INLINE omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(int globalThreadId);
 
 ////////////////////////////////////////////////////////////////////////////////
 // inlined implementation
 ////////////////////////////////////////////////////////////////////////////////
 
-#include "counter_groupi.h"
 #include "omptarget-nvptxi.h"
 #include "supporti.h"
 
 #endif

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
 extern __device__
     omptarget_nvptx_Queue<omptarget_nvptx_ThreadPrivateContext, OMP_STATE_COUNT>
         omptarget_nvptx_device_State[MAX_SM];
 
-extern __device__ __shared__
-    omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
-
 extern __device__ omptarget_nvptx_Queue<
     omptarget_nvptx_SimpleThreadPrivateContext, OMP_STATE_COUNT>
     omptarget_nvptx_device_simpleState[MAX_SM];
 
-extern __device__ __shared__ omptarget_nvptx_SimpleThreadPrivateContext
-    *omptarget_nvptx_simpleThreadPrivateContext;
-
-//
-// The team master sets the outlined function and its arguments in these
-// variables to communicate with the workers.  Since they are in shared memory,
-// there is one copy of these variables for each kernel, instance, and team.
-//
-extern volatile __device__ __shared__ omptarget_nvptx_WorkFn
-    omptarget_nvptx_workFn;
-extern __device__ __shared__ uint32_t execution_param;
-
 ////////////////////////////////////////////////////////////////////////////////
 // init entry points
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE unsigned smid() {
   unsigned id;
   asm("mov.u32 %0, %%smid;" : "=r"(id));
   return id;
if (threadId == 0) {
     int slot = smid() % MAX_SM;
     omptarget_nvptx_threadPrivateContext =
         omptarget_nvptx_device_State[slot].Dequeue();
 
     omptarget_nvptx_TeamDescr &currTeamDescr = getMyTeamDescriptor();
     omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
     // init team context
     currTeamDescr.InitTeamDescr();
-    // init counters (copy start to init)
-    workDescr.CounterGroup().Reset();
   }
   __syncthreads();
 
   omptarget_nvptx_TeamDescr &currTeamDescr = getMyTeamDescriptor();
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
 
   //
   // Initialize task descr for each thread.
   //
   omptarget_nvptx_TaskDescr *newTaskDescr =
       omptarget_nvptx_threadPrivateContext->Level1TaskDescr(threadId);
   ASSERT0(LT_FUSSY, newTaskDescr, "expected a task descr");
   newTaskDescr->InitLevelOneTaskDescr(ThreadLimit,
                                       currTeamDescr.LevelZeroTaskDescr());
   newTaskDescr->ThreadLimit() = ThreadLimit;
   // install new top descriptor
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(threadId,
                                                              newTaskDescr);
 
   // init thread private from init value
-  workDescr.CounterGroup().Init(
-      omptarget_nvptx_threadPrivateContext->Priv(threadId));
   PRINT(LD_PAR,
         "thread will execute parallel region with id %d in a team of "
         "%d threads\n",
         newTaskDescr->ThreadId(), newTaskDescr->ThreadsInTeam());
 
   if (RequiresDataSharing && threadId % WARPSIZE == 0) {
     // Warp master innitializes data sharing environment.
     unsigned WID = threadId / WARPSIZE;

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptxi.h

 
 INLINE void
 omptarget_nvptx_ThreadPrivateContext::InitThreadPrivateContext(int tid) {
   // levelOneTaskDescr is init when starting the parallel region
   // top task descr is NULL (team master version will be fixed separately)
   topTaskDescr[tid] = NULL;
   // no num threads value has been pushed
   nextRegion.tnum[tid] = 0;
-  // priv counter init to zero
-  priv[tid] = 0;
   // the following don't need to be init here; they are init when using dyn
   // sched
   // current_Event, events_Number, chunk, num_Iterations, schedule
 }
 
 ////////////////////////////////////////////////////////////////////////////////
-// Work Descriptor
-////////////////////////////////////////////////////////////////////////////////
-
-INLINE void omptarget_nvptx_WorkDescr::InitWorkDescr() {
-  cg.Clear(); // start and stop to zero too
-  // threadsInParallelTeam does not need to be init (done in start parallel)
-  hasCancel = FALSE;
-}
-
-////////////////////////////////////////////////////////////////////////////////
 // Team Descriptor
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE void omptarget_nvptx_TeamDescr::InitTeamDescr() {
   levelZeroTaskDescr.InitLevelZeroTaskDescr();
-  workDescrForActiveParallel.InitWorkDescr();
-  // omp_init_lock(criticalLock);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Get private data structure for thread
 ////////////////////////////////////////////////////////////////////////////////
 
 // Utility routines for CUDA threads
 INLINE omptarget_nvptx_TeamDescr &getMyTeamDescriptor() {

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

 #define MAX_SM 16
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // algo options
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
-// data options
-////////////////////////////////////////////////////////////////////////////////
-
-// decide if counters are 32 or 64 bit
-#define Counter unsigned long long
-
-////////////////////////////////////////////////////////////////////////////////
 // misc options (by def everythig here is device)
 ////////////////////////////////////////////////////////////////////////////////
 
 #define EXTERN extern "C" __device__
 #define INLINE __inline__ __device__
 #define NOINLINE __noinline__ __device__
 #ifndef TRUE
 #define TRUE 1
 #endif
 #ifndef FALSE
 #define FALSE 0
 #endif
 
 #endif

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

ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
           "only team master can create parallel");
 
   // set number of threads on work descriptor
   // this is different from the number of cuda threads required for the parallel
   // region
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
   workDescr.WorkTaskDescr()->CopyToWorkDescr(currTaskDescr,
                                              CudaThreadsForParallel / NumLanes);
-  // init counters (copy start to init)
-  workDescr.CounterGroup().Reset();
 }
 
 // All workers call this function.  Deactivate those not needed.
 // Fn - the outlined work function to execute.
 // returns True if this thread is active, else False.
 //
 // Only the worker threads call this routine.
 EXTERN bool __kmpc_kernel_parallel(void **WorkFn,
if (threadId < workDescr.WorkTaskDescr()->ThreadsInTeam()) {
     omptarget_nvptx_TaskDescr *newTaskDescr =
         omptarget_nvptx_threadPrivateContext->Level1TaskDescr(threadId);
     ASSERT0(LT_FUSSY, newTaskDescr, "expected a task descr");
     newTaskDescr->CopyFromWorkDescr(workDescr.WorkTaskDescr());
     // install new top descriptor
     omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(threadId,
                                                                newTaskDescr);
     // init private from int value
-    workDescr.CounterGroup().Init(
-        omptarget_nvptx_threadPrivateContext->Priv(threadId));
     PRINT(LD_PAR,
           "thread will execute parallel region with id %d in a team of "
           "%d threads\n",
           newTaskDescr->ThreadId(), newTaskDescr->NThreads());
 
     isActive = true;
   }