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 / reductionMinMax.h
diff --git a/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/reduction/reductionMinMax.h b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/reduction/reductionMinMax.h
new file mode 100755 (executable)
index 0000000..b12ca30
--- /dev/null
+++ b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/reduction/reductionMinMax.h
@@ -0,0 +1,277 @@
+#ifndef REDUCTION_MIN_MAX_H_\r
+#define REDUCTION_MIN_MAX_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
+\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 existe pour T , sinon utiliser la version in reductionMinMaxLock.h\r
+ *     (H2) tabBlock size is a power of 2\r
+ *     (H3) tabBlock is already fill with data\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Min(T* tabBlock, T* ptrDevResultat);\r
+\r
+/**\r
+ * Hyp\r
+ *     (H1) atomicMax existe pour T , sinon utiliser la version in reductionMinMaxLock.h\r
+ *     (H2) tabBlock size is a power of 2\r
+ *     (H3) tabBlock is already fill with data\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Max(T* tabBlock, T* ptrDevResultat);\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);\r
+\r
+template <typename T>\r
+static __device__ void reduction_Interblock_Max(T* tabBlock, T* ptrDevResultat);\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)\r
+    {\r
+    reduction_IntraBlock_Min(tabBlock);\r
+\r
+    __syncthreads();\r
+\r
+    reduction_Interblock_Min(tabBlock, ptrDevResultat);\r
+    }\r
+\r
+template <typename T>\r
+__device__ void reduction_Max(T* tabBlock, T* ptrDevResultat)\r
+    {\r
+    reduction_IntraBlock_Max(tabBlock);\r
+\r
+    __syncthreads();\r
+\r
+    reduction_Interblock_Max(tabBlock, ptrDevResultat);\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 (sinon use reductionMinMaxLock.h)\r
+ */\r
+template <typename T>\r
+__device__ void reduction_IntraBlock_Min(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] =MIN( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);\r
+//         }\r
+//\r
+//     __syncthreads(); // threads d'un meme block\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  (sinon use reductionMinMaxLock.h)\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(); // threads d'un meme block\r
+//\r
+//     midle /= 2;\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)    T est un type simple suppoter par atomicMin\r
+ *     (H4)    ptrDevResultat a été initialisé avec une valeur intelligente\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Interblock_Min(T* tabBlock, T* ptrDevResultat)\r
+{\r
+if (threadIdx.x == 0)\r
+{\r
+atomicMin(ptrDevResultat, tabBlock[0]); // autant d'acces que de block\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
+ *     (H4)    ptrDevResultat a été initialisé avec une valeur intelligente\r
+ */\r
+template <typename T>\r
+__device__ void reduction_Interblock_Max(T* tabBlock, T* ptrDevResultat)\r
+{\r
+if (threadIdx.x == 0)\r
+{\r
+atomicMax(ptrDevResultat, tabBlock[0]); // autant d'acces que de block\r
+}\r
+}\r
+\r
+#endif \r
+\r
+/*----------------------------------------------------------------------*\\r
+ |*                    End                                             *|\r
+ \*---------------------------------------------------------------------*/\r
GPGPU with CUDA. Course from HES-SO given by C.Bilat.