Neko 0.9.1
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
ax_helm.cu
Go to the documentation of this file.
1/*
2 Copyright (c) 2021-2024, 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 "ax_helm_kernel.h"
39#include <device/device_config.h>
40#include <device/cuda/check.h>
41
42extern "C" {
43 #include <common/neko_log.h>
44}
45
46template < const int>
47int tune(void *w, void *u, void *dx, void *dy, void *dz,
48 void *dxt, void *dyt, void *dzt, void *h1,
49 void *g11, void *g22, void *g33, void *g12,
50 void *g13, void *g23, int *nelv, int *lx);
51
52template < const int>
53int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
54 void *dxt, void *dyt, void *dzt, void *h1,
55 void *g11, void *g22, void *g33, void *g12,
56 void *g13, void *g23, int *nelv, int *lx);
57
58extern "C" {
59
63 void cuda_ax_helm(void *w, void *u, void *dx, void *dy, void *dz,
64 void *dxt, void *dyt, void *dzt, void *h1,
65 void *g11, void *g22, void *g33, void *g12,
66 void *g13, void *g23, int *nelv, int *lx) {
67
68 static int autotune[17] = { 0 };
69
70 const dim3 nthrds_1d(1024, 1, 1);
71 const dim3 nblcks_1d((*nelv), 1, 1);
72 const dim3 nthrds_kstep((*lx), (*lx), 1);
73 const dim3 nblcks_kstep((*nelv), 1, 1);
74
75 const dim3 nthrds((*lx), (*lx), 1);
76 const dim3 nblcks((*nelv), 1, 1);
77 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
78
79#define CASE_1D(LX) \
80 ax_helm_kernel_1d<real, LX, 1024> \
81 <<<nblcks_1d, nthrds_1d, 0, stream>>>((real *) w, (real *) u, \
82 (real *) dx, (real *) dy, (real *) dz, \
83 (real *) dxt, (real *) dyt, (real *) dzt, (real *) h1,\
84 (real *) g11, (real *) g22, (real *) g33, \
85 (real *) g12, (real *) g13, (real *) g23); \
86 CUDA_CHECK(cudaGetLastError());
87
88#define CASE_KSTEP(LX) \
89 ax_helm_kernel_kstep<real, LX> \
90 <<<nblcks_kstep, nthrds_kstep, 0, stream>>>((real *) w, (real *) u, \
91 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
92 (real *) g11, (real *) g22, (real *) g33, \
93 (real *) g12, (real *) g13, (real *) g23); \
94 CUDA_CHECK(cudaGetLastError());
95
96#define CASE_KSTEP_PADDED(LX) \
97 ax_helm_kernel_kstep_padded<real, LX> \
98 <<<nblcks_kstep, nthrds_kstep, 0, stream>>>((real *) w, (real *) u, \
99 (real *) dx, (real *) dy, (real *) dz, (real *) h1, \
100 (real *) g11, (real *) g22, (real *) g33, \
101 (real *) g12, (real *) g13, (real *) g23); \
102 CUDA_CHECK(cudaGetLastError());
103
104#define CASE(LX) \
105 case LX: \
106 if(autotune[LX] == 0 ) { \
107 autotune[LX]=tune<LX>( w, u, \
108 dx, dy, dz, \
109 dxt, dyt, dzt,h1, \
110 g11, g22, g33, \
111 g12, g13, g23, nelv, lx); \
112 } else if (autotune[LX] == 1 ) { \
113 CASE_1D(LX); \
114 } else if (autotune[LX] == 2 ) { \
115 CASE_KSTEP(LX); \
116 } \
117 break
118
119
120#define CASE_PADDED(LX) \
121 case LX: \
122 if(autotune[LX] == 0 ) { \
123 autotune[LX]=tune_padded<LX>(w, u, \
124 dx, dy, dz, \
125 dxt, dyt, dzt,h1, \
126 g11, g22, g33, \
127 g12, g13, g23,nelv,lx); \
128 } else if (autotune[LX] == 1 ) { \
129 CASE_1D(LX); \
130 } else if (autotune[LX] == 2 ) { \
131 CASE_KSTEP_PADDED(LX); \
132 } \
133 break
134
135#define CASE_LARGE(LX) \
136 case LX: \
137 CASE_KSTEP(LX); \
138 break
139
140#define CASE_LARGE_PADDED(LX) \
141 case LX: \
142 CASE_KSTEP_PADDED(LX); \
143 break
144
145
146
147 if ((*lx) < 12) {
148 switch(*lx) {
149 CASE(2);
150 CASE(3);
151 CASE_PADDED(4);
152 CASE(5);
153 CASE(6);
154 CASE(7);
155 CASE_PADDED(8);
156 CASE(9);
157 CASE(10);
158 CASE(11);
159 default:
160 {
161 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
162 exit(1);
163 }
164 }
165 }
166 else {
167 switch(*lx) {
168 CASE_LARGE(12);
169 CASE_LARGE(13);
170 CASE_LARGE(14);
171 CASE_LARGE(15);
172 CASE_LARGE_PADDED(16);
173 default:
174 {
175 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
176 exit(1);
177 }
178 }
179 }
180 }
181
185 void cuda_ax_helm_vector(void *au, void *av, void *aw,
186 void *u, void *v, void *w,
187 void *dx, void *dy, void *dz,
188 void *dxt, void *dyt, void *dzt,
189 void *h1, void *g11, void *g22,
190 void *g33, void *g12, void *g13,
191 void *g23, int *nelv, int *lx) {
192
193 const dim3 nthrds((*lx), (*lx), 1);
194 const dim3 nblcks((*nelv), 1, 1);
195 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
196
197#define CASE_VECTOR_KSTEP(LX) \
198 ax_helm_kernel_vector_kstep<real, LX> \
199 <<<nblcks, nthrds, 0, stream>>> ((real *) au, (real *) av, (real *) aw, \
200 (real *) u, (real *) v, (real *) w, \
201 (real *) dx, (real *) dy, (real *) dz, \
202 (real *) h1, (real *) g11, (real *) g22, \
203 (real *) g33, (real *) g12, (real *) g13, \
204 (real *) g23); \
205 CUDA_CHECK(cudaGetLastError());
206
207#define CASE_VECTOR_KSTEP_PADDED(LX) \
208 ax_helm_kernel_vector_kstep_padded<real, LX> \
209 <<<nblcks, nthrds, 0, stream>>> ((real *) au, (real *) av, (real *) aw, \
210 (real *) u, (real *) v, (real *) w, \
211 (real *) dx, (real *) dy, (real *) dz, \
212 (real *) h1, (real *) g11, (real *) g22, \
213 (real *) g33, (real *) g12, (real *) g13, \
214 (real *) g23); \
215 CUDA_CHECK(cudaGetLastError());
216
217
218#define CASE_VECTOR(LX) \
219 case LX: \
220 CASE_VECTOR_KSTEP(LX); \
221 break
222
223#define CASE_VECTOR_PADDED(LX) \
224 case LX: \
225 CASE_VECTOR_KSTEP_PADDED(LX); \
226 break
227
228 switch(*lx) {
229 CASE_VECTOR(2);
230 CASE_VECTOR(3);
231 CASE_VECTOR_PADDED(4);
232 CASE_VECTOR(5);
233 CASE_VECTOR(6);
234 CASE_VECTOR(7);
235 CASE_VECTOR_PADDED(8);
236 CASE_VECTOR(9);
237 CASE_VECTOR(10);
238 CASE_VECTOR(11);
239 CASE_VECTOR(12);
240 CASE_VECTOR(13);
241 CASE_VECTOR(14);
242 CASE_VECTOR(15);
243 CASE_VECTOR_PADDED(16);
244 default:
245 {
246 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
247 exit(1);
248 }
249 }
250 }
251
255 void cuda_ax_helm_vector_part2(void *au, void *av, void *aw,
256 void *u, void *v, void *w,
257 void *h2, void *B, int *n) {
258
259 const dim3 nthrds(1024, 1, 1);
260 const dim3 nblcks(((*n)+1024 - 1)/ 1024, 1, 1);
261 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
262
263 ax_helm_kernel_vector_part2<real>
264 <<<nblcks, nthrds, 0, stream>>> ((real *) au, (real *) av, (real *) aw,
265 (real *) u, (real *) v, (real *) w,
266 (real *) h2, (real *) B, *n);
267 }
268}
269
270template < const int LX >
271int tune(void *w, void *u, void *dx, void *dy, void *dz,
272 void *dxt, void *dyt, void *dzt, void *h1,
273 void *g11, void *g22, void *g33, void *g12,
274 void *g13, void *g23, int *nelv, int *lx) {
275 cudaEvent_t start,stop;
276 float time1,time2;
277 int retval;
278
279 const dim3 nthrds_1d(1024, 1, 1);
280 const dim3 nblcks_1d((*nelv), 1, 1);
281 const dim3 nthrds_kstep((*lx), (*lx), 1);
282 const dim3 nblcks_kstep((*nelv), 1, 1);
283 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
284
285 char *env_value = NULL;
286 char neko_log_buf[80];
287
288 env_value=getenv("NEKO_AUTOTUNE");
289
290 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
291 log_section(neko_log_buf);
292
293 if(env_value) {
294 if( !strcmp(env_value,"1D") ) {
295 CASE_1D(LX);
296 sprintf(neko_log_buf,"Set by env : 1 (1D)");
297 log_message(neko_log_buf);
298 log_end_section();
299 return 1;
300 } else if( !strcmp(env_value,"KSTEP") ) {
301 CASE_KSTEP(LX);
302 sprintf(neko_log_buf,"Set by env : 2 (KSTEP)");
303 log_message(neko_log_buf);
304 log_end_section();
305 return 2;
306 } else {
307 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
308 log_error(neko_log_buf);
309 }
310 }
311
312 cudaEventCreate(&start);
313 cudaEventCreate(&stop);
314
315 cudaEventRecord(start,0);
316
317 for(int i = 0; i < 100; i++) {
318 CASE_1D(LX);
319 }
320
321 cudaEventRecord(stop,0);
322 cudaEventSynchronize(stop);
323 cudaEventElapsedTime(&time1, start, stop);
324
325 cudaEventRecord(start,0);
326
327 for(int i = 0; i < 100; i++) {
328 CASE_KSTEP(LX);
329 }
330
331 cudaEventRecord(stop,0);
332 cudaEventSynchronize(stop);
333 cudaEventElapsedTime(&time2, start, stop);
334
335 if(time1 < time2) {
336 retval = 1;
337 } else {
338 retval = 2;
339 }
340
341 sprintf(neko_log_buf, "Chose : %d (%s)", retval,
342 (retval > 1 ? "KSTEP" : "1D"));
343 log_message(neko_log_buf);
344 log_end_section();
345 return retval;
346}
347
348template < const int LX >
349int tune_padded(void *w, void *u, void *dx, void *dy, void *dz,
350 void *dxt, void *dyt, void *dzt, void *h1,
351 void *g11, void *g22, void *g33, void *g12,
352 void *g13, void *g23, int *nelv, int *lx) {
353 cudaEvent_t start, stop;
354 float time1, time2;
355 int retval;
356
357 const dim3 nthrds_1d(1024, 1, 1);
358 const dim3 nblcks_1d((*nelv), 1, 1);
359 const dim3 nthrds_kstep((*lx), (*lx), 1);
360 const dim3 nblcks_kstep((*nelv), 1, 1);
361 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
362
363 char *env_value = NULL;
364 char neko_log_buf[80];
365
366 env_value=getenv("NEKO_AUTOTUNE");
367
368 sprintf(neko_log_buf, "Autotune Ax helm (lx: %d)", *lx);
369 log_section(neko_log_buf);
370
371 if(env_value) {
372 if( !strcmp(env_value,"1D") ) {
373 CASE_1D(LX);
374 sprintf(neko_log_buf,"Set by env : 1 (1D)");
375 log_message(neko_log_buf);
376 log_end_section();
377 return 1;
378 } else if( !strcmp(env_value,"KSTEP") ) {
379 CASE_KSTEP(LX);
380 sprintf(neko_log_buf,"Set by env : 2 (KSTEP)");
381 log_message(neko_log_buf);
382 log_end_section();
383 return 2;
384 } else {
385 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
386 log_error(neko_log_buf);
387 }
388 }
389
390 cudaEventCreate(&start);
391 cudaEventCreate(&stop);
392
393 cudaEventRecord(start,0);
394
395 for(int i = 0; i < 100; i++) {
396 CASE_1D(LX);
397 }
398
399 cudaEventRecord(stop, 0);
400 cudaEventSynchronize(stop);
401 cudaEventElapsedTime(&time1, start, stop);
402
403 cudaEventRecord(start, 0);
404
405 for(int i = 0; i < 100; i++) {
406 CASE_KSTEP_PADDED(LX);
407 }
408
409 cudaEventRecord(stop, 0);
410 cudaEventSynchronize(stop);
411 cudaEventElapsedTime(&time2, start, stop);
412
413 if(time1 < time2) {
414 retval=1;
415 } else {
416 retval=2;
417 }
418
419 sprintf(neko_log_buf, "Chose : %d (%s)", retval,
420 (retval > 1 ? "KSTEP" : "1D"));
421 log_message(neko_log_buf);
422 log_end_section();
423 return retval;
424}
CASE_VECTOR
#define CASE_VECTOR(LX)
cuda_ax_helm_vector_part2
void cuda_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.cu:255
CASE
#define CASE(LX)
CASE_KSTEP
#define CASE_KSTEP(LX)
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.cu:271
cuda_ax_helm
void cuda_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.cu:63
CASE_PADDED
#define CASE_PADDED(LX)
cuda_ax_helm_vector
void cuda_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.cu:185
CASE_LARGE
#define CASE_LARGE(LX)
CASE_1D
#define CASE_1D(LX)
CASE_KSTEP_PADDED
#define CASE_KSTEP_PADDED(LX)
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.cu:349
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
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)