Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
opr_conv1.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 "conv1_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_conv1(void *du, void *u,
49 void *vx, void *vy, void *vz,
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 *gdim, int *lx,
55 int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel);
56
57extern "C" {
58
62 void cuda_conv1(void *du, void *u,
63 void *vx, void *vy, void *vz,
64 void *dx, void *dy, void *dz,
65 void *drdx, void *dsdx, void *dtdx,
66 void *drdy, void *dsdy, void *dtdy,
67 void *drdz, void *dsdz, void *dtdz,
68 void *jacinv, int *nel, int *gdim, int *lx) {
69
70 static int autotune[17] = { 0 };
71 /* elements per block candidate chosen by the tuner */
72 static int autotune_eb[17] = { 0 };
73 /* chunk candidate chosen for the 1d variant */
74 static int autotune_ch[17] = { 0 };
75 /* warps per block candidate chosen for the dmma variant */
76 static int autotune_nw[17] = { 0 };
77 /* warps per block candidate chosen for the tma staged dmma variant */
78 static int autotune_tw[17] = { 0 };
79
80 const dim3 nthrds_1d(1024, 1, 1);
81 const dim3 nthrds_kstep((*lx), (*lx), 1);
82 const dim3 nblcks((*nel), 1, 1);
83 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
84
85#define CASE_1D(LX, C) \
86 conv1_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
87 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
88 ((real *) du, (real *) u, \
89 (real *) vx, (real *) vy, (real *) vz, \
90 (real *) dx, (real *) dy, (real *) dz, \
91 (real *) drdx, (real *) dsdx, (real *) dtdx, \
92 (real *) drdy, (real *) dsdy, (real *) dtdy, \
93 (real *) drdz, (real *) dsdz, (real *) dtdz, \
94 (real *) jacinv); \
95 CUDA_CHECK(cudaGetLastError());
96
97/* Runtime dispatch onto the tuned chunk candidate */
98#define CASE_1D_SEL(LX, SEL) \
99 switch (SEL) { \
100 case 0: CASE_1D(LX, 0); break; \
101 case 1: CASE_1D(LX, 1); break; \
102 case 2: CASE_1D(LX, 2); break; \
103 default: CASE_1D(LX, 3); break; \
104 }
105
106#define CASE_KSTEP(LX, C) \
107 conv1_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
108 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
109 ((real *) du, (real *) u, \
110 (real *) vx, (real *) vy, (real *) vz, \
111 (real *) dx, (real *) dy, (real *) dz, \
112 (real *) drdx, (real *) dsdx, (real *) dtdx, \
113 (real *) drdy, (real *) dsdy, (real *) dtdy, \
114 (real *) drdz, (real *) dsdz, (real *) dtdz, \
115 (real *) jacinv, *nel); \
116 CUDA_CHECK(cudaGetLastError());
117
118/* Runtime dispatch onto the tuned candidate */
119#define CASE_KSTEP_SEL(LX, SEL) \
120 switch (SEL) { \
121 case 0: CASE_KSTEP(LX, 0); break; \
122 case 1: CASE_KSTEP(LX, 1); break; \
123 default: CASE_KSTEP(LX, 2); break; \
124 }
125
126#define CASE_DMMA(LX, C) \
127 conv1_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
128 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
129 ((real *) du, (real *) u, \
130 (real *) vx, (real *) vy, (real *) vz, \
131 (real *) dx, (real *) dy, (real *) dz, \
132 (real *) drdx, (real *) dsdx, (real *) dtdx, \
133 (real *) drdy, (real *) dsdy, (real *) dtdy, \
134 (real *) drdz, (real *) dsdz, (real *) dtdz, \
135 (real *) jacinv, *nel); \
136 CUDA_CHECK(cudaGetLastError());
137
138/* Runtime dispatch onto the tuned warps per block candidate */
139#define CASE_DMMA_SEL(LX, SEL) \
140 switch (SEL) { \
141 case 0: CASE_DMMA(LX, 0); break; \
142 case 1: CASE_DMMA(LX, 1); break; \
143 default: CASE_DMMA(LX, 2); break; \
144 }
145
146/*
147 * The TMA staged dmma variant. Same grid and the same warps per block
148 * candidates as CASE_DMMA -- at the one lx it supports NEKO_DMMA_PACK is 1,
149 * so the grid is nel blocks -- but unlike every other launch here it carries
150 * a dynamic shared memory allocation, 71184 B, past the 48 kB a block gets for
151 * free. The kernel has to opt into that once before its first launch; see
152 * conv1_dmma_tma_optin(), which is a predictable branch after that.
153 */
154#define CASE_DMMA_TMA(LX, C) \
155 (void) conv1_dmma_tma_optin<real, LX, NEKO_DMMA_NW(C)>(); \
156 conv1_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
157 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), \
158 NEKO_CONV1_TMA_SMEM, stream>>> \
159 ((real *) du, (real *) u, \
160 (real *) vx, (real *) vy, (real *) vz, \
161 (real *) dx, (real *) dy, (real *) dz, \
162 (real *) drdx, (real *) dsdx, (real *) dtdx, \
163 (real *) drdy, (real *) dsdy, (real *) dtdy, \
164 (real *) drdz, (real *) dsdz, (real *) dtdz, \
165 (real *) jacinv); \
166 CUDA_CHECK(cudaGetLastError());
167
168/* Runtime dispatch onto the tuned warps per block candidate */
169#define CASE_DMMA_TMA_SEL(LX, SEL) \
170 switch (SEL) { \
171 case 0: CASE_DMMA_TMA(LX, 0); break; \
172 case 1: CASE_DMMA_TMA(LX, 1); break; \
173 default: CASE_DMMA_TMA(LX, 2); break; \
174 }
175
176#define CASE(LX) \
177 case LX: \
178 if(autotune[LX] == 0 ) { \
179 autotune[LX]=tune_conv1<LX>(du, u, \
180 vx, vy, vz, \
181 dx, dy, dz, \
182 drdx, dsdx, dtdx, \
183 drdy, dsdy, dtdy, \
184 drdz, dsdz, dtdz, \
185 jacinv, nel, gdim, lx, &autotune_eb[LX], \
186 &autotune_ch[LX], &autotune_nw[LX], \
187 &autotune_tw[LX]); \
188 } else if (autotune[LX] == 1 ) { \
189 CASE_1D_SEL(LX, autotune_ch[LX]); \
190 } else if (autotune[LX] == 2 ) { \
191 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
192 } else if (autotune[LX] == 3 ) { \
193 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
194 } else if (autotune[LX] == 4 ) { \
195 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
196 } \
197 break
198
199#define CASE_LARGE(LX) \
200 case LX: \
201 CASE_KSTEP(LX, 0); \
202 break
203
204
205 if ((*lx) < 11) {
206 switch(*lx) {
207 CASE(2);
208 CASE(3);
209 CASE(4);
210 CASE(5);
211 CASE(6);
212 CASE(7);
213 CASE(8);
214 CASE(9);
215 CASE(10);
216 default:
217 {
218 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
219 exit(1);
220 }
221 }
222 }
223 else {
224 switch(*lx) {
225 CASE_LARGE(11);
226 CASE_LARGE(12);
227 CASE_LARGE(13);
228 CASE_LARGE(14);
229 CASE_LARGE(15);
230 CASE_LARGE(16);
231 default:
232 {
233 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
234 exit(1);
235 }
236 }
237 }
238 }
239}
240
241template < const int LX >
242int tune_conv1(void *du, void *u,
243 void *vx, void *vy, void *vz,
244 void *dx, void *dy, void *dz,
245 void *drdx, void *dsdx, void *dtdx,
246 void *drdy, void *dsdy, void *dtdy,
247 void *drdz, void *dsdz, void *dtdz,
248 void *jacinv, int *nel, int *gdim, int *lx,
249 int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel) {
252 int best1 = 0;
255 int best3 = 0;
257 int best4 = 0;
258 const int rounds = neko_tune_rounds();
259 const int iters = neko_tune_iters();
260 const int sweep = neko_eb_sweep();
261 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
262 /* Whether the pointers are bulk copy aligned is a property of this call
263 rather than of the kernel, so it is checked rather than assumed, see
264 dmma_tma_kernel.h. The shared memory gate is a device query too: this
265 variant asks for more than a block gets by default */
266 const bool tma = dmma_tma_conv1_lx_supported<LX>() &&
269 drdx, dsdx, dtdx,
270 drdy, dsdy, dtdy,
271 drdz, dsdz, dtdz);
272 int best = 0;
273 int retval;
274
275 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
277 }
278 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
280 }
281 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
284 }
285
286 const dim3 nthrds_1d(1024, 1, 1);
287 const dim3 nthrds_kstep((*lx), (*lx), 1);
288 const dim3 nblcks((*nel), 1, 1);
289 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
290
291 char *env_value = NULL;
292 char neko_log_buf[80];
293
294 env_value=getenv("NEKO_AUTOTUNE");
295
296 sprintf(neko_log_buf, "Autotune conv1 (lx: %d)", *lx);
298
299 *eb_sel = 0;
300 *ch_sel = 0;
301 *nw_sel = 0;
302 *tw_sel = 0;
303
304 if(env_value) {
305 if( !strcmp(env_value,"1D") ) {
308 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
312 return 1;
313 } else if( !strcmp(env_value,"KSTEP") ) {
314 *eb_sel = neko_eb_env();
316 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
320 return 2;
321 } else if( !strcmp(env_value,"DMMA") ) {
322 if (dmma) {
323 const int c = neko_dmma_env();
324 *nw_sel = c;
325 CASE_DMMA_SEL(LX, c);
326 sprintf(neko_log_buf,"Set by env : 3 (DMMA, %d warps)",
327 NEKO_DMMA_NW(c));
330 return 3;
331 } else {
332 sprintf(neko_log_buf, "DMMA strategy not available for this config");
334 }
335 } else if( !strcmp(env_value,"DMMA_TMA") ) {
336 if (tma) {
337 const int c = neko_dmma_tma_env();
338 *tw_sel = c;
340 sprintf(neko_log_buf,"Set by env : 4 (DMMA_TMA, %d warps)",
341 NEKO_DMMA_NW(c));
344 return 4;
345 } else {
347 "DMMA_TMA strategy not available for this config");
349 }
350 } else {
351 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
353 }
354 }
355
358
359 /* Warm every variant before timing anything: each specialisation has to be
360 resident and the clocks at steady state, or whichever is timed first is
361 measured on a colder part */
362 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
363 CASE_1D(LX, 0);
364 CASE_1D(LX, 1);
365 CASE_1D(LX, 2);
366 CASE_1D(LX, 3);
367 CASE_KSTEP(LX, 0);
368 if (sweep) {
369 CASE_KSTEP(LX, 1);
370 CASE_KSTEP(LX, 2);
371 }
372 if (dmma) {
373 CASE_DMMA(LX, 0);
374 CASE_DMMA(LX, 1);
375 CASE_DMMA(LX, 2);
376 }
377 if (tma) {
378 CASE_DMMA_TMA(LX, 0);
379 CASE_DMMA_TMA(LX, 1);
380 CASE_DMMA_TMA(LX, 2);
381 }
382 }
383
384 /* Interleaved rounds, best time per variant */
385 for (int r = 0; r < rounds; r++) {
391 if (sweep) {
394 }
395 if (dmma) {
399 }
400 if (tma) {
404 }
405 }
406
410
415 *eb_sel = best;
416 *ch_sel = best1;
417 *nw_sel = best3;
418 *tw_sel = best4;
419
420 if (time1[best1] < time2[best]) {
421 retval = 1;
422 } else {
423 retval = 2;
424 }
425
426 /* The dmma variants join the comparison only where they exist, their
427 candidates are left at NEKO_TUNE_INIT otherwise */
428 float tbest = (retval == 1) ? time1[best1] : time2[best];
429
430 if (time3[best3] < tbest) {
431 retval = 3;
432 tbest = time3[best3];
433 }
434 if (time4[best4] < tbest) {
435 retval = 4;
436 }
437
438 /* Leave the chosen kernel's output in place: the tuner stands in for a real
439 evaluation and the variants do not sum in the same order */
440 if (retval == 1) {
442 } else if (retval == 2) {
444 } else if (retval == 3) {
446 } else {
448 }
449
450 if (retval == 1) {
451 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
453 } else if (retval == 2) {
454 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
456 } else if (retval == 3) {
457 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
459 } else {
460 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
462 }
465 return retval;
466}
__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)
__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
__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
__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
__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
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__ 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
__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 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
__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
__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
__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
__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
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ vz
__global__ void const T *__restrict__ const T *__restrict__ vx
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ vy
#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
#define NEKO_DMMA_CANDIDATES
static bool cuda_have_dmma()
static int neko_dmma_env()
#define NEKO_DMMA_PACK(LX)
#define NEKO_TUNE_LOG_DMMA(LX, T3)
#define NEKO_DMMA_NW(C)
static bool cuda_have_tma_conv1()
#define NEKO_TUNE_LOG_DMMA_TMA(LX, T4)
static int neko_dmma_tma_env()
static bool dmma_tma_conv1_aligned(const void *du, const void *u, const void *vx, const void *vy, const void *vz, const void *jacinv, const void *drdx, const void *dsdx, const void *dtdx, const void *drdy, const void *dsdy, const void *dtdy, const void *drdz, const void *dsdz, const void *dtdz)
void log_error(char *msg)
void log_message(char *msg)
void log_end_section()
void log_section(char *msg)
#define CASE_DMMA(LX, C)
#define CASE_DMMA_SEL(LX, SEL)
#define CASE_KSTEP_SEL(LX, SEL)
#define CASE(LX)
void cuda_conv1(void *du, void *u, void *vx, void *vy, void *vz, 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 *gdim, int *lx)
Definition opr_conv1.cu:62
#define CASE_1D_SEL(LX, SEL)
#define CASE_KSTEP(LX, C)
#define CASE_LARGE(LX)
#define CASE_DMMA_TMA_SEL(LX, SEL)
#define CASE_DMMA_TMA(LX, C)
#define CASE_1D(LX, C)
int tune_conv1(void *du, void *u, void *vx, void *vy, void *vz, 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 *gdim, int *lx, int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel)
Definition opr_conv1.cu:242