Neko 1.99.7
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 const int sweep = neko_eb_sweep();
176 int best = 0;
177 int retval;
178
179 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
181 }
182 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
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 cudaStream_t stream = (cudaStream_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 convect_scalar (lx: %d)", *lx);
198
199 *eb_sel = 0;
200 *ch_sel = 0;
201
202 if(env_value) {
203 if( !strcmp(env_value,"1D") ) {
206 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
210 return 1;
211 } else if( !strcmp(env_value,"KSTEP") ) {
212 *eb_sel = neko_eb_env();
214 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
218 return 2;
219 } else {
220 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
222 }
223 }
224
227 /* Warm every variant before timing anything: each specialisation has to be
228 resident and the clocks at steady state, or whichever is timed first is
229 measured on a colder part */
230 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
231 CASE_1D(LX, 0);
232 CASE_1D(LX, 1);
233 CASE_1D(LX, 2);
234 CASE_1D(LX, 3);
235 CASE_KSTEP(LX, 0);
236 if (sweep) {
237 CASE_KSTEP(LX, 1);
238 CASE_KSTEP(LX, 2);
239 }
240 }
241
242 /* Interleaved rounds, best time per variant */
243 for (int r = 0; r < rounds; r++) {
249 if (sweep) {
252 }
253 }
254
258 *eb_sel = best;
259 *ch_sel = best1;
260
261 if (time1[best1] < time2[best]) {
262 retval = 1;
263 } else {
264 retval = 2;
265 }
266
267 /* Leave the chosen kernel's output in place: the tuner stands in for a real
268 evaluation and the variants do not sum in the same order */
269 if (retval == 1) {
271 } else {
273 }
274
275 if (retval == 1) {
276 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
278 } else {
279 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
281 }
284 return retval;
285}
__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__ u
__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__ 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__ 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
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
static int neko_eb_env()
static int neko_tune_rounds()
#define NEKO_TUNE_LOG(LX, T1, T2)
#define NEKO_TUNE_INIT
#define NEKO_TUNE_BEST(T, BEST, N)
static int neko_tune_iters()
static int neko_chunks_env()
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_KSTEP(LX, C)
#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)
#define CASE_1D(LX, C)
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)