--- /dev/null
+#ifndef LOCK_MIXTE_H_\r
+#define LOCK_MIXTE_H_\r
+\r
+#include "cudaTools.h"\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |* Declaration *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |* Public *|\r
+ \*-------------------------------------*/\r
+\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 (host) :\r
+ *\r
+ * LockMixte lock();\r
+ * kernel<<<dg,db>>>(...,lock,...); // pas besoin de memory managment MM\r
+ *\r
+ * Use (device):\r
+ *\r
+ * lock.lock();\r
+ * doSomething();\r
+ * lock.unlock();\r
+ *\r
+ */\r
+class LockMixte\r
+ {\r
+\r
+ /*--------------------------------------*\\r
+ |* Constructor *|\r
+ \*-------------------------------------*/\r
+\r
+ public :\r
+\r
+ LockMixte(void)\r
+ {\r
+ int state = 0;\r
+ HANDLE_ERROR(cudaMalloc((void**) &ptrDev_mutex, sizeof(int)));\r
+ HANDLE_ERROR(cudaMemcpy(ptrDev_mutex, &state, sizeof(int), cudaMemcpyHostToDevice));\r
+ }\r
+\r
+ /**\r
+ * Observation:\r
+ *\r
+ * Comme Lock est passer par valeur au kernel,ie par copie,deux objets au total sont vivants, un sur le cpu et sun sur le gpu.\r
+ * Le destructeur sera donc appeler 2x,la premiere fois, lors de la fin du kernel sur le gpu, la seconde fois sur le cpu.\r
+ * Comme cudaFree peut etre appeler sur le cpu ou gpu, le code ci-dessous est valide, mais sans le HANDLE_ERROR qui n'a aucun sens\r
+ * (affichage d'un message d'error) sur le GPU et qui est donc enlever ici.\r
+ * Heureusement cudaFree peut être appeler plusieurs fois de suite sans probleme,mais un seul appel suffirait!\r
+ *\r
+ * Attention:\r
+ *\r
+ * Sur le GPU, il ne faut jamais passer Lock par valeur, car sinon, la premiere instance detruite, detruit ptrDev_mutex !!\r
+ * et le mutex devient alors inutilisable!\r
+ */\r
+ ~LockMixte(void)\r
+ {\r
+ // HANDLE_ERROR(cudaFree(ptrDev_mutex)); // HANDLE_ERROR : pas possible car executer aussi sur le GPU\r
+ cudaFree(ptrDev_mutex);\r
+ }\r
+\r
+ /*--------------------------------------*\\r
+ |* Methodes *|\r
+ \*-------------------------------------*/\r
+\r
+ public :\r
+\r
+ __device__ 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(ptrDev_mutex, 0, 1) != 0);\r
+ }\r
+\r
+ __device__ 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(ptrDev_mutex, 0);\r
+ }\r
+\r
+ /*--------------------------------------*\\r
+ |* Attributs *|\r
+ \*-------------------------------------*/\r
+\r
+ private :\r
+\r
+ int* ptrDev_mutex; // Espace adressage GPU, en global memory GM\r
+ };\r
+\r
+#endif \r
+\r
+/*----------------------------------------------------------------------*\\r
+ |* End *|\r
+ \*---------------------------------------------------------------------*/\r
projects / GPU.git / blobdiff