1 #ifndef REDUCTION_MIN_MAX_H_
2 #define REDUCTION_MIN_MAX_H_
5 #define MIN(X,Y) ((X)<(Y)?(X):(Y))
9 #define MAX(X,Y) ((X)>(Y)?(X):(Y))
13 /*----------------------------------------------------------------------*\
15 \*---------------------------------------------------------------------*/
17 /*--------------------------------------*\
19 \*-------------------------------------*/
23 /*--------------------------------------*\
25 \*-------------------------------------*/
29 * (H1) atomicMin existe pour T , sinon utiliser la version in reductionMinMaxLock.h
30 * (H2) tabBlock size is a power of 2
31 * (H3) tabBlock is already fill with data
34 __device__
void reduction_Min(T
* tabBlock
, T
* ptrDevResultat
);
38 * (H1) atomicMax existe pour T , sinon utiliser la version in reductionMinMaxLock.h
39 * (H2) tabBlock size is a power of 2
40 * (H3) tabBlock is already fill with data
43 __device__
void reduction_Max(T
* tabBlock
, T
* ptrDevResultat
);
45 /*--------------------------------------*\
47 \*-------------------------------------*/
54 static __device__
void reduction_IntraBlock_Min(T
* tabBlock
);
57 static __device__
void reduction_IntraBlock_Max(T
* tabBlock
);
64 static __device__
void reduction_Interblock_Min(T
* tabBlock
, T
* ptrDevResultat
);
67 static __device__
void reduction_Interblock_Max(T
* tabBlock
, T
* ptrDevResultat
);
69 /*----------------------------------------------------------------------*\
71 \*---------------------------------------------------------------------*/
75 /*--------------------------------------*\
77 \*-------------------------------------*/
80 __device__
void reduction_Min(T
* tabBlock
, T
* ptrDevResultat
)
82 reduction_IntraBlock_Min(tabBlock
);
86 reduction_Interblock_Min(tabBlock
, ptrDevResultat
);
90 __device__
void reduction_Max(T
* tabBlock
, T
* ptrDevResultat
)
92 reduction_IntraBlock_Max(tabBlock
);
96 reduction_Interblock_Max(tabBlock
, ptrDevResultat
);
99 /*--------------------------------------*\
101 \*-------------------------------------*/
105 /*-----------------*\
107 \*-----------------*/
112 * (H1) length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
113 * (H2) gride et Block de dim 1
114 * (H3) T est un type simple suppoter par atomicMin (sinon use reductionMinMaxLock.h)
116 template <typename T
>
117 __device__
void reduction_IntraBlock_Min(T
* tabBlock
)
120 // int midle = blockDim.x / 2;
121 // int tidLocal = threadIdx.x;
123 // // int tidLocal = Indice1D::tidLocal()
125 // while (midle >= 1)
128 // if (tidLocal < midle)
130 // tabBlock[tidLocal] =MIN( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);
133 // __syncthreads(); // threads d'un meme block
140 int midle
= blockDim
.x
/ 2;
141 int tidLocal
= threadIdx
.x
;
143 //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases
147 if (tidLocal
< midle
)
149 tabBlock
[tidLocal
] =MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ midle
]);
158 // Utilisation des 32 thread d'un warp pour finir la reduction
161 // no __syncthreads() necessary after exah of the following lines as long as we acces the data via a pointzer decalred as volatile
162 // because teh 32 therad in each warp execute in a locked-step with each other
163 volatile T
* ptrData
=tabBlock
;
165 ptrData
[tidLocal
]=MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 32]);
166 ptrData
[tidLocal
]=MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 16]);
167 ptrData
[tidLocal
]=MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 8]);
168 ptrData
[tidLocal
]=MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 4]);
169 ptrData
[tidLocal
]=MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 2]);
170 ptrData
[tidLocal
]=MIN( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 1]);
176 * (H1) length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
177 * (H2) gride et Block de dim 1
178 * (H3) T est un type simple suppoter par atomicMax (sinon use reductionMinMaxLock.h)
180 template <typename T
>
181 __device__
void reduction_IntraBlock_Max(T
* tabBlock
)
184 // int midle = blockDim.x / 2;
185 // int tidLocal = threadIdx.x;
187 // // int tidLocal = Indice1D::tidLocal()
189 // while (midle >= 1)
192 // if (tidLocal < midle)
194 // tabBlock[tidLocal] =MAX( tabBlock[tidLocal] ,tabBlock[tidLocal + midle]);
197 // __syncthreads(); // threads d'un meme block
203 int midle
= blockDim
.x
/ 2;
204 int tidLocal
= threadIdx
.x
;
206 //a 64 on ne divise plus et on a besoin de 32 thread pour finir de reduire le 64 premières cases
210 if (tidLocal
< midle
)
212 tabBlock
[tidLocal
] =MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ midle
]);
221 // Utilisation des 32 thread d'un warp pour finir la reduction
224 // no __syncthreads() necessary after exah of the following lines as long as we acces the data via a pointzer decalred as volatile
225 // because teh 32 therad in each warp execute in a locked-step with each other
226 volatile float* ptrData
=tabBlock
;
228 ptrData
[tidLocal
]=MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 32]);
229 ptrData
[tidLocal
]=MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 16]);
230 ptrData
[tidLocal
]=MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 8]);
231 ptrData
[tidLocal
]=MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 4]);
232 ptrData
[tidLocal
]=MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 2]);
233 ptrData
[tidLocal
]=MAX( tabBlock
[tidLocal
], tabBlock
[tidLocal
+ 1]);
237 /*-----------------*\
239 \*-----------------*/
243 * (H1) length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
244 * (H2) gride et Block de dim 1
245 * (H3) T est un type simple suppoter par atomicMin
246 * (H4) ptrDevResultat a été initialisé avec une valeur intelligente
248 template <typename T
>
249 __device__
void reduction_Interblock_Min(T
* tabBlock
, T
* ptrDevResultat
)
251 if (threadIdx
.x
== 0)
253 atomicMin(ptrDevResultat
, tabBlock
[0]); // autant d'acces que de block
259 * (H1) length(tabBlock) est une puissance de 2 ! ie blockDim.x est une puissance de 2
260 * (H2) gride et Block de dim 1
261 * (H3) T est un type simple suppoter par atomicMax
262 * (H4) ptrDevResultat a été initialisé avec une valeur intelligente
264 template <typename T
>
265 __device__
void reduction_Interblock_Max(T
* tabBlock
, T
* ptrDevResultat
)
267 if (threadIdx
.x
== 0)
269 atomicMax(ptrDevResultat
, tabBlock
[0]); // autant d'acces que de block
275 /*----------------------------------------------------------------------*\
277 \*---------------------------------------------------------------------*/