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 /* Candidates of the kstep sweep, one -- the unblocked shape -- when the
261 elements per block sweep is off */
262 const int eb_cand = neko_eb_sweep() ? NEKO_EB_CANDIDATES : 1;
263 /* Geometry pinned by each formulation's own variable, -1 to sweep it */
264 const int ch_pin = neko_chunks_pin();
265 const int eb_pin = neko_eb_pin();
266 const int nw_pin = neko_dmma_pin();
267 const int tw_pin = neko_dmma_tma_pin();
268 /* Formulation pinned by NEKO_AUTOTUNE, as the identifier this returns */
269 int strat = 0;
270 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
271 /* Whether the pointers are bulk copy aligned is a property of this call
272 rather than of the kernel, so it is checked rather than assumed, see
273 dmma_tma_kernel.h. The shared memory gate is a device query too: this
274 variant asks for more than a block gets by default */
275 const bool tma = dmma_tma_conv1_lx_supported<LX>() &&
278 drdx, dsdx, dtdx,
279 drdy, dsdy, dtdy,
280 drdz, dsdz, dtdz);
281 int best = 0;
282 int retval;
283
284 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
286 }
287 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
289 }
290 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
293 }
294
295 const dim3 nthrds_1d(1024, 1, 1);
296 const dim3 nthrds_kstep((*lx), (*lx), 1);
297 const dim3 nblcks((*nel), 1, 1);
298 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
299
300 char *env_value = NULL;
301 char neko_log_buf[80];
302
303 env_value=getenv("NEKO_AUTOTUNE");
304
305 sprintf(neko_log_buf, "Autotune conv1 (lx: %d)", *lx);
307
308 *eb_sel = 0;
309 *ch_sel = 0;
310 *nw_sel = 0;
311 *tw_sel = 0;
312
313 /*
314 * NEKO_AUTOTUNE names a formulation, and that is all it does: the sweep
315 * below is narrowed to that one kernel family, but its geometry -- the
316 * chunk size, the elements per block, the warps per block -- is still
317 * measured candidate against candidate. A formulation this build or this
318 * device does not have is reported and ignored, leaving the full sweep.
319 */
320 if(env_value) {
321 if( !strcmp(env_value,"1D") ) {
322 strat = 1;
323 } else if( !strcmp(env_value,"KSTEP") ) {
324 strat = 2;
325 } else if( !strcmp(env_value,"DMMA") ) {
326 if (dmma) {
327 strat = 3;
328 } else {
329 sprintf(neko_log_buf, "DMMA strategy not available for this config");
331 }
332 } else if( !strcmp(env_value,"DMMA_TMA") ) {
333 if (tma) {
334 strat = 4;
335 } else {
337 "DMMA_TMA strategy not available for this config");
339 }
340 } else {
341 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
343 }
344 }
345
346 /* Geometry of the pinned formulation, if its own variable fixes that too.
347 Both pinned leaves nothing to measure, so the kernel is launched once and
348 reported, which is what pinning has always done */
349 const int pin = (strat == 1) ? ch_pin : (strat == 2) ? eb_pin :
350 (strat == 3) ? nw_pin : (strat == 4) ? tw_pin : -1;
351
352 if (pin >= 0) {
353 switch (strat) {
354 case 1:
355 *ch_sel = pin;
357 sprintf(neko_log_buf, "Set by env : 1 (1D, %d chunk)",
359 break;
360 case 2:
361 *eb_sel = pin;
363 sprintf(neko_log_buf, "Set by env : 2 (KSTEP, %d elem/block)",
364 NEKO_EB_SEL(LX, pin));
365 break;
366 case 3:
367 *nw_sel = pin;
369 sprintf(neko_log_buf, "Set by env : 3 (DMMA, %d warps)",
371 break;
372 default:
373 *tw_sel = pin;
375 sprintf(neko_log_buf, "Set by env : 4 (DMMA_TMA, %d warps)",
377 break;
378 }
381 return strat;
382 }
383
384 if (strat) {
385 sprintf(neko_log_buf, "Set by env : %d (%s)", strat, env_value);
387 }
388
389 /* Formulations the sweep considers, see NEKO_TUNE_FOR() */
390 const bool try_1d = (strat == 0 || strat == 1);
391 const bool try_kstep = (strat == 0 || strat == 2);
392 const bool try_dmma = dmma && (strat == 0 || strat == 3);
393 const bool try_tma = tma && (strat == 0 || strat == 4);
394
397
398 /* Warm every variant before timing anything: each specialisation has to be
399 resident and the clocks at steady state, or whichever is timed first is
400 measured on a colder part */
401 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
403 CASE_1D_SEL(LX, c);
404 }
406 CASE_KSTEP_SEL(LX, c);
407 }
409 CASE_DMMA_SEL(LX, c);
410 }
413 }
414 }
415
416 /* Interleaved rounds, best time per variant */
417 for (int r = 0; r < rounds; r++) {
420 }
423 }
426 }
429 }
430 }
431
435
440 *eb_sel = best;
441 *ch_sel = best1;
442 *nw_sel = best3;
443 *tw_sel = best4;
444
445 if (time1[best1] < time2[best]) {
446 retval = 1;
447 } else {
448 retval = 2;
449 }
450
451 /* The dmma variants join the comparison only where they exist, their
452 candidates are left at NEKO_TUNE_INIT otherwise */
453 float tbest = (retval == 1) ? time1[best1] : time2[best];
454
455 if (time3[best3] < tbest) {
456 retval = 3;
457 tbest = time3[best3];
458 }
459 if (time4[best4] < tbest) {
460 retval = 4;
461 }
462
463 /* Leave the chosen kernel's output in place: the tuner stands in for a real
464 evaluation and the variants do not sum in the same order */
465 if (retval == 1) {
467 } else if (retval == 2) {
469 } else if (retval == 3) {
471 } else {
473 }
474
475 if (retval == 1) {
476 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
478 } else if (retval == 2) {
479 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
481 } else if (retval == 3) {
482 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
484 } else {
485 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
487 }
490 return retval;
491}
__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
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
#define NEKO_DMMA_CANDIDATES
static bool cuda_have_dmma()
#define NEKO_DMMA_PACK(LX)
#define NEKO_TUNE_LOG_DMMA(LX, T3)
#define NEKO_DMMA_NW(C)
static int neko_dmma_pin()
static bool cuda_have_tma_conv1()
#define NEKO_TUNE_LOG_DMMA_TMA(LX, T4)
static int neko_dmma_tma_pin()
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_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_LARGE(LX)
#define CASE_DMMA_TMA_SEL(LX, SEL)
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