Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
opr_convect_scalar.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>
39#include "elem_block_tune.h"
41#include <device/cuda/check.h>
42
43extern "C" {
44 #include <common/neko_log.h>
45}
46
47template < const int >
48int tune_convect_scalar(void *du, void *u,
49 void *cr, void *cs, void *ct,
50 void *dx, void *dy, void *dz,
51 int *nel, int *lx, int *eb_sel, int *ch_sel);
52
53extern "C" {
54
58 void cuda_convect_scalar(void *du, void *u,
59 void *cr, void *cs, void *ct,
60 void *dx, void *dy, void *dz,
61 int *nel, int *lx) {
62
63 static int autotune[17] = { 0 };
64 /* elements per block candidate chosen by the tuner */
65 static int autotune_eb[17] = { 0 };
66 /* chunk candidate chosen for the 1d variant */
67 static int autotune_ch[17] = { 0 };
68
69 const dim3 nthrds_1d(1024, 1, 1);
70 const dim3 nthrds_kstep((*lx), (*lx), 1);
71 const dim3 nblcks((*nel), 1, 1);
72 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
73
74#define CASE_1D(LX, C) \
75 convect_scalar_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
76 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
77 ((real *) du, (real *) u, \
78 (real *) cr, (real *) cs, (real *) ct, \
79 (real *) dx, (real *) dy, (real *) dz); \
80 CUDA_CHECK(cudaGetLastError());
81
82/* Runtime dispatch onto the tuned chunk candidate */
83#define CASE_1D_SEL(LX, SEL) \
84 switch (SEL) { \
85 case 0: CASE_1D(LX, 0); break; \
86 case 1: CASE_1D(LX, 1); break; \
87 case 2: CASE_1D(LX, 2); break; \
88 default: CASE_1D(LX, 3); break; \
89 }
90
91#define CASE_KSTEP(LX, C) \
92 convect_scalar_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
93 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
94 ((real *) du, (real *) u, \
95 (real *) cr, (real *) cs, (real *) ct, \
96 (real *) dx, (real *) dy, (real *) dz, *nel); \
97 CUDA_CHECK(cudaGetLastError());
98
99/* Runtime dispatch onto the tuned candidate */
100#define CASE_KSTEP_SEL(LX, SEL) \
101 switch (SEL) { \
102 case 0: CASE_KSTEP(LX, 0); break; \
103 case 1: CASE_KSTEP(LX, 1); break; \
104 default: CASE_KSTEP(LX, 2); break; \
105 }
106
107#define CASE(LX) \
108 case LX: \
109 if(autotune[LX] == 0 ) { \
110 autotune[LX]=tune_convect_scalar<LX>(du, u, \
111 cr, cs, ct, \
112 dx, dy, dz, \
113 nel, lx, &autotune_eb[LX], \
114 &autotune_ch[LX]); \
115 } else if (autotune[LX] == 1 ) { \
116 CASE_1D_SEL(LX, autotune_ch[LX]); \
117 } else if (autotune[LX] == 2 ) { \
118 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
119 } \
120 break
121
122#define CASE_LARGE(LX) \
123 case LX: \
124 CASE_KSTEP(LX, 0); \
125 break
126
127
128 if ((*lx) < 11) {
129 switch(*lx) {
130 CASE(2);
131 CASE(3);
132 CASE(4);
133 CASE(5);
134 CASE(6);
135 CASE(7);
136 CASE(8);
137 CASE(9);
138 CASE(10);
139 default:
140 {
141 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
142 exit(1);
143 }
144 }
145 }
146 else {
147 switch(*lx) {
148 CASE_LARGE(11);
149 CASE_LARGE(12);
150 CASE_LARGE(13);
151 CASE_LARGE(14);
152 CASE_LARGE(15);
153 CASE_LARGE(16);
154 default:
155 {
156 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
157 exit(1);
158 }
159 }
160 }
161 }
162}
163
164template < const int LX >
165int tune_convect_scalar(void *du, void *u,
166 void *cr, void *cs, void *ct,
167 void *dx, void *dy, void *dz,
168 int *nel, int *lx, int *eb_sel, int *ch_sel) {
171 int best1 = 0;
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++) {
188 }
189 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
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 cudaStream_t stream = (cudaStream_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 convect_scalar (lx: %d)", *lx);
205
206 *eb_sel = 0;
207 *ch_sel = 0;
208
209 /*
210 * NEKO_AUTOTUNE names a formulation, and that is all it does: the sweep
211 * below is narrowed to that one kernel family, but its geometry -- the
212 * chunk size, the elements per block -- is still measured candidate
213 * against candidate.
214 */
215 if(env_value) {
216 if( !strcmp(env_value,"1D") ) {
217 strat = 1;
218 } else if( !strcmp(env_value,"KSTEP") ) {
219 strat = 2;
220 } else {
221 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
223 }
224 }
225
226 /* Geometry of the pinned formulation, if its own variable fixes that too.
227 Both pinned leaves nothing to measure, so the kernel is launched once and
228 reported, which is what pinning has always done */
229 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin : -1;
230
231 if (pin >= 0) {
232 switch (strat) {
233 case 1:
234 *ch_sel = pin;
236 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
238 break;
239 default:
240 *eb_sel = pin;
242 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
243 NEKO_EB_SEL(LX, pin));
244 break;
245 }
248 return strat;
249 }
250
251 if (strat) {
252 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
254 }
255
256 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
257 const bool try_1d = (strat == 0 || strat == 1);
258 const bool try_kstep = (strat == 0 || strat == 2);
259
262 /* Warm every variant before timing anything: each specialisation has to be
263 resident and the clocks at steady state, or whichever is timed first is
264 measured on a colder part */
265 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
267 CASE_1D_SEL(LX, c);
268 }
270 CASE_KSTEP_SEL(LX, c);
271 }
272 }
273
274 /* Interleaved rounds, best time per variant */
275 for (int r = 0; r < rounds; r++) {
278 }
281 }
282 }
283
287 *eb_sel = best;
288 *ch_sel = best1;
289
290 if (time1[best1] < time2[best]) {
291 retval = 1;
292 } else {
293 retval = 2;
294 }
295
296 /* Leave the chosen kernel's output in place: the tuner stands in for a real
297 evaluation and the variants do not sum in the same order */
298 if (retval == 1) {
300 } else {
302 }
303
304 if (retval == 1) {
305 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
307 } else {
308 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
310 }
313 return retval;
314}
__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)
const int i
__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
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
__global__ void const T *__restrict__ const T *__restrict__ cr
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ cs
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ ct
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
#define NEKO_EB_SEL(LX, SEL)
Definition elem_block.h:108
#define NEKO_CHUNKS_SEL(LX, SEL)
Definition elem_block.h:147
static int neko_chunks_pin()
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
static int neko_tune_rounds()
static int neko_eb_pin()
#define NEKO_TUNE_LOG(LX, T1, T2)
#define NEKO_TUNE_INIT
#define NEKO_TUNE_BEST(T, BEST, N)
static int neko_tune_iters()
#define NEKO_TUNE_FOR(C, ON, PIN, N)
static int neko_eb_sweep()
#define NEKO_TUNE_WARMUP
void log_error(char *msg)
void log_message(char *msg)
void log_end_section()
void log_section(char *msg)
#define CASE_KSTEP_SEL(LX, SEL)
#define CASE(LX)
#define CASE_1D_SEL(LX, SEL)
#define CASE_LARGE(LX)
void cuda_convect_scalar(void *du, void *u, void *cr, void *cs, void *ct, void *dx, void *dy, void *dz, int *nel, int *lx)
int tune_convect_scalar(void *du, void *u, void *cr, void *cs, void *ct, void *dx, void *dy, void *dz, int *nel, int *lx, int *eb_sel, int *ch_sel)