9.2
general documentation
cs_blas_hip.h
Go to the documentation of this file.
1#ifndef CS_BLAS_HIP_H
2#define CS_BLAS_HIP_H
3
4/*============================================================================
5 * BLAS (Basic Linear Algebra Subroutine) functions using HIP.
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_blas_hip_warp_reduce_sum(volatile T *stmp,
73 size_t tid)
74{
75 if (stride == 1) {
76#if CS_HIP_WARP_SIZE == 64
77 if (blockSize >= 128) stmp[tid] += stmp[tid + 64];
78#endif
79 if (blockSize >= 64) stmp[tid] += stmp[tid + 32];
80 if (blockSize >= 32) stmp[tid] += stmp[tid + 16];
81 if (blockSize >= 16) stmp[tid] += stmp[tid + 8];
82 if (blockSize >= 8) stmp[tid] += stmp[tid + 4];
83 if (blockSize >= 4) stmp[tid] += stmp[tid + 2];
84 if (blockSize >= 2) stmp[tid] += stmp[tid + 1];
85
86 }
87 else {
88
89#if CS_HIP_WARP_SIZE == 64
90 if (blockSize >= 128) {
91 #pragma unroll
92 for (size_t i = 0; i < stride; i++)
93 stmp[tid*stride + i] += stmp[(tid + 64)*stride + i];
94 }
95#endif
96 if (blockSize >= 64) {
97 #pragma unroll
98 for (size_t i = 0; i < stride; i++)
99 stmp[tid*stride + i] += stmp[(tid + 32)*stride + i];
100 }
101 if (blockSize >= 32) {
102 #pragma unroll
103 for (size_t i = 0; i < stride; i++)
104 stmp[tid*stride + i] += stmp[(tid + 16)*stride + i];
105 }
106 if (blockSize >= 16) {
107 #pragma unroll
108 for (size_t i = 0; i < stride; i++)
109 stmp[tid*stride + i] += stmp[(tid + 8)*stride + i];
110 }
111 if (blockSize >= 8) {
112 #pragma unroll
113 for (size_t i = 0; i < stride; i++)
114 stmp[tid*stride + i] += stmp[(tid + 4)*stride + i];
115 }
116 if (blockSize >= 4) {
117 #pragma unroll
118 for (size_t i = 0; i < stride; i++)
119 stmp[tid*stride + i] += stmp[(tid + 2)*stride + i];
120 }
121 if (blockSize >= 2) {
122 #pragma unroll
123 for (size_t i = 0; i < stride; i++)
124 stmp[tid*stride + i] += stmp[(tid + 1)*stride + i];
125 }
126 }
127}
128
129/*----------------------------------------------------------------------------*/
143/*----------------------------------------------------------------------------*/
144
145template <size_t blockSize, size_t stride, typename T>
146__device__ static void __forceinline__
147cs_blas_hip_block_reduce_sum(T *stmp,
148 size_t tid,
149 T *sum_block)
150{
151 __syncthreads();
152
153 /* Loop explicitely unrolled */
154
155 if (stride == 1) { /* Scalar data */
156
157 if (blockSize >= 1024) {
158 if (tid < 512) {
159 stmp[tid] += stmp[tid + 512];
160 }
161 __syncthreads();
162 }
163 if (blockSize >= 512) {
164 if (tid < 256) {
165 stmp[tid] += stmp[tid + 256];
166 }
167 __syncthreads();
168 }
169 if (blockSize >= 256) {
170 if (tid < 128) {
171 stmp[tid] += stmp[tid + 128];
172 } __syncthreads();
173 }
174#if CS_HIP_WARP_SIZE == 32
175 if (blockSize >= 128) {
176 if (tid < 64) {
177 stmp[tid] += stmp[tid + 64];
178 } __syncthreads();
179 }
180
181 if (tid < 32) {
182 cs_blas_hip_warp_reduce_sum<blockSize, stride>(stmp, tid);
183 }
184#else
185 if (tid < 64) {
186 cs_blas_hip_warp_reduce_sum<blockSize, stride>(stmp, tid);
187 }
188#endif
189
190 // Output: b_res for this block
191
192 if (tid == 0) sum_block[blockIdx.x] = stmp[0];
193
194 }
195
196 else { /* Vector data */
197
198 if (blockSize >= 1024) {
199 if (tid < 512) {
200 #pragma unroll
201 for (size_t i = 0; i < stride; i++)
202 stmp[tid*stride + i] += stmp[(tid + 512)*stride + i];
203 }
204 __syncthreads();
205 }
206 if (blockSize >= 512) {
207 if (tid < 256) {
208 #pragma unroll
209 for (size_t i = 0; i < stride; i++)
210 stmp[tid*stride + i] += stmp[(tid + 256)*stride + i];
211 }
212 __syncthreads();
213 }
214 if (blockSize >= 256) {
215 if (tid < 128) {
216 #pragma unroll
217 for (size_t i = 0; i < stride; i++)
218 stmp[tid*stride + i] += stmp[(tid + 128)*stride + i];
219 } __syncthreads();
220 }
221#if CS_HIP_WARP_SIZE == 32
222 if (blockSize >= 128) {
223 if (tid < 64) {
224 #pragma unroll
225 for (size_t i = 0; i < stride; i++)
226 stmp[tid*stride + i] += stmp[(tid + 64)*stride + i];
227 } __syncthreads();
228 }
229
230 if (tid < 32)
231 cs_blas_hip_warp_reduce_sum<blockSize, stride>(stmp, tid);
232#else
233 if (tid < 64)
234 cs_blas_hip_warp_reduce_sum<blockSize, stride>(stmp, tid);
235#endif
236
237 // Output: b_res for this block
238
239 if (tid == 0) {
240 #pragma unroll
241 for (size_t i = 0; i < stride; i++)
242 sum_block[(blockIdx.x)*stride + i] = stmp[i];
243 }
244
245 }
246}
247
248/*----------------------------------------------------------------------------*/
258/*----------------------------------------------------------------------------*/
259
260template <size_t blockSize, size_t stride, typename T>
261__global__ static void
262cs_blas_hip_reduce_single_block(size_t n,
263 T *g_idata,
264 T *g_odata)
265{
266 __shared__ T sdata[blockSize * stride];
267
268 size_t tid = threadIdx.x;
269 T r_s[stride];
270
271 if (stride == 1) {
272
273 r_s[0] = 0;
274 sdata[tid] = 0.;
275
276 for (size_t i = threadIdx.x; i < n; i+= blockSize)
277 r_s[0] += g_idata[i];
278
279 sdata[tid] = r_s[0];
280 __syncthreads();
281
282 for (size_t j = blockSize/2; j > CS_HIP_WARP_SIZE; j /= 2) {
283 if (tid < j) {
284 sdata[tid] += sdata[tid + j];
285 }
286 __syncthreads();
287 }
288
289 if (tid < CS_HIP_WARP_SIZE) cs_blas_hip_warp_reduce_sum<blockSize, stride>(sdata, tid);
290 if (tid == 0) *g_odata = sdata[0];
291
292 }
293 else {
294
295 #pragma unroll
296 for (size_t k = 0; k < stride; k++) {
297 r_s[k] = 0;
298 sdata[tid*stride + k] = 0.;
299 }
300
301 for (size_t i = threadIdx.x; i < n; i+= blockSize) {
302 #pragma unroll
303 for (size_t k = 0; k < stride; k++) {
304 r_s[k] += g_idata[i*stride + k];
305 }
306 }
307
308 #pragma unroll
309 for (size_t k = 0; k < stride; k++)
310 sdata[tid*stride + k] = r_s[k];
311 __syncthreads();
312
313 for (size_t j = blockSize/2; j > CS_HIP_WARP_SIZE; j /= 2) {
314 if (tid < j) {
315 #pragma unroll
316 for (size_t k = 0; k < stride; k++)
317 sdata[tid*stride + k] += sdata[(tid + j)*stride + k];
318 }
319 __syncthreads();
320 }
321
322 if (tid < CS_HIP_WARP_SIZE) cs_blas_hip_warp_reduce_sum<blockSize, stride>(sdata, tid);
323 if (tid == 0) {
324 #pragma unroll
325 for (size_t k = 0; k < stride; k++)
326 g_odata[k] = sdata[k];
327 }
328
329 }
330}
331
332#endif /* defined(__HIPCC__) */
333
334/*============================================================================
335 * Public function prototypes
336 *============================================================================*/
337
338/*----------------------------------------------------------------------------*/
339/*
340 * \brief Finalize HIP BLAS API.
341 *
342 * This frees resources such as the hipBLAS handle, if used.
343 */
344/*----------------------------------------------------------------------------*/
345
346extern "C" void
347cs_blas_hip_finalize(void);
348
349#if defined(__HIPCC__)
350
351/*----------------------------------------------------------------------------*/
352/*
353 * \brief Return the absolute sum of vector values using HIP.
354 *
355 * \param[in] stream associated HIP stream
356 * \param[in] n size of array x
357 * \param[in] x array of floating-point values (on device)
358 *
359 * \return sum of absolute array values
360 */
361/*----------------------------------------------------------------------------*/
362
363double
364cs_blas_hip_asum(hipStream_t stream,
365 cs_lnum_t n,
366 const cs_real_t x[]);
367
368/*----------------------------------------------------------------------------*/
369/*
370 * \brief Return the dot product of 2 vectors: x.y using HIP.
371 *
372 * \param[in] stream associated HIP stream
373 * \param[in] n size of arrays x and y
374 * \param[in] x array of floating-point values (on device)
375 * \param[in] y array of floating-point values (on device)
376 *
377 * \return dot product
378 */
379/*----------------------------------------------------------------------------*/
380
381double
382cs_blas_hip_dot(hipStream_t stream,
383 cs_lnum_t n,
384 const cs_real_t x[],
385 const cs_real_t y[]);
386
387#if defined(HAVE_HIPBLAS)
388
389/*----------------------------------------------------------------------------*/
390/*
391 * \brief Return the absolute sum of vector values using hipBLAS.
392 *
393 * \param[in] stream associated HIP stream
394 * \param[in] n size of arrays x and y
395 * \param[in] x array of floating-point values (on device)
396 *
397 * \return sum of absolute array values
398 */
399/*----------------------------------------------------------------------------*/
400
401double
402cs_blas_hipblas_asum(hipStream_t stream,
403 cs_lnum_t n,
404 const cs_real_t x[]);
405
406/*----------------------------------------------------------------------------*/
407/*
408 * \brief Return the dot product of 2 vectors: x.y using hipBLAS.
409 *
410 * \param[in] stream associated HIP stream
411 * \param[in] n size of arrays x and y
412 * \param[in] x array of floating-point values (on device)
413 * \param[in] y array of floating-point values (on device)
414 *
415 * \return dot product
416 */
417/*----------------------------------------------------------------------------*/
418
419double
420cs_blas_hipblas_dot(hipStream_t stream,
421 cs_lnum_t n,
422 const cs_real_t x[],
423 const cs_real_t y[]);
424
425#endif /* defined(HAVE_HIPBLAS) */
426
427/*----------------------------------------------------------------------------
428 * Compute x <- alpha.x
429 *
430 * This function may be set to use either hipBLAS or a local kernel.
431 *
432 * \param[in] stream associated HIP stream
433 * \param[in] n number of elements
434 * \param[in] alpha constant value (on device)
435 * \param[in] x vector of elements (on device)
436 *----------------------------------------------------------------------------*/
437
438void
439cs_blas_hip_scal(hipStream_t stream,
440 cs_lnum_t n,
441 const cs_real_t *alpha,
442 cs_real_t *x);
443
444#endif /* defined(__HIPCC__)
445
446/*----------------------------------------------------------------------------*/
447
448#endif /* CS_BLAS_HIP_H */
cs_base.h
cs_base_hip.h
cs_blas_hip_finalize
void cs_blas_hip_finalize(void)
Finalize HIP BLAS API.
Definition: cs_blas_hip.cpp:343
cs_defs.h
cs_real_t
double cs_real_t
Floating-point value.
Definition: cs_defs.h:332
cs_lnum_t
int cs_lnum_t
local mesh entity id
Definition: cs_defs.h:325
k
@ k
Definition: cs_field_pointer.h:68