+#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