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