9.2
general documentation
cs_hip_reduce.h
Go to the documentation of this file.
1#ifndef CS_HIP_REDUCE_H
2#define CS_HIP_REDUCE_H
3
4/*============================================================================
5 * BLAS (Basic Linear Algebra Subroutine) functions
6 *============================================================================*/
7
8/*
9 This file is part of code_saturne, a general-purpose CFD tool.
10
11 Copyright (C) 1998-2026 EDF S.A.
12
13 This program is free software; you can redistribute it and/or modify it under
14 the terms of the GNU General Public License as published by the Free Software
15 Foundation; either version 2 of the License, or (at your option) any later
16 version.
17
18 This program is distributed in the hope that it will be useful, but WITHOUT
19 ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
20 FOR A PARTICULAR PURPOSE. See the GNU General Public License for more
21 details.
22
23 You should have received a copy of the GNU General Public License along with
24 this program; if not, write to the Free Software Foundation, Inc., 51 Franklin
25 Street, Fifth Floor, Boston, MA 02110-1301, USA.
26*/
27
28/*----------------------------------------------------------------------------*/
29
30#include "base/cs_defs.h"
31
32/*----------------------------------------------------------------------------
33 * External library headers
34 *----------------------------------------------------------------------------*/
35
36/*----------------------------------------------------------------------------
37 * Local headers
38 *----------------------------------------------------------------------------*/
39
40#include "base/cs_base.h"
41#include "base/cs_base_hip.h"
42
43/*============================================================================
44 * Macro definitions
45 *============================================================================*/
46
47/*============================================================================
48 * Type definitions
49 *============================================================================*/
50
51/*============================================================================
52 * Templated function definitions
53 *============================================================================*/
54
55#if defined(__HIPCC__)
56
57/*----------------------------------------------------------------------------*/
68/*----------------------------------------------------------------------------*/
69
70template <size_t blockSize, size_t stride, typename T>
71__device__ static void __forceinline__
72cs_hip_reduce_warp_reduce_sum(volatile T *stmp,
73 size_t tid)
74{
75 if (stride == 1) {
76
77#if CS_HIP_WARP_SIZE == 64
78 if (blockSize >= 128) stmp[tid] += stmp[tid + 64];
79#endif
80 if (blockSize >= 64) stmp[tid] += stmp[tid + 32];
81 if (blockSize >= 32) stmp[tid] += stmp[tid + 16];
82 if (blockSize >= 16) stmp[tid] += stmp[tid + 8];
83 if (blockSize >= 8) stmp[tid] += stmp[tid + 4];
84 if (blockSize >= 4) stmp[tid] += stmp[tid + 2];
85 if (blockSize >= 2) stmp[tid] += stmp[tid + 1];
86
87 }
88 else {
89
90#if CS_HIP_WARP_SIZE == 64
91 if (blockSize >= 128) {
92 #pragma unroll
93 for (size_t i = 0; i < stride; i++)
94 stmp[tid*stride + i] += stmp[(tid + 64)*stride + i];
95 }
96#endif
97 if (blockSize >= 64) {
98 #pragma unroll
99 for (size_t i = 0; i < stride; i++)
100 stmp[tid*stride + i] += stmp[(tid + 32)*stride + i];
101 }
102 if (blockSize >= 32) {
103 #pragma unroll
104 for (size_t i = 0; i < stride; i++)
105 stmp[tid*stride + i] += stmp[(tid + 16)*stride + i];
106 }
107 if (blockSize >= 16) {
108 #pragma unroll
109 for (size_t i = 0; i < stride; i++)
110 stmp[tid*stride + i] += stmp[(tid + 8)*stride + i];
111 }
112 if (blockSize >= 8) {
113 #pragma unroll
114 for (size_t i = 0; i < stride; i++)
115 stmp[tid*stride + i] += stmp[(tid + 4)*stride + i];
116 }
117 if (blockSize >= 4) {
118 #pragma unroll
119 for (size_t i = 0; i < stride; i++)
120 stmp[tid*stride + i] += stmp[(tid + 2)*stride + i];
121 }
122 if (blockSize >= 2) {
123 #pragma unroll
124 for (size_t i = 0; i < stride; i++)
125 stmp[tid*stride + i] += stmp[(tid + 1)*stride + i];
126 }
127 }
128}
129
130/*----------------------------------------------------------------------------*/
148/*----------------------------------------------------------------------------*/
149
150template <size_t blockSize, size_t stride, typename T>
151__device__ static void __forceinline__
152cs_hip_reduce_block_reduce_sum(T *stmp,
153 size_t tid,
154 T *sum_block)
155{
156 __syncthreads();
157
158 /* Loop explicitely unrolled */
159
160 if (stride == 1) { /* Scalar data */
161
162 if (blockSize >= 1024) {
163 if (tid < 512) {
164 stmp[tid] += stmp[tid + 512];
165 }
166 __syncthreads();
167 }
168 if (blockSize >= 512) {
169 if (tid < 256) {
170 stmp[tid] += stmp[tid + 256];
171 }
172 __syncthreads();
173 }
174 if (blockSize >= 256) {
175 if (tid < 128) {
176 stmp[tid] += stmp[tid + 128];
177 } __syncthreads();
178 }
179#if CS_HIP_WARP_SIZE == 32
180 if (blockSize >= 128) {
181 if (tid < 64) {
182 stmp[tid] += stmp[tid + 64];
183 } __syncthreads();
184 }
185
186 if (tid < 32) {
187 cs_hip_reduce_warp_reduce_sum<blockSize, stride>(stmp, tid);
188 }
189#else
190 if (tid < 64) {
191 cs_hip_reduce_warp_reduce_sum<blockSize, stride>(stmp, tid);
192 }
193#endif
194
195 // Output: b_res for this block
196
197 if (tid == 0) sum_block[blockIdx.x] = stmp[0];
198
199 }
200
201 else { /* Vector data */
202
203 if (blockSize >= 1024) {
204 if (tid < 512) {
205 #pragma unroll
206 for (size_t i = 0; i < stride; i++)
207 stmp[tid*stride + i] += stmp[(tid + 512)*stride + i];
208 }
209 __syncthreads();
210 }
211 if (blockSize >= 512) {
212 if (tid < 256) {
213 #pragma unroll
214 for (size_t i = 0; i < stride; i++)
215 stmp[tid*stride + i] += stmp[(tid + 256)*stride + i];
216 }
217 __syncthreads();
218 }
219 if (blockSize >= 256) {
220 if (tid < 128) {
221 #pragma unroll
222 for (size_t i = 0; i < stride; i++)
223 stmp[tid*stride + i] += stmp[(tid + 128)*stride + i];
224 } __syncthreads();
225 }
226#if CS_HIP_WARP_SIZE == 32
227 if (blockSize >= 128) {
228 if (tid < 64) {
229 #pragma unroll
230 for (size_t i = 0; i < stride; i++)
231 stmp[tid*stride + i] += stmp[(tid + 64)*stride + i];
232 } __syncthreads();
233 }
234
235 if (tid < 32)
236 cs_hip_reduce_warp_reduce_sum<blockSize, stride>(stmp, tid);
237#else
238 if (tid < 64)
239 cs_hip_reduce_warp_reduce_sum<blockSize, stride>(stmp, tid);
240#endif
241
242 // Output: b_res for this block
243
244 if (tid == 0) {
245 #pragma unroll
246 for (size_t i = 0; i < stride; i++)
247 sum_block[(blockIdx.x)*stride + i] = stmp[i];
248 }
249
250 }
251}
252
253/*----------------------------------------------------------------------------*/
263/*----------------------------------------------------------------------------*/
264
265template <size_t blockSize, size_t stride, typename T>
266__global__ static void
267cs_hip_reduce_sum_single_block(size_t n,
268 T *g_idata,
269 T *g_odata)
270{
271 __shared__ T sdata[blockSize * stride];
272
273 size_t tid = threadIdx.x;
274 T r_s[stride];
275
276 if (stride == 1) {
277
278 r_s[0] = 0;
279 sdata[tid] = 0.;
280
281 for (size_t i = threadIdx.x; i < n; i+= blockSize)
282 r_s[0] += g_idata[i];
283
284 sdata[tid] = r_s[0];
285 __syncthreads();
286
287 for (size_t j = blockSize/2; j > CS_HIP_WARP_SIZE; j /= 2) {
288 if (tid < j) {
289 sdata[tid] += sdata[tid + j];
290 }
291 __syncthreads();
292 }
293
294 if (tid < CS_HIP_WARP_SIZE) cs_hip_reduce_warp_reduce_sum<blockSize, stride>(sdata, tid);
295 if (tid == 0) *g_odata = sdata[0];
296
297 }
298 else {
299
300 #pragma unroll
301 for (size_t k = 0; k < stride; k++) {
302 r_s[k] = 0;
303 sdata[tid*stride + k] = 0.;
304 }
305
306 for (size_t i = threadIdx.x; i < n; i+= blockSize) {
307 #pragma unroll
308 for (size_t k = 0; k < stride; k++) {
309 r_s[k] += g_idata[i*stride + k];
310 }
311 }
312
313 #pragma unroll
314 for (size_t k = 0; k < stride; k++)
315 sdata[tid*stride + k] = r_s[k];
316 __syncthreads();
317
318 for (size_t j = blockSize/2; j > CS_HIP_WARP_SIZE; j /= 2) {
319 if (tid < j) {
320 #pragma unroll
321 for (size_t k = 0; k < stride; k++)
322 sdata[tid*stride + k] += sdata[(tid + j)*stride + k];
323 }
324 __syncthreads();
325 }
326
327 if (tid < CS_HIP_WARP_SIZE) cs_hip_reduce_warp_reduce_sum<blockSize, stride>(sdata, tid);
328 if (tid == 0) {
329 #pragma unroll
330 for (size_t k = 0; k < stride; k++)
331 g_odata[k] = sdata[k];
332 }
333
334 }
335}
336
337/*----------------------------------------------------------------------------*/
348/*----------------------------------------------------------------------------*/
349
350template <size_t blockSize, typename R, typename T>
351__device__ static void __forceinline__
352cs_hip_reduce_warp_reduce(volatile T *stmp,
353 size_t tid)
354{
355 R reducer;
356
357#if CS_HIP_WARP_SIZE == 64
358 if (blockSize >= 128) reducer.combine(stmp[tid], stmp[tid + 64]);
359#endif
360 if (blockSize >= 64) reducer.combine(stmp[tid], stmp[tid + 32]);
361 if (blockSize >= 32) reducer.combine(stmp[tid], stmp[tid + 16]);
362 if (blockSize >= 16) reducer.combine(stmp[tid], stmp[tid + 8]);
363 if (blockSize >= 8) reducer.combine(stmp[tid], stmp[tid + 4]);
364 if (blockSize >= 4) reducer.combine(stmp[tid], stmp[tid + 2]);
365 if (blockSize >= 2) reducer.combine(stmp[tid], stmp[tid + 1]);
366}
367
368/*----------------------------------------------------------------------------*/
383/*----------------------------------------------------------------------------*/
384
385template <size_t blockSize, typename R, typename T>
386__device__ static void __forceinline__
387cs_hip_reduce_block_reduce(T *stmp,
388 size_t tid,
389 T *rd_block)
390{
391 R reducer;
392 __syncthreads();
393
394 /* Loop explicitely unrolled */
395
396 if (blockSize >= 1024) {
397 if (tid < 512) {
398 reducer.combine(stmp[tid], stmp[tid + 512]);
399 }
400 __syncthreads();
401 }
402 if (blockSize >= 512) {
403 if (tid < 256) {
404 reducer.combine(stmp[tid], stmp[tid + 256]);
405 }
406 __syncthreads();
407 }
408 if (blockSize >= 256) {
409 if (tid < 128) {
410 reducer.combine(stmp[tid], stmp[tid + 128]);
411 } __syncthreads();
412 }
413#if CS_HIP_WARP_SIZE == 32
414 if (blockSize >= 128) {
415 if (tid < 64) {
416 reducer.combine(stmp[tid], stmp[tid + 64]);
417 } __syncthreads();
418 }
419
420 if (tid < 32) {
421 cs_hip_reduce_warp_reduce<blockSize, R>(stmp, tid);
422 }
423#else
424 if (tid < 64) {
425 cs_hip_reduce_warp_reduce<blockSize, R>(stmp, tid);
426 }
427#endif
428
429 // Output: rd_block for this block
430
431 if (tid == 0) rd_block[blockIdx.x] = stmp[0];
432}
433
434/*----------------------------------------------------------------------------*/
446/*----------------------------------------------------------------------------*/
447
448template <size_t blockSize, typename R, typename T>
449__global__ static void
450cs_hip_reduce_single_block(size_t n,
451 T *g_idata,
452 T *g_odata)
453{
454 extern __shared__ int p_stmp[];
455 T *sdata = reinterpret_cast<T *>(p_stmp);
456 R reducer;
457
458 size_t tid = threadIdx.x;
459 T r_s[1];
460
461 reducer.identity(r_s[0]);
462 reducer.identity(sdata[tid]);
463
464 for (size_t i = threadIdx.x; i < n; i+= blockSize)
465 reducer.combine(r_s[0], g_idata[i]);
466
467 sdata[tid] = r_s[0];
468 __syncthreads();
469
470 for (size_t j = blockSize/2; j > CS_HIP_WARP_SIZE; j /= 2) {
471 if (tid < j) {
472 reducer.combine(sdata[tid], sdata[tid + j]);
473 }
474 __syncthreads();
475 }
476
477 if (tid < CS_HIP_WARP_SIZE) cs_hip_reduce_warp_reduce<blockSize, R>(sdata, tid);
478 if (tid == 0) *g_odata = sdata[0];
479}
480
481#endif /* defined(__HIPCC__) */
482
483/*============================================================================
484 * Public function prototypes
485 *============================================================================*/
486
487/*----------------------------------------------------------------------------*/
488
489#endif /* CS_HIP_REDUCE_H */
cs_base.h
cs_base_hip.h
cs_defs.h
k
@ k
Definition: cs_field_pointer.h:68