projects / GPU.git / blobdiff
? search:
summary | shortlog | log | commit | commitdiff | tree
raw | inline | side by side
Ajout de l'ensemble du workspace.
[GPU.git] / WCudaMSE / BilatTools_Cuda / src / core / cudatools / header / device / reduction / reductionMinMaxLock.h
diff --git a/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/reduction/reductionMinMaxLock.h b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/reduction/reductionMinMaxLock.h
new file mode 100755 (executable)
index 0000000..cf9b713
--- /dev/null
+++ b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/reduction/reductionMinMaxLock.h
@@ -0,0 +1,304 @@
+#ifndef REDUCTION_MIN_MAX_LOCKH_\r
+#define REDUCTION_MIN_MAX_LOCKH_\r
+\r
+#include "Lock.h"\r
+\r
+#ifndef MIN\r
+#define MIN(X,Y) ((X)<(Y)?(X):(Y))\r
+#endif\r
+\r
+#ifndef MAX\r
+#define MAX(X,Y) ((X)>(Y)?(X):(Y))\r
+#endif\r
+\r
+//TODO : mieux : use foncteur\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    Declaration                                     *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+/*--------------------------------------*\\r
+ |*            Imported                *|\r
+ \*-------------------------------------*/\r
+\r
+#include "Indice1D.h"\r
+\r
+/*--------------------------------------*\\r
+ |*            Public                  *|\r
+ \*-------------------------------------*/\r
+\r
+/**\r
+ * Hyp\r
+ *     (H1) atomicMin n'existe pas pour T , sinon utiliser la version in reductionMinMax.h\r
+ *     (H2) tabBlock size is a power of 2\r
+ *     (H3) tabBlock is already fill with data\r
+ *\r
+ *  Usage LockGPU\r
+ *\r
+ *     // variable global .cu\r
+ *     __device__ int mutex=0; // Attention à l'initialisation\r
+ *\r
+ *     // Variable local inside kernel (same .cu as variable mutex)\r
+ *     LockGPU lock=LockGPU(&mutex);\r
+ *\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);\r
+\r
+/**\r
+ * Hyp\r
+ *     (H1) atomicMax n'existe pas pour T , sinon utiliser la version in reductionMinMax.h\r
+ *     (H2) tabBlock size is a power of 2\r
+ *     (H3) tabBlock is already fill with data\r
+ *\r
+ *  Usage LockGPU\r
+ *\r
+ *     // variable global .cu\r
+ *     __device__ int mutex=0; // Attention à l'initialisation\r
+ *\r
+ *     // Variable local inside kernel\r
+ *     LockGPU lock=LockGPU(&mutex);\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);\r
+\r
+/*--------------------------------------*\\r
+ |*            Private                 *|\r
+ \*-------------------------------------*/\r
+\r
+/*-----------------*\\r
+ |*    Intra      *|\r
+ \*-----------------*/\r
+\r
+template <typename T>\r
+static __device__ void reduction_IntraBlock_Min(T* tabBlock);\r
+\r
+template <typename T>\r
+static __device__ void reduction_IntraBlock_Max(T* tabBlock);\r
+\r
+/*-----------------*\\r
+ |*    Inter      *|\r
+ \*-----------------*/\r
+\r
+template <typename T>\r
+static __device__ void reduction_Interblock_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);\r
+\r
+template <typename T>\r
+static __device__ void reduction_Interblock_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);\r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    Implementation                                  *|\r
+ \*---------------------------------------------------------------------*/\r
+\r
+\r
+\r
+/*--------------------------------------*\\r
+ |*            Public                  *|\r
+ \*-------------------------------------*/\r
+\r
+template <typename T>\r
+__device__ void reduction_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)\r
+    {\r
+    reduction_IntraBlock_Min(tabBlock);\r
+\r
+    __syncthreads();\r
+\r
+    reduction_Interblock_Min(tabBlock, ptrDevResultat, ptrlock);\r
+    }\r
+\r
+template <typename T>\r
+__device__ void reduction_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)\r
+    {\r
+    reduction_IntraBlock_Max(tabBlock);\r
+\r
+    __syncthreads();\r
+\r
+    reduction_Interblock_Max(tabBlock, ptrDevResultat,ptrlock);\r
+    }\r
+\r
+/*--------------------------------------*\\r
+ |*            Private                 *|\r
+ \*-------------------------------------*/\r
+\r
+\r
+\r
+/*-----------------*\\r
+ |*    Intra      *|\r
+ \*-----------------*/\r
+\r
+\r
+/**\r
+ * Hyp :\r
+ *     (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2\r
+ *     (H2)    gride et Block de dim 1\r
+ *     (H3)    T est un type simple suppoter par atomicMin\r
+ */\r
+template <typename T>\r
+__device__ void reduction_IntraBlock_Min(T* tabBlock)\r
+    {\r
+//    int midle = blockDim.x / 2;\r
+//    int tidLocal = threadIdx.x;\r
+//\r
+//    // int tidLocal =   Indice1D::tidLocal()\r
+//\r
+//    while (midle >= 1)\r
+//     {\r
+//\r
+//     if (tidLocal < midle)\r
+//         {\r
+//         tabBlock[tidLocal] =MIN( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);\r
+//         }\r
+//\r
+//     __syncthreads();\r
+//\r
+//     //midle /= 2;\r
+//             midle>>=1;\r
+//     }\r
+\r
+    //v2\r
+        int midle = blockDim.x / 2;\r
+        int tidLocal = threadIdx.x;\r
+\r
+        //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases\r
+        while (midle >= 64)\r
+       {\r
+\r
+       if (tidLocal < midle)\r
+           {\r
+           tabBlock[tidLocal] =MIN( tabBlock[tidLocal], tabBlock[tidLocal + midle]);\r
+}\r
+\r
+       __syncthreads();\r
+\r
+       //midle /= 2;\r
+       midle>>=1;\r
+       }\r
+\r
+    // Utilisation des 32 thread d'un warp pour finir la reduction\r
+        if(tidLocal<32)\r
+       {\r
+       // no __syncthreads() necessary after exah of the following lines as long as  we acces the data via a pointzer decalred as volatile\r
+       // because teh 32 therad in each warp execute in a locked-step with each other\r
+       volatile T* ptrData=tabBlock;\r
+\r
+       ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 32]);\r
+ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 16]);\r
+ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 8]);\r
+ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 4]);\r
+ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 2]);\r
+ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 1]);\r
+}\r
+    }\r
+\r
+/**\r
+ * Hyp :\r
+ *     (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2\r
+ *     (H2)    gride et Block de dim 1\r
+ *     (H3)    T est un type simple suppoter par atomicMax\r
+ */\r
+template <typename T>\r
+__device__ void reduction_IntraBlock_Max(T* tabBlock)\r
+    {\r
+    // v1\r
+//    int midle = blockDim.x / 2;\r
+//    int tidLocal = threadIdx.x;\r
+//\r
+//    // int tidLocal =   Indice1D::tidLocal()\r
+//\r
+//    while (midle >= 1)\r
+//     {\r
+//\r
+//     if (tidLocal < midle)\r
+//         {\r
+//         tabBlock[tidLocal] =MAX( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);\r
+//         }\r
+//\r
+//     __syncthreads();\r
+//\r
+//     midle /= 2;\r
+//     }\r
+\r
+\r
+    //v2\r
+              int midle = blockDim.x / 2;\r
+              int tidLocal = threadIdx.x;\r
+\r
+              //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases\r
+              while (midle >= 64)\r
+               {\r
+\r
+               if (tidLocal < midle)\r
+                   {\r
+                   tabBlock[tidLocal] =MAX( tabBlock[tidLocal], tabBlock[tidLocal + midle]);\r
+   }\r
+\r
+               __syncthreads();\r
+\r
+               //midle /= 2;\r
+               midle>>=1;\r
+               }\r
+\r
+          // Utilisation des 32 thread d'un warp pour finir la reduction\r
+              if(tidLocal<32)\r
+               {\r
+               // no __syncthreads() necessary after exah of the following lines as long as  we acces the data via a pointzer decalred as volatile\r
+               // because teh 32 therad in each warp execute in a locked-step with each other\r
+               volatile float* ptrData=tabBlock;\r
+\r
+               ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 32]);\r
+   ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 16]);\r
+   ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 8]);\r
+   ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 4]);\r
+   ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 2]);\r
+   ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 1]);\r
+   }\r
+    }\r
+\r
+/*-----------------*\\r
+ |*    Inter      *|\r
+ \*-----------------*/\r
+\r
+/**\r
+ * Hyp :\r
+ *     (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2\r
+ *     (H2)    gride et Block de dim 1\r
+ *     (H3)    ptrDevResultat a été initialisé avec une valeur intelligente\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Interblock_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)\r
+    {\r
+    if (threadIdx.x == 0)\r
+       {\r
+       ptrlock->lock();\r
+\r
+       *ptrDevResultat=MIN(*ptrDevResultat, tabBlock[0]); // autant d'acces que de block\r
+\r
+       ptrlock->unlock();\r
+       }\r
+    }\r
+\r
+/**\r
+ * Hyp :\r
+ *     (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2\r
+ *     (H2)    gride et Block de dim 1\r
+ *     (H3)    ptrDevResultat a été initialisé avec une valeur intelligente\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Interblock_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)\r
+    {\r
+    if (threadIdx.x == 0)\r
+       {\r
+       ptrlock->lock();\r
+\r
+       *ptrDevResultat=MAX(*ptrDevResultat, tabBlock[0]); // autant d'acces que de block\r
+\r
+       ptrlock->unlock();\r
+       }\r
+    }\r
+\r
+\r
+#endif \r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    End                                             *|\r
+ \*---------------------------------------------------------------------*/\r
GPGPU with CUDA. Course from HES-SO given by C.Bilat.