Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
opr_lambda2.cu
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 "lambda2_kernel.h"
39#include "elem_block_tune.h"
40#include <device/device_config.h>
41#include <device/cuda/check.h>
42
43extern "C" {
44 #include <common/neko_log.h>
45}
46
47template < const int >
48int tune_lambda2(void *ux, void *uy, void *uz, void *u,
49 void *dx, void *dy, void *dz,
50 void *drdx, void *dsdx, void *dtdx,
51 void *drdy, void *dsdy, void *dtdy,
52 void *drdz, void *dsdz, void *dtdz,
53 void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel);
54
55extern "C" {
56
60 void cuda_lambda2(void *lambda2, void *u, void *v, void *w,
61 void *dx, void *dy, void *dz,
62 void *drdx, void *dsdx, void *dtdx,
63 void *drdy, void *dsdy, void *dtdy,
64 void *drdz, void *dsdz, void *dtdz,
65 void *jacinv, int *nel, int *lx) {
66
67 static int autotune[17] = { 0 };
68 /* elements per block candidate chosen by the tuner */
69 static int autotune_eb[17] = { 0 };
70 /* chunk candidate chosen for the 1d variant */
71 static int autotune_ch[17] = { 0 };
72
73 const dim3 nthrds_1d(1024, 1, 1);
74 const dim3 nthrds_kstep((*lx), (*lx), 1);
75 const dim3 nblcks((*nel), 1, 1);
76 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
77
78#define CASE_1D(LX, C) \
79 lambda2_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
80 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
81 ((real *) lambda2, (real *) u, (real *) v, (real *) w, \
82 (real *) dx, (real *) dy, (real *) dz, \
83 (real *) drdx, (real *) dsdx, (real *) dtdx, \
84 (real *) drdy, (real *) dsdy, (real *) dtdy, \
85 (real *) drdz, (real *) dsdz, (real *) dtdz, \
86 (real *) jacinv); \
87 CUDA_CHECK(cudaGetLastError());
88
89/* Runtime dispatch onto the tuned chunk candidate */
90#define CASE_1D_SEL(LX, SEL) \
91 switch (SEL) { \
92 case 0: CASE_1D(LX, 0); break; \
93 case 1: CASE_1D(LX, 1); break; \
94 case 2: CASE_1D(LX, 2); break; \
95 default: CASE_1D(LX, 3); break; \
96 }
97
98
99#define CASE_KSTEP(LX, C) \
100 lambda2_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
101 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
102 ((real *) lambda2, (real *) u, (real *) v, (real *) w, \
103 (real *) dx, (real *) dy, (real *) dz, \
104 (real *) drdx, (real *) dsdx, (real *) dtdx, \
105 (real *) drdy, (real *) dsdy, (real *) dtdy, \
106 (real *) drdz, (real *) dsdz, (real *) dtdz, \
107 (real *) jacinv, *nel); \
108 CUDA_CHECK(cudaGetLastError());
109
110/* Runtime dispatch onto the tuned candidate */
111#define CASE_KSTEP_SEL(LX, SEL) \
112 switch (SEL) { \
113 case 0: CASE_KSTEP(LX, 0); break; \
114 case 1: CASE_KSTEP(LX, 1); break; \
115 default: CASE_KSTEP(LX, 2); break; \
116 }
117
118#define CASE(LX) \
119 case LX: \
120 if(autotune[LX] == 0 ) { \
121 autotune[LX]=tune_lambda2<LX>(lambda2, u, v, w, \
122 dx, dy, dz, \
123 drdx, dsdx, dtdx, \
124 drdy, dsdy, dtdy, \
125 drdz, dsdz, dtdz, \
126 jacinv, nel, lx, &autotune_eb[LX], \
127 &autotune_ch[LX]); \
128 } else if (autotune[LX] == 1 ) { \
129 CASE_1D_SEL(LX, autotune_ch[LX]); \
130 } else if (autotune[LX] == 2 ) { \
131 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
132 } \
133 break
134
135 switch(*lx) {
136 CASE(2);
137 CASE(3);
138 CASE(4);
139 CASE(5);
140 CASE(6);
141 CASE(7);
142 CASE(8);
143 CASE(9);
144 CASE(10);
145 CASE(11);
146 CASE(12);
147 default:
148 {
149 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
150 exit(1);
151 }
152 }
153 }
154}
155
156template < const int LX >
157int tune_lambda2(void *lambda2, void *u, void *v, void *w,
158 void *dx, void *dy, void *dz,
159 void *drdx, void *dsdx, void *dtdx,
160 void *drdy, void *dsdy, void *dtdy,
161 void *drdz, void *dsdz, void *dtdz,
162 void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel) {
163 cudaEvent_t start,stop;
164 float time1[NEKO_CHUNKS_CANDIDATES];
165 int best1 = 0;
166 float time2[NEKO_EB_CANDIDATES];
167 const int rounds = neko_tune_rounds();
168 const int iters = neko_tune_iters();
169 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
170 elements per block sweep is off */
171 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
172 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
173 const int ch_pin = neko_chunks_pin();
174 const int eb_pin = neko_eb_pin();
175 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
176 int strat = 0;
177 int best = 0;
178 int retval;
179
180 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
181 time2[c] = NEKO_TUNE_INIT;
182 }
183 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
184 time1[c] = NEKO_TUNE_INIT;
185 }
186
187 const dim3 nthrds_1d(1024, 1, 1);
188 const dim3 nthrds_kstep((*lx), (*lx), 1);
189 const dim3 nblcks((*nel), 1, 1);
190 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
191
192 char *env_value = NULL;
193 char neko_log_buf[80];
194
195 env_value=getenv("NEKO_AUTOTUNE");
196
197 sprintf(neko_log_buf, "Autotune lambda2 (lx: %d)", *lx);
198 log_section(neko_log_buf);
199
200 /*
201 * NEKO_AUTOTUNE names a formulation, and that is all it does: the sweep
202 * below is narrowed to that one kernel family, but its geometry -- the
203 * chunk size, the elements per block -- is still measured candidate
204 * against candidate, and reported.
205 */
206 if(env_value) {
207 if( !strcmp(env_value,"1D") ) {
208 strat = 1;
209 } else if( !strcmp(env_value,"KSTEP") ) {
210 strat = 2;
211 } else {
212 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
213 log_error(neko_log_buf);
214 }
215 }
216
217 /* Geometry of the pinned formulation, if its own variable fixes that too.
218 Both pinned leaves nothing to measure, so the kernel is launched once and
219 reported, which is what pinning has always done */
220 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin : -1;
221
222 if (pin >= 0) {
223 switch (strat) {
224 case 1:
225 *ch_sel = pin;
226 CASE_1D_SEL(LX, pin);
227 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
228 NEKO_CHUNKS_SEL(LX, pin));
229 break;
230 default:
231 *eb_sel = pin;
232 CASE_KSTEP_SEL(LX, pin);
233 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
234 NEKO_EB_SEL(LX, pin));
235 break;
236 }
237 log_message(neko_log_buf);
238 log_end_section();
239 return strat;
240 }
241
242 if (strat) {
243 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
244 log_message(neko_log_buf);
245 }
246
247 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
248 const bool try_1d = (strat == 0 || strat == 1);
249 const bool try_kstep = (strat == 0 || strat == 2);
250
251 cudaEventCreate(&start);
252 cudaEventCreate(&stop);
253 /* Warm every variant before timing anything: each specialisation has to be
254 resident and the clocks at steady state, or whichever is timed first is
255 measured on a colder part */
256 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
257 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
258 CASE_1D_SEL(LX, c);
259 }
260 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
261 CASE_KSTEP_SEL(LX, c);
262 }
263 }
264
265 /* Interleaved rounds, best time per variant */
266 for (int r = 0; r < rounds; r++) {
267 NEKO_TUNE_FOR(c, try_1d, ch_pin, NEKO_CHUNKS_CANDIDATES) {
268 NEKO_TUNE_TIME(time1, CASE_1D_SEL, LX, c, iters);
269 }
270 NEKO_TUNE_FOR(c, try_kstep, eb_pin, eb_cand) {
271 NEKO_TUNE_TIME(time2, CASE_KSTEP_SEL, LX, c, iters);
272 }
273 }
274
275 NEKO_TUNE_LOG(LX, time1, time2);
276 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
277 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
278 *eb_sel = best;
279 *ch_sel = best1;
280
281 if (time1[best1] < time2[best]) {
282 retval = 1;
283 } else {
284 retval = 2;
285 }
286
287 /* Leave the chosen kernel's output in place: the tuner stands in for a real
288 evaluation and the variants do not sum in the same order */
289 if (retval == 1) {
290 CASE_1D_SEL(LX, best1);
291 } else {
292 CASE_KSTEP_SEL(LX, best);
293 }
294
295 if (retval == 1) {
296 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
297 NEKO_CHUNKS_SEL(LX, best1));
298 } else {
299 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
300 NEKO_EB_SEL(LX, best));
301 }
302 log_message(neko_log_buf);
303 log_end_section();
304 return retval;
305}
306
ale_add_kinematics_kernel
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)
Definition ale_kinematics_kernel.h:52
dtdy
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdy
Definition ax_helm_full_kernel.h:60
dtdx
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdx
Definition ax_helm_full_kernel.h:59
dtdz
__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__ 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__ dtdz
Definition ax_helm_full_kernel.h:61
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
jacinv
__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__ 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__ jacinv
Definition ax_helm_full_kernel.h:62
i
const int i
Definition ax_helm_full_kernel.h:95
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
dsdz
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdz
Definition ax_helm_full_kernel.h:58
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
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
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
drdz
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdz
Definition ax_helm_full_kernel.h:55
drdx
__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__ const T *__restrict__ drdx
Definition ax_helm_full_kernel.h:53
dsdx
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdx
Definition ax_helm_full_kernel.h:56
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition ax_helm_full_kernel.h:47
dsdy
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdy
Definition ax_helm_full_kernel.h:57
drdy
__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__ const T *__restrict__ const T *__restrict__ drdy
Definition ax_helm_full_kernel.h:54
NEKO_CHUNKS_CANDIDATES
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
NEKO_EB_CANDIDATES
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
NEKO_EB_SEL
#define NEKO_EB_SEL(LX, SEL)
Definition elem_block.h:108
NEKO_CHUNKS_SEL
#define NEKO_CHUNKS_SEL(LX, SEL)
Definition elem_block.h:147
neko_chunks_pin
static int neko_chunks_pin()
Definition elem_block_tune.h:112
NEKO_TUNE_TIME
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
Definition elem_block_tune.h:163
neko_tune_rounds
static int neko_tune_rounds()
Definition elem_block_tune.h:128
neko_eb_pin
static int neko_eb_pin()
Definition elem_block_tune.h:95
NEKO_TUNE_LOG
#define NEKO_TUNE_LOG(LX, T1, T2)
Definition elem_block_tune.h:193
NEKO_TUNE_INIT
#define NEKO_TUNE_INIT
Definition elem_block_tune.h:72
NEKO_TUNE_BEST
#define NEKO_TUNE_BEST(T, BEST, N)
Definition elem_block_tune.h:185
neko_tune_iters
static int neko_tune_iters()
Definition elem_block_tune.h:136
NEKO_TUNE_FOR
#define NEKO_TUNE_FOR(C, ON, PIN, N)
Definition elem_block_tune.h:157
neko_eb_sweep
static int neko_eb_sweep()
Definition elem_block_tune.h:75
NEKO_TUNE_WARMUP
#define NEKO_TUNE_WARMUP
Definition elem_block_tune.h:73
uy
__global__ void T *__restrict__ uy
Definition opgrad_kernel.h:130
uz
__global__ void T *__restrict__ T *__restrict__ uz
Definition opgrad_kernel.h:131
device_config.h
elem_block_tune.h
lambda2_kernel.h
lambda2
A simulation component that computes lambda2 The values are stored in the field registry under the na...
Definition lambda2.f90:37
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)
CASE_KSTEP_SEL
#define CASE_KSTEP_SEL(LX, SEL)
CASE
#define CASE(LX)
CASE_1D_SEL
#define CASE_1D_SEL(LX, SEL)
tune_lambda2
int tune_lambda2(void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel)
Definition opr_lambda2.cu:157
cuda_lambda2
void cuda_lambda2(void *lambda2, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *jacinv, int *nel, int *lx)
Definition opr_lambda2.cu:60