6 /*----------------------------------------------------------------------*\
8 \*---------------------------------------------------------------------*/
10 /*--------------------------------------*\
12 \*-------------------------------------*/
19 * A Mutex will act like something of a traffic signal that governs access to some resources.
20 * 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.
21 * Therefore, the thread is free to lock the memory and make whatever changes it desires,free of interference from other threads.
22 * To lock the memory location in nquestion, the thread writes a 1 to the mutex.
23 * This will act as a "red light" for potentially competing threads.
24 * 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.
27 * On utilise un lock lorsque l'opération qui doit être synchroniser ne possède pas d'opérateur atomic (comme atomicADD, ...),
28 * ou lorsqui'il s'agit de plusieurs opérations à synchroniser (joue alors le role de section critique)
32 * Lock ne laisse aucune trace coté host, il s'instancie only coté device: Code moins invasif
33 * LockMixte laisse une trace coté host. Code plus invasif
38 * kernel<<<dg,db>>>(...,lock,...); // pas besoin de memory managment MM
50 /*--------------------------------------*\
52 \*-------------------------------------*/
59 HANDLE_ERROR(cudaMalloc((void**) &ptrDev_mutex
, sizeof(int)));
60 HANDLE_ERROR(cudaMemcpy(ptrDev_mutex
, &state
, sizeof(int), cudaMemcpyHostToDevice
));
66 * 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.
67 * Le destructeur sera donc appeler 2x,la premiere fois, lors de la fin du kernel sur le gpu, la seconde fois sur le cpu.
68 * 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
69 * (affichage d'un message d'error) sur le GPU et qui est donc enlever ici.
70 * Heureusement cudaFree peut être appeler plusieurs fois de suite sans probleme,mais un seul appel suffirait!
74 * Sur le GPU, il ne faut jamais passer Lock par valeur, car sinon, la premiere instance detruite, detruit ptrDev_mutex !!
75 * et le mutex devient alors inutilisable!
79 // HANDLE_ERROR(cudaFree(ptrDev_mutex)); // HANDLE_ERROR : pas possible car executer aussi sur le GPU
80 cudaFree(ptrDev_mutex
);
83 /*--------------------------------------*\
85 \*-------------------------------------*/
89 __device__
void lock(void)
92 // if (*ptrDev_mutex==0) {*ptrDev_mutex==1}
93 // But must be thread safe!
96 // atomicCAS = atomic Compare And Swap
97 // Prototype : c atomicCAS(ptr,a,b)
98 // Action : compare ptr avec a, si egale affecte b à ptr, renvoie ptr
100 // Tant que ptrDev_mutex!=0 le thread cuda boucle sur while
101 // Des qu'il vaut 0, il met le mutex à 1 et lock se termine
102 while (atomicCAS(ptrDev_mutex
, 0, 1) != 0);
105 __device__
void unlock(void)
108 // Put 1 in the mutex
113 // Aucun thread en competition ici. L'affectation n'a pas besoin d'être atomique.
114 // Solution satisfaisante.
116 // Solution 2 (prefered for symetric approach)
118 // Une solution atomique
120 // Echange et renvoie la valeur originale
121 atomicExch(ptrDev_mutex
, 0);
124 /*--------------------------------------*\
126 \*-------------------------------------*/
130 int* ptrDev_mutex
; // Espace adressage GPU, en global memory GM
135 /*----------------------------------------------------------------------*\
137 \*---------------------------------------------------------------------*/