4 /*----------------------------------------------------------------------*\
6 \*---------------------------------------------------------------------*/
8 /*--------------------------------------*\
10 \*-------------------------------------*/
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.
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)
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
31 * Use (Device side only)
34 * // Global variable of .cu
35 * __device__ int mutex=0; // Attention � l'initialisation
37 * // variable local inside a kernel (same .cu as variable mutex)
38 * Lock lock=Lock(&mutex);
48 /*--------------------------------------*\
50 \*-------------------------------------*/
55 * Toutes les instance se partagent la meme adresse du mutex ptrDevMutexGM
58 Lock(int* ptrDevMutexGM
)
60 this->ptrDevMutexGM
= ptrDevMutexGM
;
63 /*--------------------------------------*\
65 \*-------------------------------------*/
73 // if (*ptrDev_mutex==0) {*ptrDev_mutex==1}
74 // But must be thread safe!
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
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);
95 // Aucun thread en competition ici. L'affectation n'a pas besoin d'�tre atomique.
96 // Solution satisfaisante.
98 // Solution 2 (prefered for symetric approach)
100 // Une solution atomique
102 // Echange et renvoie la valeur originale
103 atomicExch(ptrDevMutexGM
, 0);
106 /*--------------------------------------*\
108 \*-------------------------------------*/
117 /*----------------------------------------------------------------------*\
119 \*---------------------------------------------------------------------*/