WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/device/reduction/reductionMinMaxLock.h

   1 #ifndef REDUCTION_MIN_MAX_LOCKH_
   2 #define REDUCTION_MIN_MAX_LOCKH_
   3
   4 #include "Lock.h"
   5
   6 #ifndef MIN
   7 #define MIN(X,Y) ((X)<(Y)?(X):(Y))
   8 #endif
   9
  10 #ifndef MAX
  11 #define MAX(X,Y) ((X)>(Y)?(X):(Y))
  12 #endif
  13
  14 //TODO : mieux : use foncteur
  15
  16 /*----------------------------------------------------------------------*\
  17  |*                     Declaration                                     *|
  18  \*---------------------------------------------------------------------*/
  19
  20 /*--------------------------------------*\
  21  |*             Imported                *|
  22  \*-------------------------------------*/
  23
  24 #include "Indice1D.h"
  25
  26 /*--------------------------------------*\
  27  |*             Public                  *|
  28  \*-------------------------------------*/
  29
  30 /**
  31  * Hyp
  32  *      (H1) atomicMin n'existe pas pour T , sinon utiliser la version in reductionMinMax.h
  33  *      (H2) tabBlock size is a power of 2
  34  *      (H3) tabBlock is already fill with data
  35  *
  36  *  Usage LockGPU
  37  *
  38  *      // variable global .cu
  39  *      __device__ int mutex=0; // Attention à l'initialisation
  40  *
  41  *      // Variable local inside kernel (same .cu as variable mutex)
  42  *      LockGPU lock=LockGPU(&mutex);
  43  *
  44  */
  45 template <typename T>
  46 __device__ void reduction_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);
  47
  48 /**
  49  * Hyp
  50  *      (H1) atomicMax n'existe pas pour T , sinon utiliser la version in reductionMinMax.h
  51  *      (H2) tabBlock size is a power of 2
  52  *      (H3) tabBlock is already fill with data
  53  *
  54  *  Usage LockGPU
  55  *
  56  *      // variable global .cu
  57  *      __device__ int mutex=0; // Attention à l'initialisation
  58  *
  59  *      // Variable local inside kernel
  60  *      LockGPU lock=LockGPU(&mutex);
  61  */
  62 template <typename T>
  63 __device__ void reduction_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);
  64
  65 /*--------------------------------------*\
  66  |*             Private                 *|
  67  \*-------------------------------------*/
  68
  69 /*-----------------*\
  70  |*     Intra      *|
  71  \*-----------------*/
  72
  73 template <typename T>
  74 static __device__ void reduction_IntraBlock_Min(T* tabBlock);
  75
  76 template <typename T>
  77 static __device__ void reduction_IntraBlock_Max(T* tabBlock);
  78
  79 /*-----------------*\
  80  |*     Inter      *|
  81  \*-----------------*/
  82
  83 template <typename T>
  84 static __device__ void reduction_Interblock_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);
  85
  86 template <typename T>
  87 static __device__ void reduction_Interblock_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock);
  88
  89 /*----------------------------------------------------------------------*\
  90  |*                     Implementation                                  *|
  91  \*---------------------------------------------------------------------*/
  92
  93
  94
  95 /*--------------------------------------*\
  96  |*             Public                  *|
  97  \*-------------------------------------*/
  98
  99 template <typename T>
 100 __device__ void reduction_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)
 101     {
 102     reduction_IntraBlock_Min(tabBlock);
 103
 104     __syncthreads();
 105
 106     reduction_Interblock_Min(tabBlock, ptrDevResultat, ptrlock);
 107     }
 108
 109 template <typename T>
 110 __device__ void reduction_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)
 111     {
 112     reduction_IntraBlock_Max(tabBlock);
 113
 114     __syncthreads();
 115
 116     reduction_Interblock_Max(tabBlock, ptrDevResultat,ptrlock);
 117     }
 118
 119 /*--------------------------------------*\
 120  |*             Private                 *|
 121  \*-------------------------------------*/
 122
 123
 124
 125 /*-----------------*\
 126  |*     Intra      *|
 127  \*-----------------*/
 128
 129
 130 /**
 131  * Hyp :
 132  *      (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
 133  *      (H2)    gride et Block de dim 1
 134  *      (H3)    T est un type simple suppoter par atomicMin
 135  */
 136 template <typename T>
 137 __device__ void reduction_IntraBlock_Min(T* tabBlock)
 138     {
 139 //    int midle = blockDim.x / 2;
 140 //    int tidLocal = threadIdx.x;
 141 //
 142 //    // int tidLocal =   Indice1D::tidLocal()
 143 //
 144 //    while (midle >= 1)
 145 //      {
 146 //
 147 //      if (tidLocal < midle)
 148 //          {
 149 //          tabBlock[tidLocal] =MIN( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);
 150 //          }
 151 //
 152 //      __syncthreads();
 153 //
 154 //      //midle /= 2;
 155 //              midle>>=1;
 156 //      }
 157
 158     //v2
 159         int midle = blockDim.x / 2;
 160         int tidLocal = threadIdx.x;
 161
 162         //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases
 163         while (midle >= 64)
 164         {
 165
 166         if (tidLocal < midle)
 167             {
 168             tabBlock[tidLocal] =MIN( tabBlock[tidLocal], tabBlock[tidLocal + midle]);
 169 }
 170
 171         __syncthreads();
 172
 173         //midle /= 2;
 174         midle>>=1;
 175         }
 176
 177     // Utilisation des 32 thread d'un warp pour finir la reduction
 178         if(tidLocal<32)
 179         {
 180         // no __syncthreads() necessary after exah of the following lines as long as  we acces the data via a pointzer decalred as volatile
 181         // because teh 32 therad in each warp execute in a locked-step with each other
 182         volatile T* ptrData=tabBlock;
 183
 184         ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 32]);
 185 ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 16]);
 186 ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 8]);
 187 ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 4]);
 188 ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 2]);
 189 ptrData[tidLocal]=MIN( tabBlock[tidLocal], tabBlock[tidLocal + 1]);
 190 }
 191     }
 192
 193 /**
 194  * Hyp :
 195  *      (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
 196  *      (H2)    gride et Block de dim 1
 197  *      (H3)    T est un type simple suppoter par atomicMax
 198  */
 199 template <typename T>
 200 __device__ void reduction_IntraBlock_Max(T* tabBlock)
 201     {
 202     // v1
 203 //    int midle = blockDim.x / 2;
 204 //    int tidLocal = threadIdx.x;
 205 //
 206 //    // int tidLocal =   Indice1D::tidLocal()
 207 //
 208 //    while (midle >= 1)
 209 //      {
 210 //
 211 //      if (tidLocal < midle)
 212 //          {
 213 //          tabBlock[tidLocal] =MAX( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);
 214 //          }
 215 //
 216 //      __syncthreads();
 217 //
 218 //      midle /= 2;
 219 //      }
 220
 221
 222     //v2
 223               int midle = blockDim.x / 2;
 224               int tidLocal = threadIdx.x;
 225
 226               //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases
 227               while (midle >= 64)
 228                 {
 229
 230                 if (tidLocal < midle)
 231                     {
 232                     tabBlock[tidLocal] =MAX( tabBlock[tidLocal], tabBlock[tidLocal + midle]);
 233    }
 234
 235                 __syncthreads();
 236
 237                 //midle /= 2;
 238                 midle>>=1;
 239                 }
 240
 241           // Utilisation des 32 thread d'un warp pour finir la reduction
 242               if(tidLocal<32)
 243                 {
 244                 // no __syncthreads() necessary after exah of the following lines as long as  we acces the data via a pointzer decalred as volatile
 245                 // because teh 32 therad in each warp execute in a locked-step with each other
 246                 volatile float* ptrData=tabBlock;
 247
 248                 ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 32]);
 249    ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 16]);
 250    ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 8]);
 251    ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 4]);
 252    ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 2]);
 253    ptrData[tidLocal]=MAX( tabBlock[tidLocal], tabBlock[tidLocal + 1]);
 254    }
 255     }
 256
 257 /*-----------------*\
 258  |*     Inter      *|
 259  \*-----------------*/
 260
 261 /**
 262  * Hyp :
 263  *      (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
 264  *      (H2)    gride et Block de dim 1
 265  *      (H3)    ptrDevResultat a été initialisé avec une valeur intelligente
 266  */
 267 template <typename T>
 268 __device__ void reduction_Interblock_Min(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)
 269     {
 270     if (threadIdx.x == 0)
 271         {
 272         ptrlock->lock();
 273
 274         *ptrDevResultat=MIN(*ptrDevResultat, tabBlock[0]); // autant d'acces que de block
 275
 276         ptrlock->unlock();
 277         }
 278     }
 279
 280 /**
 281  * Hyp :
 282  *      (H1)    length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
 283  *      (H2)    gride et Block de dim 1
 284  *      (H3)    ptrDevResultat a été initialisé avec une valeur intelligente
 285  */
 286 template <typename T>
 287 __device__ void reduction_Interblock_Max(T* tabBlock, T* ptrDevResultat,Lock* ptrlock)
 288     {
 289     if (threadIdx.x == 0)
 290         {
 291         ptrlock->lock();
 292
 293         *ptrDevResultat=MAX(*ptrDevResultat, tabBlock[0]); // autant d'acces que de block
 294
 295         ptrlock->unlock();
 296         }
 297     }
 298
 299
 300 #endif
 301
 302 /*----------------------------------------------------------------------*\
 303  |*                     End                                             *|
 304  \*---------------------------------------------------------------------*/