Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
opr_dudxyz.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 "dudxyz_kernel.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_dudxyz(void *du, void *u,
49 void *dr, void *ds, void *dt,
50 void *dx, void *dy, void *dz,
51 void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel);
52
53extern "C" {
54
58 void cuda_dudxyz(void *du, void *u,
59 void *dr, void *ds, void *dt,
60 void *dx, void *dy, void *dz,
61 void *jacinv, int *nel, int *lx) {
62
63 static int autotune[16] = { 0 };
64 /* elements per block candidate chosen by the tuner */
65 static int autotune_eb[16] = { 0 };
66 /* chunk candidate chosen for the 1d variant */
67 static int autotune_ch[16] = { 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 dudxyz_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
76 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
77 ((real *) du, (real *) u, \
78 (real *) dr, (real *) ds, (real *) dt, \
79 (real *) dx, (real *) dy, (real *) dz, \
80 (real *) jacinv); \
81 CUDA_CHECK(cudaGetLastError());
82
83/* Runtime dispatch onto the tuned chunk candidate */
84#define CASE_1D_SEL(LX, SEL) \
85 switch (SEL) { \
86 case 0: CASE_1D(LX, 0); break; \
87 case 1: CASE_1D(LX, 1); break; \
88 case 2: CASE_1D(LX, 2); break; \
89 default: CASE_1D(LX, 3); break; \
90 }
91
92#define CASE_KSTEP(LX, C) \
93 dudxyz_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
94 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
95 ((real *) du, (real *) u, \
96 (real *) dr, (real *) ds, (real *) dt, \
97 (real *) dx, (real *) dy, (real *) dz, \
98 (real *) jacinv, *nel); \
99 CUDA_CHECK(cudaGetLastError());
100
101/* Runtime dispatch onto the tuned candidate */
102#define CASE_KSTEP_SEL(LX, SEL) \
103 switch (SEL) { \
104 case 0: CASE_KSTEP(LX, 0); break; \
105 case 1: CASE_KSTEP(LX, 1); break; \
106 default: CASE_KSTEP(LX, 2); break; \
107 }
108
109 #define CASE(LX) \
110 case LX: \
111 if(autotune[LX] == 0 ) { \
112 autotune[LX]=tune_dudxyz<LX>(du, u, \
113 dr, ds, dt, \
114 dx, dy, dz, \
115 jacinv, nel, lx, &autotune_eb[LX], \
116 &autotune_ch[LX]); \
117 } else if (autotune[LX] == 1 ) { \
118 CASE_1D_SEL(LX, autotune_ch[LX]); \
119 } else if (autotune[LX] == 2 ) { \
120 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
121 } \
122 break
123
124#define CASE_LARGE(LX) \
125 case LX: \
126 CASE_KSTEP(LX, 0); \
127 break
128
129
130 if ((*lx) < 11) {
131 switch(*lx) {
132 CASE(2);
133 CASE(3);
134 CASE(4);
135 CASE(5);
136 CASE(6);
137 CASE(7);
138 CASE(8);
139 CASE(9);
140 CASE(10);
141 default:
142 {
143 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
144 exit(1);
145 }
146 }
147 }
148 else {
149 switch(*lx) {
150 CASE_LARGE(11);
151 CASE_LARGE(12);
152 CASE_LARGE(13);
153 CASE_LARGE(14);
154 CASE_LARGE(15);
155 CASE_LARGE(16);
156 default:
157 {
158 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
159 exit(1);
160 }
161 }
162 }
163 }
164}
165
166template < const int LX >
167int tune_dudxyz(void *du, void *u,
168 void *dr, void *ds, void *dt,
169 void *dx, void *dy, void *dz,
170 void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel) {
173 int best1 = 0;
175 const int rounds = neko_tune_rounds();
176 const int iters = neko_tune_iters();
177 const int sweep = neko_eb_sweep();
178 int best = 0;
179 int retval;
180
181 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
183 }
184 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
186 }
187
188 const dim3 nthrds_1d(1024, 1, 1);
189 const dim3 nthrds_kstep((*lx), (*lx), 1);
190 const dim3 nblcks((*nel), 1, 1);
191 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
192
193 char *env_value = NULL;
194 char neko_log_buf[80];
195
196 env_value=getenv("NEKO_AUTOTUNE");
197
198 sprintf(neko_log_buf, "Autotune dudxyz (lx: %d)", *lx);
200
201 *eb_sel = 0;
202 *ch_sel = 0;
203
204 if(env_value) {
205 if( !strcmp(env_value,"1D") ) {
208 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
212 return 1;
213 } else if( !strcmp(env_value,"KSTEP") ) {
214 *eb_sel = neko_eb_env();
216 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
220 return 2;
221 } else {
222 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
224 }
225 }
226
229
230 /* Warm every variant before timing anything: each specialisation has to be
231 resident and the clocks at steady state, or whichever is timed first is
232 measured on a colder part */
233 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
234 CASE_1D(LX, 0);
235 CASE_1D(LX, 1);
236 CASE_1D(LX, 2);
237 CASE_1D(LX, 3);
238 CASE_KSTEP(LX, 0);
239 if (sweep) {
240 CASE_KSTEP(LX, 1);
241 CASE_KSTEP(LX, 2);
242 }
243 }
244
245 /* Interleaved rounds, best time per variant */
246 for (int r = 0; r < rounds; r++) {
252 if (sweep) {
255 }
256 }
257
261 *eb_sel = best;
262 *ch_sel = best1;
263
264 if (time1[best1] < time2[best]) {
265 retval = 1;
266 } else {
267 retval = 2;
268 }
269
270 /* Leave the chosen kernel's output in place: the tuner stands in for a real
271 evaluation and the variants do not sum in the same order */
272 if (retval == 1) {
274 } else {
276 }
277
278 if (retval == 1) {
279 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
281 } else {
282 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
284 }
287 return retval;
288}
__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 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
__global__ void const T *__restrict__ const T *__restrict__ dr
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ ds
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
#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)
void cuda_dudxyz(void *du, void *u, void *dr, void *ds, void *dt, void *dx, void *dy, void *dz, void *jacinv, int *nel, int *lx)
Definition opr_dudxyz.cu:58
#define CASE_KSTEP_SEL(LX, SEL)
#define CASE(LX)
#define CASE_1D_SEL(LX, SEL)
int tune_dudxyz(void *du, void *u, void *dr, void *ds, void *dt, void *dx, void *dy, void *dz, void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel)
#define CASE_KSTEP(LX, C)
#define CASE_LARGE(LX)
#define CASE_1D(LX, C)