1 #ifndef REDUCTION_MIN_MAX_LOCKH_
2 #define REDUCTION_MIN_MAX_LOCKH_
7 #define MIN(X,Y) ((X)<(Y)?(X):(Y))
11 #define MAX(X,Y) ((X)>(Y)?(X):(Y))
14 //TODO : mieux : use foncteur
16 /*----------------------------------------------------------------------*\
18 \*---------------------------------------------------------------------*/
20 /*--------------------------------------*\
22 \*-------------------------------------*/
26 /*--------------------------------------*\
28 \*-------------------------------------*/
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
38 * // variable global .cu
39 * __device__ int mutex=0; // Attention à l'initialisation
41 * // Variable local inside kernel (same .cu as variable mutex)
42 * LockGPU lock=LockGPU(&mutex);
46 __device__
void reduction_Min(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
);
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
56 * // variable global .cu
57 * __device__ int mutex=0; // Attention à l'initialisation
59 * // Variable local inside kernel
60 * LockGPU lock=LockGPU(&mutex);
63 __device__
void reduction_Max(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
);
65 /*--------------------------------------*\
67 \*-------------------------------------*/
74 static __device__
void reduction_IntraBlock_Min(T
* tabBlock
);
77 static __device__
void reduction_IntraBlock_Max(T
* tabBlock
);
84 static __device__
void reduction_Interblock_Min(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
);
87 static __device__
void reduction_Interblock_Max(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
);
89 /*----------------------------------------------------------------------*\
91 \*---------------------------------------------------------------------*/
95 /*--------------------------------------*\
97 \*-------------------------------------*/
100 __device__
void reduction_Min(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
)
102 reduction_IntraBlock_Min(tabBlock
);
106 reduction_Interblock_Min(tabBlock
, ptrDevResultat
, ptrlock
);
109 template <typename T
>
110 __device__
void reduction_Max(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
)
112 reduction_IntraBlock_Max(tabBlock
);
116 reduction_Interblock_Max(tabBlock
, ptrDevResultat
,ptrlock
);
119 /*--------------------------------------*\
121 \*-------------------------------------*/
125 /*-----------------*\
127 \*-----------------*/
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
136 template <typename T
>
137 __device__
void reduction_IntraBlock_Min(T
* tabBlock
)
139 // int midle = blockDim.x / 2;
140 // int tidLocal = threadIdx.x;
142 // // int tidLocal = Indice1D::tidLocal()
144 // while (midle >= 1)
147 // if (tidLocal < midle)
149 // tabBlock[tidLocal] =MIN( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);
159 int midle
= blockDim
.x
/ 2;
160 int tidLocal
= threadIdx
.x
;
162 //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases
166 if (tidLocal
< midle
)
168 tabBlock
[tidLocal
] =MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ midle
]);
177 // Utilisation des 32 thread d'un warp pour finir la reduction
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
;
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]);
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
199 template <typename T
>
200 __device__
void reduction_IntraBlock_Max(T
* tabBlock
)
203 // int midle = blockDim.x / 2;
204 // int tidLocal = threadIdx.x;
206 // // int tidLocal = Indice1D::tidLocal()
208 // while (midle >= 1)
211 // if (tidLocal < midle)
213 // tabBlock[tidLocal] =MAX( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);
223 int midle
= blockDim
.x
/ 2;
224 int tidLocal
= threadIdx
.x
;
226 //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases
230 if (tidLocal
< midle
)
232 tabBlock
[tidLocal
] =MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ midle
]);
241 // Utilisation des 32 thread d'un warp pour finir la reduction
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
;
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]);
257 /*-----------------*\
259 \*-----------------*/
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
267 template <typename T
>
268 __device__
void reduction_Interblock_Min(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
)
270 if (threadIdx
.x
== 0)
274 *ptrDevResultat
=MIN(*ptrDevResultat
, tabBlock
[0]); // autant d'acces que de block
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
286 template <typename T
>
287 __device__
void reduction_Interblock_Max(T
* tabBlock
, T
* ptrDevResultat
,Lock
* ptrlock
)
289 if (threadIdx
.x
== 0)
293 *ptrDevResultat
=MAX(*ptrDevResultat
, tabBlock
[0]); // autant d'acces que de block
302 /*----------------------------------------------------------------------*\
304 \*---------------------------------------------------------------------*/