X-Git-Url: http://git.euphorik.ch/?p=GPU.git;a=blobdiff_plain;f=WCudaMSE%2FBilatTools_Cuda%2Fsrc%2Fcore%2Fcudatools%2Fheader%2Fdevice%2Fsynchronisation%2FLock.h;fp=WCudaMSE%2FBilatTools_Cuda%2Fsrc%2Fcore%2Fcudatools%2Fheader%2Fdevice%2Fsynchronisation%2FLock.h;h=97534a0ab1248db0c06e63d9819aac655675b5dc;hp=4c6787d4f00079cfcf4a6463269f7c52943cf5e7;hb=19015d26dfb874d075516772ef531ee5e42fa213;hpb=8cb42977ea2d8904795381957474b65f15f46526 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 index 4c6787d..97534a0 100755 --- a/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/synchronisation/Lock.h +++ b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/synchronisation/Lock.h @@ -20,19 +20,19 @@ * 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. * * Interet - * On utilise un lock lorsque l'opération qui doit être synchroniser ne possède pas d'opérateur atomic (comme atomicADD, ...), - * ou lorsqui'il s'agit de plusieurs opérations à synchroniser (joue alors le role de section critique) + * On utilise un lock lorsque l'op�ration qui doit �tre synchroniser ne poss�de pas d'op�rateur atomic (comme atomicADD, ...), + * ou lorsqui'il s'agit de plusieurs op�rations � synchroniser (joue alors le role de section critique) * * Note : * - * Lock ne laisse aucune trace coté host, il s'instancie only coté device: Code moins invasif - * LockMixte laisse une trace coté host. Code plus invasif + * Lock ne laisse aucune trace cot� host, il s'instancie only cot� device: Code moins invasif + * LockMixte laisse une trace cot� host. Code plus invasif * * Use (Device side only) * * * // Global variable of .cu - * __device__ int mutex=0; // Attention à l'initialisation + * __device__ int mutex=0; // Attention � l'initialisation * * // variable local inside a kernel (same .cu as variable mutex) * Lock lock=Lock(&mutex); @@ -76,10 +76,10 @@ class Lock // Solution: // atomicCAS = atomic Compare And Swap // Prototype : c atomicCAS(ptr,a,b) - // Action : compare ptr avec a, si egale affecte b à ptr, renvoie ptr + // Action : compare ptr avec a, si egale affecte b � ptr, renvoie ptr // Tant que ptrDev_mutex!=0 le thread cuda boucle sur while - // Des qu'il vaut 0, il met le mutex à 1 et lock se termine + // Des qu'il vaut 0, il met le mutex � 1 et lock se termine while (atomicCAS(ptrDevMutexGM, 0, 1) != 0); } @@ -92,7 +92,7 @@ class Lock // Solution 1: // // *ptrDev_mutex=0; - // Aucun thread en competition ici. L'affectation n'a pas besoin d'être atomique. + // Aucun thread en competition ici. L'affectation n'a pas besoin d'�tre atomique. // Solution satisfaisante. // // Solution 2 (prefered for symetric approach) @@ -112,7 +112,7 @@ class Lock int* ptrDevMutexGM; }; -#endif +#endif /*----------------------------------------------------------------------*\ |* End *|