Neko 1.99.3
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-2026, 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 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
289
290 char *env_value = NULL;
291 char neko_log_buf[80];
292
293 env_value=getenv("NEKO_AUTOTUNE");
294
295 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
296 log_section(neko_log_buf);
297
298 if(env_value) {
299 if( !strcmp(env_value,"1D") ) {
300 CASE_1D(LX);
301 sprintf(neko_log_buf,"Set by env : 1 (1D)");
302 log_message(neko_log_buf);
303 log_end_section();
304 return 1;
305 } else if( !strcmp(env_value,"KSTEP") ) {
306 CASE_KSTEP(LX);
307 sprintf(neko_log_buf,"Set by env : 2 (KSTEP)");
308 log_message(neko_log_buf);
309 log_end_section();
310 return 2;
311 } else {
312 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
313 log_error(neko_log_buf);
314 }
315 }
316
317 HIP_CHECK(hipEventCreate(&start));
318 HIP_CHECK(hipEventCreate(&stop));
319
320 /* Warmup */
321 for(int i = 0; i < 10; i++) {
322 CASE_1D(LX);
323 }
324
325 HIP_CHECK(hipEventRecord(start, stream));
326
327 for(int i = 0; i < 100; i++) {
328 CASE_1D(LX);
329 }
330
331 HIP_CHECK(hipEventRecord(stop, stream));
332 HIP_CHECK(hipEventSynchronize(stop));
333 HIP_CHECK(hipEventElapsedTime(&time1, start, stop));
334
335 HIP_CHECK(hipEventRecord(start, stream));
336
337 for(int i = 0; i < 100; i++) {
338 CASE_KSTEP(LX);
339 }
340
341 HIP_CHECK(hipEventRecord(stop, stream));
342 HIP_CHECK(hipEventSynchronize(stop));
343 HIP_CHECK(hipEventElapsedTime(&time2, start, stop));
344
345 if(time1 < time2) {
346 retval = 1;
347 } else {
348 retval = 2;
349 }
350
351 sprintf(neko_log_buf, "Chose : %d (%s)", retval,
352 (retval > 1 ? "KSTEP" : "1D"));
353 log_message(neko_log_buf);
354 log_end_section();
355 return retval;
356}
357
358template < const int LX >
359int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
360 void *dxt, void *dyt, void *dzt, void *h1,
361 void *g11, void *g22, void *g33, void *g12,
362 void *g13, void *g23, int *nelv, int *lx) {
363 hipEvent_t start, stop;
364 float time1, time2;
365 int retval;
366
367 const dim3 nthrds_1d(1024, 1, 1);
368 const dim3 nblcks_1d((*nelv), 1, 1);
369 const dim3 nthrds_kstep((*lx), (*lx), 1);
370 const dim3 nblcks_kstep((*nelv), 1, 1);
371 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
372
373 char *env_value = NULL;
374 char neko_log_buf[80];
375
376 env_value=getenv("NEKO_AUTOTUNE");
377
378 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
379 log_section(neko_log_buf);
380
381 if(env_value) {
382 if( !strcmp(env_value,"1D") ) {
383 CASE_1D(LX);
384 sprintf(neko_log_buf,"Set by env : 1 (1D)");
385 log_message(neko_log_buf);
386 log_end_section();
387 return 1;
388 } else if( !strcmp(env_value,"KSTEP") ) {
389 CASE_KSTEP(LX);
390 sprintf(neko_log_buf,"Set by env : 2 (KSTEP)");
391 log_message(neko_log_buf);
392 log_end_section();
393 return 2;
394 } else {
395 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
396 log_error(neko_log_buf);
397 }
398 }
399
400 HIP_CHECK(hipEventCreate(&start));
401 HIP_CHECK(hipEventCreate(&stop));
402
403 /* Warmup */
404 for(int i = 0; i < 10; i++) {
405 CASE_1D(LX);
406 }
407
408 HIP_CHECK(hipEventRecord(start, stream));
409
410 for(int i = 0; i < 100; i++) {
411 CASE_1D(LX);
412 }
413
414 HIP_CHECK(hipEventRecord(stop, stream));
415 HIP_CHECK(hipEventSynchronize(stop));
416 HIP_CHECK(hipEventElapsedTime(&time1, start, stop));
417
418 HIP_CHECK(hipEventRecord(start, stream));
419
420 for(int i = 0; i < 100; i++) {
421 CASE_KSTEP_PADDED(LX);
422 }
423
424 HIP_CHECK(hipEventRecord(stop, stream));
425 HIP_CHECK(hipEventSynchronize(stop));
426 HIP_CHECK(hipEventElapsedTime(&time2, start, stop));
427
428 if(time1 < time2) {
429 retval=1;
430 } else {
431 retval=2;
432 }
433
434 sprintf(neko_log_buf, "Chose : %d (%s)", retval,
435 (retval > 1 ? "KSTEP" : "1D"));
436 log_message(neko_log_buf);
437 log_end_section();
438 return retval;
439}
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:359
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)