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