WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/synchronisation/Lock.h

   1 #ifndef LOCK_CUDA_H_
   2 #define LOCK_CUDA_H_
   3
   4 /*----------------------------------------------------------------------*\
   5  |*                     Declaration                                     *|
   6  \*---------------------------------------------------------------------*/
   7
   8 /*--------------------------------------*\
   9  |*             Public                  *|
  10  \*-------------------------------------*/
  11
  12 /**
  13  * General :
  14  *
  15  *      A Mutex will act like something of a traffic signal that governs access to some resources.
  16  *      When a thread reads a 0 from the mutex,it interprets this valu as a "green lignt" indicating that no other thread is using the memory.
  17  *      Therefore, the thread is free to lock the memory and make whatever changes it desires,free of interference from other threads.
  18  *      To lock the memory location in nquestion, the thread writes a 1 to the mutex.
  19  *      This will act as a "red light" for potentially competing threads.
  20  *      The competing threads must then wait until the owner has written a 0 to the mutex beforfe they can attempt to modify the locked memory.
  21  *
  22  * Interet
  23  *      On utilise un lock lorsque l'op�ration qui doit �tre synchroniser ne poss�de pas d'op�rateur atomic (comme atomicADD, ...),
  24  *      ou lorsqui'il s'agit de plusieurs op�rations � synchroniser (joue alors le role de section critique)
  25  *
  26  * Note :
  27  *
  28  *      Lock ne laisse aucune trace cot� host, il s'instancie only cot� device: Code moins invasif
  29  *      LockMixte laisse une trace cot� host. Code plus invasif
  30  *
  31  * Use (Device side only)
  32  *
  33  *
  34  *      // Global variable of .cu
  35  *      __device__ int mutex=0; // Attention � l'initialisation
  36  *
  37  *      // variable local inside a kernel (same .cu as variable mutex)
  38  *      Lock lock=Lock(&mutex);
  39  *
  40  *      lock.lock();
  41  *      doSomething();
  42  *      lock.unlock();
  43  *
  44  */
  45 class Lock
  46     {
  47
  48         /*--------------------------------------*\
  49          |*             Constructor             *|
  50          \*-------------------------------------*/
  51
  52     public:
  53
  54         /**
  55          * Toutes les instance se partagent la meme adresse du mutex ptrDevMutexGM
  56          */
  57         __device__
  58         Lock(int* ptrDevMutexGM)
  59             {
  60             this->ptrDevMutexGM = ptrDevMutexGM;
  61             }
  62
  63         /*--------------------------------------*\
  64          |*             Methodes                *|
  65          \*-------------------------------------*/
  66
  67     public:
  68
  69         __device__
  70         void lock(void)
  71             {
  72             // Goal :
  73             //          if (*ptrDev_mutex==0) {*ptrDev_mutex==1}
  74             //          But must be thread safe!
  75             //
  76             // Solution:
  77             //          atomicCAS  = atomic Compare And Swap
  78             //          Prototype :     c atomicCAS(ptr,a,b)
  79             //          Action :        compare ptr avec a, si egale affecte b � ptr, renvoie ptr
  80
  81             // Tant que ptrDev_mutex!=0 le thread cuda boucle sur while
  82             // Des qu'il vaut 0, il met le mutex � 1 et lock se termine
  83             while (atomicCAS(ptrDevMutexGM, 0, 1) != 0);
  84             }
  85
  86         __device__
  87         void unlock(void)
  88             {
  89             // Goal :
  90             //          Put 1 in the mutex
  91             //
  92             // Solution 1:
  93             //
  94             //          *ptrDev_mutex=0;
  95             //          Aucun thread en competition ici. L'affectation n'a pas besoin d'�tre atomique.
  96             //          Solution satisfaisante.
  97             //
  98             // Solution 2 (prefered for symetric approach)
  99             //
 100             //          Une solution atomique
 101
 102             // Echange et renvoie la valeur originale
 103             atomicExch(ptrDevMutexGM, 0);
 104             }
 105
 106         /*--------------------------------------*\
 107          |*             Attributs               *|
 108          \*-------------------------------------*/
 109
 110     private:
 111
 112         int* ptrDevMutexGM;
 113     };
 114
 115 #endif
 116
 117 /*----------------------------------------------------------------------*\
 118  |*                     End                                             *|
 119  \*---------------------------------------------------------------------*/