Neko 1.99.1
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
ax_helm.hip
Go to the documentation of this file.
1/*
2 Copyright (c) 2021-2025, The Neko Authors
3 All rights reserved.
4
5 Redistribution and use in source and binary forms, with or without
6 modification, are permitted provided that the following conditions
7 are met:
8
9 * Redistributions of source code must retain the above copyright
10 notice, this list of conditions and the following disclaimer.
11
12 * Redistributions in binary form must reproduce the above
13 copyright notice, this list of conditions and the following
14 disclaimer in the documentation and/or other materials provided
15 with the distribution.
16
17 * Neither the name of the authors nor the names of its
18 contributors may be used to endorse or promote products derived
19 from this software without specific prior written permission.
20
21 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32 POSSIBILITY OF SUCH DAMAGE.
33*/
34
35#include <string.h>
36#include <stdlib.h>
37#include <stdio.h>
38#include <hip/hip_runtime.h>
39#include <device/device_config.h>
40#include <device/hip/check.h>
41#include "ax_helm_kernel.h"
42
43extern "C" {
44 #include <common/neko_log.h>
45}
46
47template < const int>
48int tune(void *w, void *u, void *dx, void *dy, void *dz,
49 void *dxt, void *dyt, void *dzt, void *h1,
50 void *g11, void *g22, void *g33, void *g12,
51 void *g13, void *g23, int *nelv, int *lx);
52
53template < const int>
54int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
55 void *dxt, void *dyt, void *dzt, void *h1,
56 void *g11, void *g22, void *g33, void *g12,
57 void *g13, void *g23, int *nelv, int *lx);
58
59extern "C" {
60
64 void hip_ax_helm(void *w, void *u, void *dx, void *dy, void *dz,
65 void *dxt, void *dyt, void *dzt, void *h1,
66 void *g11, void *g22, void *g33, void *g12,
67 void *g13, void *g23, int *nelv, int *lx) {
68
69 static int autotune[13] = { 0 };
70
71 const dim3 nthrds_1d(1024, 1, 1);
72 const dim3 nblcks_1d((*nelv), 1, 1);
73 const dim3 nthrds_kstep((*lx), (*lx), 1);
74 const dim3 nblcks_kstep((*nelv), 1, 1);
75
76#define CASE_1D(LX) \
77 hipLaunchKernelGGL(HIP_KERNEL_NAME( ax_helm_kernel_1d<real, LX, 1024> ), \
78 nblcks_1d, nthrds_1d, 0, \
79 (hipStream_t) glb_cmd_queue, \
80 (real *) w, (real *) u, \
81 (real *) dx, (real *) dy, (real *) dz, \
82 (real *) dxt, (real *) dyt, (real *) dzt, (real *) h1, \
83 (real *) g11, (real *) g22, (real *) g33, \
84 (real *) g12, (real *) g13, (real *) g23); \
85 HIP_CHECK(hipGetLastError());
86
87#define CASE_KSTEP(LX) \
88 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_kernel_kstep<real, LX> ), \
89 nblcks_kstep, nthrds_kstep, 0, \
90 (hipStream_t) glb_cmd_queue, \
91 (real *) w, (real *) u, \
92 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
93 (real *) g11, (real *) g22, (real *) g33, \
94 (real *) g12, (real *) g13, (real *) g23); \
95 HIP_CHECK(hipGetLastError());
96
97
98
99#define CASE_KSTEP_PADDED(LX) \
100 hipLaunchKernelGGL( HIP_KERNEL_NAME(ax_helm_kernel_kstep_padded<real, LX> ),\
101 nblcks_kstep, nthrds_kstep, 0, \
102 (hipStream_t) glb_cmd_queue, \
103 (real *) w, (real *) u, \
104 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
105 (real *) g11, (real *) g22, (real *) g33, \
106 (real *) g12, (real *) g13, (real *) g23); \
107 HIP_CHECK(hipGetLastError());
108
109#define CASE(LX) \
110 case LX: \
111 if(autotune[LX] == 0 ) { \
112 autotune[LX]=tune<LX>( w, u, \
113 dx, dy, dz, \
114 dxt, dyt, dzt,h1, \
115 g11, g22, g33, \
116 g12, g13, g23, nelv, lx); \
117 } else if (autotune[LX] == 1 ) { \
118 CASE_1D(LX); \
119 } else if (autotune[LX] == 2 ) { \
120 CASE_KSTEP(LX); \
121 } \
122 break
123
124
125#define CASE_PADDED(LX) \
126 case LX: \
127 if(autotune[LX] == 0 ) { \
128 autotune[LX]=tune_padded<LX>(w, u, \
129 dx, dy, dz, \
130 dxt, dyt, dzt,h1, \
131 g11, g22, g33, \
132 g12, g13, g23,nelv,lx); \
133 } else if (autotune[LX] == 1 ) { \
134 CASE_1D(LX); \
135 } else if (autotune[LX] == 2 ) { \
136 CASE_KSTEP_PADDED(LX); \
137 } \
138 break
139
140#define CASE_LARGE(LX) \
141 case LX: \
142 CASE_KSTEP(LX); \
143 break
144
145#define CASE_LARGE_PADDED(LX) \
146 case LX: \
147 CASE_KSTEP_PADDED(LX); \
148 break
149
150 if ((*lx) < 13) {
151 switch(*lx) {
152 CASE(2);
153 CASE(3);
154 CASE(4);
155 CASE(5);
156 CASE(6);
157 CASE(7);
158 CASE_PADDED(8);
159 CASE(9);
160 CASE(10);
161 CASE(11);
162 CASE(12);
163 default:
164 {
165 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
166 exit(1);
167 }
168 }
169 }
170 else {
171 switch(*lx) {
172 CASE_LARGE(12);
173 CASE_LARGE(13);
174 CASE_LARGE(14);
175 CASE_LARGE(15);
176 CASE_LARGE_PADDED(16);
177 default:
178 {
179 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
180 exit(1);
181 }
182 }
183 }
184 }
188 void hip_ax_helm_vector(void *au, void *av, void *aw,
189 void *u, void *v, void *w,
190 void *dx, void *dy, void *dz,
191 void *dxt, void *dyt, void *dzt,
192 void *h1, void *g11, void *g22,
193 void *g33, void *g12, void *g13,
194 void *g23, int *nelv, int *lx) {
195
196 const dim3 nthrds((*lx), (*lx), 1);
197 const dim3 nblcks((*nelv), 1, 1);
198
199#define CASE_VECTOR_KSTEP(LX) \
200 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_kernel_vector_kstep<real, LX> ), \
201 nblcks, nthrds, 0, \
202 (hipStream_t) glb_cmd_queue, \
203 (real *) au, (real *) av, (real *) aw, \
204 (real *) u, (real *) v, (real *) w, \
205 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
206 (real *) g11, (real *) g22, (real *) g33, \
207 (real *) g12, (real *) g13, (real *) g23); \
208 HIP_CHECK(hipGetLastError());
209
210#define CASE_VECTOR_KSTEP_PADDED(LX) \
211 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_kernel_vector_kstep_padded<real, LX> ), \
212 nblcks, nthrds, 0, \
213 (hipStream_t) glb_cmd_queue, \
214 (real *) au, (real *) av, (real *) aw, \
215 (real *) u, (real *) v, (real *) w, \
216 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
217 (real *) g11, (real *) g22, (real *) g33, \
218 (real *) g12, (real *) g13, (real *) g23); \
219 HIP_CHECK(hipGetLastError());
220
221#define CASE_VECTOR(LX) \
222 case LX: \
223 CASE_VECTOR_KSTEP(LX); \
224 break
225
226#define CASE_VECTOR_PADDED(LX) \
227 case LX: \
228 CASE_VECTOR_KSTEP_PADDED(LX); \
229 break
230
231 switch(*lx) {
232 CASE_VECTOR(2);
233 CASE_VECTOR(3);
234 CASE_VECTOR_PADDED(4);
235 CASE_VECTOR(5);
236 CASE_VECTOR(6);
237 CASE_VECTOR(7);
238 CASE_VECTOR_PADDED(8);
239 CASE_VECTOR(9);
240 CASE_VECTOR(10);
241 CASE_VECTOR(11);
242 CASE_VECTOR(12);
243 CASE_VECTOR(13);
244 CASE_VECTOR(14);
245 CASE_VECTOR(15);
246 CASE_VECTOR_PADDED(16);
247 default:
248 {
249 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
250 exit(1);
251 }
252 }
253 }
254
258 void hip_ax_helm_vector_part2(void *au, void *av, void *aw,
259 void *u, void *v, void *w,
260 void *h2, void *B, int *n) {
261
262 const dim3 nthrds(1024, 1, 1);
263 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
264 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
265
266 hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_kernel_vector_part2<real> ),
267 nblcks, nthrds, 0, stream,
268 (real *) au, (real *) av, (real *) aw,
269 (real *) u, (real *) v, (real *) w,
270 (real *) h2, (real *) B, *n);
271 }
272
273}
274
275template < const int LX >
276int tune(void *w, void *u, void *dx, void *dy, void *dz,
277 void *dxt, void *dyt, void *dzt, void *h1,
278 void *g11, void *g22, void *g33, void *g12,
279 void *g13, void *g23, int *nelv, int *lx) {
280 hipEvent_t start,stop;
281 float time1,time2;
282 int retval;
283
284 const dim3 nthrds_1d(1024, 1, 1);
285 const dim3 nblcks_1d((*nelv), 1, 1);
286 const dim3 nthrds_kstep((*lx), (*lx), 1);
287 const dim3 nblcks_kstep((*nelv), 1, 1);
288
289 char *env_value = NULL;
290 char neko_log_buf[80];
291
292 env_value=getenv("NEKO_AUTOTUNE");
293
294 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
295 log_section(neko_log_buf);
296
297 if(env_value) {
298 if( !strcmp(env_value,"1D") ) {
299 CASE_1D(LX);
300 sprintf(neko_log_buf,"Set by env : 1 (1D)");
301 log_message(neko_log_buf);
302 log_end_section();
303 return 1;
304 } else if( !strcmp(env_value,"KSTEP") ) {
305 CASE_KSTEP(LX);
306 sprintf(neko_log_buf,"Set by env : 2 (KSTEP)");
307 log_message(neko_log_buf);
308 log_end_section();
309 return 2;
310 } else {
311 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
312 log_error(neko_log_buf);
313 }
314 }
315
316 HIP_CHECK(hipEventCreate(&start));
317 HIP_CHECK(hipEventCreate(&stop));
318
319 HIP_CHECK(hipEventRecord(start,0));
320
321 for(int i = 0; i < 100; i++) {
322 CASE_1D(LX);
323 }
324
325 HIP_CHECK(hipEventRecord(stop,0));
326 HIP_CHECK(hipEventSynchronize(stop));
327 HIP_CHECK(hipEventElapsedTime(&time1, start, stop));
328
329 HIP_CHECK(hipEventRecord(start,0));
330
331 for(int i = 0; i < 100; i++) {
332 CASE_KSTEP(LX);
333 }
334
335 HIP_CHECK(hipEventRecord(stop,0));
336 HIP_CHECK(hipEventSynchronize(stop));
337 HIP_CHECK(hipEventElapsedTime(&time2, start, stop));
338
339 if(time1 < time2) {
340 retval = 1;
341 } else {
342 retval = 2;
343 }
344
345 sprintf(neko_log_buf, "Chose : %d (%s)", retval,
346 (retval > 1 ? "KSTEP" : "1D"));
347 log_message(neko_log_buf);
348 log_end_section();
349 return retval;
350}
351
352template < const int LX >
353int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
354 void *dxt, void *dyt, void *dzt, void *h1,
355 void *g11, void *g22, void *g33, void *g12,
356 void *g13, void *g23, int *nelv, int *lx) {
357 hipEvent_t start, stop;
358 float time1, time2;
359 int retval;
360
361 const dim3 nthrds_1d(1024, 1, 1);
362 const dim3 nblcks_1d((*nelv), 1, 1);
363 const dim3 nthrds_kstep((*lx), (*lx), 1);
364 const dim3 nblcks_kstep((*nelv), 1, 1);
365
366 char *env_value = NULL;
367 char neko_log_buf[80];
368
369 env_value=getenv("NEKO_AUTOTUNE");
370
371 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
372 log_section(neko_log_buf);
373
374 if(env_value) {
375 if( !strcmp(env_value,"1D") ) {
376 CASE_1D(LX);
377 sprintf(neko_log_buf,"Set by env : 1 (1D)");
378 log_message(neko_log_buf);
379 log_end_section();
380 return 1;
381 } else if( !strcmp(env_value,"KSTEP") ) {
382 CASE_KSTEP(LX);
383 sprintf(neko_log_buf,"Set by env : 2 (KSTEP)");
384 log_message(neko_log_buf);
385 log_end_section();
386 return 2;
387 } else {
388 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
389 log_error(neko_log_buf);
390 }
391 }
392
393 HIP_CHECK(hipEventCreate(&start));
394 HIP_CHECK(hipEventCreate(&stop));
395
396 HIP_CHECK(hipEventRecord(start,0));
397
398 for(int i = 0; i < 100; i++) {
399 CASE_1D(LX);
400 }
401
402 HIP_CHECK(hipEventRecord(stop, 0));
403 HIP_CHECK(hipEventSynchronize(stop));
404 HIP_CHECK(hipEventElapsedTime(&time1, start, stop));
405
406 HIP_CHECK(hipEventRecord(start, 0));
407
408 for(int i = 0; i < 100; i++) {
409 CASE_KSTEP_PADDED(LX);
410 }
411
412 HIP_CHECK(hipEventRecord(stop, 0));
413 HIP_CHECK(hipEventSynchronize(stop));
414 HIP_CHECK(hipEventElapsedTime(&time2, start, stop));
415
416 if(time1 < time2) {
417 retval=1;
418 } else {
419 retval=2;
420 }
421
422 sprintf(neko_log_buf, "Chose : %d (%s)", retval,
423 (retval > 1 ? "KSTEP" : "1D"));
424 log_message(neko_log_buf);
425 log_end_section();
426 return retval;
427}
CASE_VECTOR
#define CASE_VECTOR(LX)
CASE
#define CASE(LX)
CASE_KSTEP
#define CASE_KSTEP(LX)
hip_ax_helm_vector_part2
void hip_ax_helm_vector_part2(void *au, void *av, void *aw, void *u, void *v, void *w, void *h2, void *B, int *n)
Definition ax_helm.hip:258
tune
int tune(void *w, void *u, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.hip:276
CASE_PADDED
#define CASE_PADDED(LX)
CASE_LARGE
#define CASE_LARGE(LX)
CASE_1D
#define CASE_1D(LX)
CASE_KSTEP_PADDED
#define CASE_KSTEP_PADDED(LX)
hip_ax_helm_vector
void hip_ax_helm_vector(void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.hip:188
hip_ax_helm
void hip_ax_helm(void *w, void *u, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.hip:64
CASE_LARGE_PADDED
#define CASE_LARGE_PADDED(LX)
tune_padded
int tune_padded(void *w, void *u, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *g11, void *g22, void *g33, void *g12, void *g13, void *g23, int *nelv, int *lx)
Definition ax_helm.hip:353
CASE_VECTOR_PADDED
#define CASE_VECTOR_PADDED(LX)
aw
__global__ void T *__restrict__ T *__restrict__ aw
Definition ax_helm_full_kernel.h:45
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
i
const int i
Definition ax_helm_full_kernel.h:95
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition ax_helm_full_kernel.h:49
av
__global__ void T *__restrict__ av
Definition ax_helm_full_kernel.h:44
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
dz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition ax_helm_full_kernel.h:51
h1
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ h1
Definition ax_helm_full_kernel.h:52
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition ax_helm_full_kernel.h:50
dzt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
Definition cdtp_kernel.h:112
dyt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
Definition cdtp_kernel.h:111
dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
Definition cdtp_kernel.h:110
dirichlet_apply_scalar_kernel
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
Definition dirichlet_kernel.h:42
device_config.h
real
double real
Definition device_config.h:12
ax_helm_kernel.h
g23
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g23
Definition ax_helm_kernel.h:161
g22
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g22
Definition ax_helm_kernel.h:157
g13
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g13
Definition ax_helm_kernel.h:160
g12
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g12
Definition ax_helm_kernel.h:159
g33
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g33
Definition ax_helm_kernel.h:158
g11
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ g11
Definition ax_helm_kernel.h:156
HIP_CHECK
#define HIP_CHECK(err)
Definition check.h:8
neko_log.h
log_error
void log_error(char *msg)
log_message
void log_message(char *msg)
log_end_section
void log_end_section()
log_section
void log_section(char *msg)