Ajout de l'ensemble du workspace.

[GPU.git] / WCudaMSE / BilatTools_Cuda / src / core / cudatools / header / device / synchronisation / Lock.h

diff --git a/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/synchronisation/Lock.h b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/synchronisation/Lock.h
new file mode 100755 (executable)
index 0000000..4c6787d
--- /dev/null
+++ b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/synchronisation/Lock.h
@@ -0,0 +1,119 @@
+#ifndef LOCK_CUDA_H_\r
+#define LOCK_CUDA_H_\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    Declaration                                     *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |*            Public                  *|\r
+ \*-------------------------------------*/\r
+\r
+/**\r
+ * General :\r
+ *\r
+ *     A Mutex will act like something of a traffic signal that governs access to some resources.\r
+ *     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.\r
+ *     Therefore, the thread is free to lock the memory and make whatever changes it desires,free of interference from other threads.\r
+ *     To lock the memory location in nquestion, the thread writes a 1 to the mutex.\r
+ *     This will act as a "red light" for potentially competing threads.\r
+ *     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.\r
+ *\r
+ * Interet\r
+ *     On utilise un lock lorsque l'opération qui doit être synchroniser ne possède pas d'opérateur atomic (comme atomicADD, ...),\r
+ *     ou lorsqui'il s'agit de plusieurs opérations à synchroniser (joue alors le role de section critique)\r
+ *\r
+ * Note :\r
+ *\r
+ *     Lock ne laisse aucune trace coté host, il s'instancie only coté device: Code moins invasif\r
+ *     LockMixte laisse une trace coté host. Code plus invasif\r
+ *\r
+ * Use (Device side only)\r
+ *\r
+ *\r
+ *     // Global variable of .cu\r
+ *     __device__ int mutex=0; // Attention à l'initialisation\r
+ *\r
+ *     // variable local inside a kernel (same .cu as variable mutex)\r
+ *     Lock lock=Lock(&mutex);\r
+ *\r
+ *     lock.lock();\r
+ *     doSomething();\r
+ *     lock.unlock();\r
+ *\r
+ */\r
+class Lock\r
+    {\r
+\r
+       /*--------------------------------------*\\r
+        |*             Constructor             *|\r
+        \*-------------------------------------*/\r
+\r
+    public:\r
+\r
+       /**\r
+        * Toutes les instance se partagent la meme adresse du mutex ptrDevMutexGM\r
+        */\r
+       __device__\r
+       Lock(int* ptrDevMutexGM)\r
+           {\r
+           this->ptrDevMutexGM = ptrDevMutexGM;\r
+           }\r
+\r
+       /*--------------------------------------*\\r
+        |*             Methodes                *|\r
+        \*-------------------------------------*/\r
+\r
+    public:\r
+\r
+       __device__\r
+       void lock(void)\r
+           {\r
+           // Goal :\r
+           //          if (*ptrDev_mutex==0) {*ptrDev_mutex==1}\r
+           //          But must be thread safe!\r
+           //\r
+           // Solution:\r
+           //          atomicCAS  = atomic Compare And Swap\r
+           //          Prototype :     c atomicCAS(ptr,a,b)\r
+           //          Action :        compare ptr avec a, si egale affecte b à ptr, renvoie ptr\r
+\r
+           // Tant que ptrDev_mutex!=0 le thread cuda boucle sur while\r
+           // Des qu'il vaut 0, il met le mutex à 1 et lock se termine\r
+           while (atomicCAS(ptrDevMutexGM, 0, 1) != 0);\r
+           }\r
+\r
+       __device__\r
+       void unlock(void)\r
+           {\r
+           // Goal :\r
+           //          Put 1 in the mutex\r
+           //\r
+           // Solution 1:\r
+           //\r
+           //          *ptrDev_mutex=0;\r
+           //          Aucun thread en competition ici. L'affectation n'a pas besoin d'être atomique.\r
+           //          Solution satisfaisante.\r
+           //\r
+           // Solution 2 (prefered for symetric approach)\r
+           //\r
+           //          Une solution atomique\r
+\r
+           // Echange et renvoie la valeur originale\r
+           atomicExch(ptrDevMutexGM, 0);\r
+           }\r
+\r
+       /*--------------------------------------*\\r
+        |*             Attributs               *|\r
+        \*-------------------------------------*/\r
+\r
+    private:\r
+\r
+       int* ptrDevMutexGM;\r
+    };\r
+\r
+#endif \r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    End                                             *|\r
+ \*---------------------------------------------------------------------*/\r