Neko 1.99.7
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 const int sweep = neko_eb_sweep();
176 int best = 0;
177 int retval;
178
179 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
180 time2[c] = NEKO_TUNE_INIT;
181 }
182 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
183 time1[c] = NEKO_TUNE_INIT;
184 }
185
186 const dim3 nthrds_1d(1024, 1, 1);
187 const dim3 nthrds_kstep((*lx), (*lx), 1);
188 const dim3 nblcks((*nel), 1, 1);
189 const hipStream_t stream = (hipStream_t) glb_cmd_queue;
190
191 char *env_value = NULL;
192 char neko_log_buf[80];
193
194 env_value=getenv("NEKO_AUTOTUNE");
195
196 sprintf(neko_log_buf, "Autotune lambda2 (lx: %d)", *lx);
197 log_section(neko_log_buf);
198
199 if(env_value) {
200 if( !strcmp(env_value,"1D") ) {
201 *ch_sel = neko_chunks_env();
202 CASE_1D_SEL(LX, *ch_sel);
203 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
204 NEKO_CHUNKS_SEL(LX, *ch_sel));
205 log_message(neko_log_buf);
206 log_end_section();
207 return 1;
208 } else if( !strcmp(env_value,"KSTEP") ) {
209 *eb_sel = neko_eb_env();
210 CASE_KSTEP_SEL(LX, *eb_sel);
211 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
212 NEKO_EB_SEL(LX, *eb_sel));
213 log_message(neko_log_buf);
214 log_end_section();
215 return 2;
216 } else {
217 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
218 log_error(neko_log_buf);
219 }
220 }
221
222 HIP_CHECK(hipEventCreate(&start));
223 HIP_CHECK(hipEventCreate(&stop));
224 /* Warm every variant before timing anything: each specialisation has to be
225 resident and the clocks at steady state, or whichever is timed first is
226 measured on a colder part */
227 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
228 CASE_1D(LX, 0);
229 CASE_1D(LX, 1);
230 CASE_1D(LX, 2);
231 CASE_1D(LX, 3);
232 CASE_KSTEP(LX, 0);
233 if (sweep) {
234 CASE_KSTEP(LX, 1);
235 CASE_KSTEP(LX, 2);
236 }
237 }
238
239 /* Interleaved rounds, best time per variant */
240 for (int r = 0; r < rounds; r++) {
241 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
242 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
243 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
244 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
245 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 0, iters);
246 if (sweep) {
247 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 1, iters);
248 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 2, iters);
249 }
250 }
251
252 NEKO_TUNE_LOG(LX, time1, time2);
253 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
254 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
255 *eb_sel = best;
256 *ch_sel = best1;
257
258 if (time1[best1] < time2[best]) {
259 retval = 1;
260 } else {
261 retval = 2;
262 }
263
264 /* Leave the chosen kernel's output in place: the tuner stands in for a real
265 evaluation and the variants do not sum in the same order */
266 if (retval == 1) {
267 CASE_1D_SEL(LX, best1);
268 } else {
269 CASE_KSTEP_SEL(LX, best);
270 }
271
272 if (retval == 1) {
273 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
274 NEKO_CHUNKS_SEL(LX, best1));
275 } else {
276 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
277 NEKO_EB_SEL(LX, best));
278 }
279 log_message(neko_log_buf);
280 log_end_section();
281 return retval;
282}
283
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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_TUNE_TIME
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
Definition elem_block_tune.h:119
neko_eb_env
static int neko_eb_env()
Definition elem_block_tune.h:79
neko_tune_rounds
static int neko_tune_rounds()
Definition elem_block_tune.h:102
NEKO_TUNE_LOG
#define NEKO_TUNE_LOG(LX, T1, T2)
Definition elem_block_tune.h:138
NEKO_TUNE_INIT
#define NEKO_TUNE_INIT
Definition elem_block_tune.h:65
NEKO_TUNE_BEST
#define NEKO_TUNE_BEST(T, BEST, N)
Definition elem_block_tune.h:130
neko_tune_iters
static int neko_tune_iters()
Definition elem_block_tune.h:110
neko_chunks_env
static int neko_chunks_env()
Definition elem_block_tune.h:91
neko_eb_sweep
static int neko_eb_sweep()
Definition elem_block_tune.h:68
NEKO_TUNE_WARMUP
#define NEKO_TUNE_WARMUP
Definition elem_block_tune.h:66
uy
__global__ void T *__restrict__ uy
Definition opgrad_kernel.h:128
uz
__global__ void T *__restrict__ T *__restrict__ uz
Definition opgrad_kernel.h:129
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
CASE_KSTEP
#define CASE_KSTEP(LX, C)
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
CASE_1D
#define CASE_1D(LX, C)