[OPENMP][NVPTX]Improved lock/critical constructs, NFC.

[ABataev]

Summary

Improved support for critical constructs + omp_..._lock... constructs.

[gtbercea]

LG

This is actually a bug fix.

libomptarget/deviceRTLs/nvptx/src/critical.cu

 // This file contains the implementation of critical with KMPC interface
 //
 //===----------------------------------------------------------------------===//
 
 #include <stdio.h>
 
 #include "omptarget-nvptx.h"
 
-EXTERN
-void __kmpc_critical(kmp_Indent *loc, int32_t global_tid,
+EXTERN void __kmpc_critical(kmp_Indent *loc, int32_t global_tid,
                             kmp_CriticalName *lck) {
   PRINT0(LD_IO, "call to kmpc_critical()\n");
-  omptarget_nvptx_TeamDescr &teamDescr = getMyTeamDescriptor();
-  omp_set_lock(teamDescr.CriticalLock());
+  omp_set_lock((omp_lock_t *)lck);
 }
 
-EXTERN
-void __kmpc_end_critical(kmp_Indent *loc, int32_t global_tid,
+EXTERN void __kmpc_end_critical(kmp_Indent *loc, int32_t global_tid,
                                 kmp_CriticalName *lck) {
   PRINT0(LD_IO, "call to kmpc_end_critical()\n");
-  omptarget_nvptx_TeamDescr &teamDescr = getMyTeamDescriptor();
-  omp_unset_lock(teamDescr.CriticalLock());
+  omp_unset_lock((omp_lock_t *)lck);
 }

libomptarget/deviceRTLs/nvptx/src/libcall.cu

 // locks
 ////////////////////////////////////////////////////////////////////////////////
 
 #define __OMP_SPIN 1000
 #define UNSET 0
 #define SET 1
 
 EXTERN void omp_init_lock(omp_lock_t *lock) {
-  *lock = UNSET;
+  omp_unset_lock(lock);
   PRINT0(LD_IO, "call omp_init_lock()\n");
 }
 
 EXTERN void omp_destroy_lock(omp_lock_t *lock) {
+  omp_unset_lock(lock);
   PRINT0(LD_IO, "call omp_destroy_lock()\n");
 }
 
 EXTERN void omp_set_lock(omp_lock_t *lock) {
   // int atomicCAS(int* address, int compare, int val);
   // (old == compare ? val : old)
-  int compare = UNSET;
-  int val = SET;
 
   // TODO: not sure spinning is a good idea here..
-  while (atomicCAS(lock, compare, val) != UNSET) {
-
+  while (atomicCAS(lock, UNSET, SET) != UNSET) {
     clock_t start = clock();
     clock_t now;
     for (;;) {
       now = clock();
       clock_t cycles = now > start ? now - start : now + (0xffffffff - start);
       if (cycles >= __OMP_SPIN * blockIdx.x) {
         break;
       }
     }
   } // wait for 0 to be the read value
 
   PRINT0(LD_IO, "call omp_set_lock()\n");
 }
 
 EXTERN void omp_unset_lock(omp_lock_t *lock) {
-  int compare = SET;
-  int val = UNSET;
-  int old = atomicCAS(lock, compare, val);
+  (void)atomicExch(lock, UNSET);
 
   PRINT0(LD_IO, "call omp_unset_lock()\n");
 }
 
 EXTERN int omp_test_lock(omp_lock_t *lock) {
   // int atomicCAS(int* address, int compare, int val);
   // (old == compare ? val : old)
-  int compare = UNSET;
-  int val = SET;
-
-  int ret = atomicCAS(lock, compare, val);
+  int ret = atomicAdd(lock, 0);
 
   PRINT(LD_IO, "call omp_test_lock() return %d\n", ret);
 
   return ret;
 }
 
 // for xlf Fotran
 // Fotran, the return is LOGICAL type

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

 public:
   // access to data
   INLINE omptarget_nvptx_TaskDescr *LevelZeroTaskDescr() {
     return &levelZeroTaskDescr;
   }
   INLINE omptarget_nvptx_WorkDescr &WorkDescr() {
     return workDescrForActiveParallel;
   }
-  INLINE omp_lock_t *CriticalLock() { return &criticalLock; }
   INLINE uint64_t *getLastprivateIterBuffer() { return &lastprivateIterBuffer; }
 
   // init
   INLINE void InitTeamDescr();
 
   INLINE __kmpc_data_sharing_slot *RootS(int wid, bool IsMasterThread) {
     // If this is invoked by the master thread of the master warp then intialize
     // it with a smaller slot.
INLINE __kmpc_data_sharing_slot *GetPreallocatedSlotAddr(int wid) {
     return (__kmpc_data_sharing_slot *)&worker_rootS[wid];
   }
 
 private:
   omptarget_nvptx_TaskDescr
       levelZeroTaskDescr; // icv for team master initial thread
   omptarget_nvptx_WorkDescr
       workDescrForActiveParallel; // one, ONLY for the active par
-  omp_lock_t criticalLock;
   uint64_t lastprivateIterBuffer;
 
   __align__(16)
       __kmpc_data_sharing_worker_slot_static worker_rootS[WARPSIZE];
   __align__(16) __kmpc_data_sharing_master_slot_static master_rootS[1];
 };
 
 ////////////////////////////////////////////////////////////////////////////////