Neko 1.99.9
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"
40#include <device/device_config.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 int *nw_sel, int *tw_sel);
53
54extern "C" {
55
59 void cuda_dudxyz(void *du, void *u,
60 void *dr, void *ds, void *dt,
61 void *dx, void *dy, void *dz,
62 void *jacinv, int *nel, int *lx) {
63
64 static int autotune[16] = { 0 };
65 /* elements per block candidate chosen by the tuner */
66 static int autotune_eb[16] = { 0 };
67 /* chunk candidate chosen for the 1d variant */
68 static int autotune_ch[16] = { 0 };
69 /* warps per block candidate chosen for the dmma variant */
70 static int autotune_nw[16] = { 0 };
71 /* warps per block candidate chosen for the tma staged dmma variant */
72 static int autotune_tw[16] = { 0 };
73
74 const dim3 nthrds_1d(1024, 1, 1);
75 const dim3 nthrds_kstep((*lx), (*lx), 1);
76 const dim3 nblcks((*nel), 1, 1);
77 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
78
79#define CASE_1D(LX, C) \
80 dudxyz_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
81 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
82 ((real *) du, (real *) u, \
83 (real *) dr, (real *) ds, (real *) dt, \
84 (real *) dx, (real *) dy, (real *) dz, \
85 (real *) jacinv); \
86 CUDA_CHECK(cudaGetLastError());
87
88/* Runtime dispatch onto the tuned chunk candidate */
89#define CASE_1D_SEL(LX, SEL) \
90 switch (SEL) { \
91 case 0: CASE_1D(LX, 0); break; \
92 case 1: CASE_1D(LX, 1); break; \
93 case 2: CASE_1D(LX, 2); break; \
94 default: CASE_1D(LX, 3); break; \
95 }
96
97#define CASE_KSTEP(LX, C) \
98 dudxyz_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
99 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
100 ((real *) du, (real *) u, \
101 (real *) dr, (real *) ds, (real *) dt, \
102 (real *) dx, (real *) dy, (real *) dz, \
103 (real *) jacinv, *nel); \
104 CUDA_CHECK(cudaGetLastError());
105
106/* Runtime dispatch onto the tuned candidate */
107#define CASE_KSTEP_SEL(LX, SEL) \
108 switch (SEL) { \
109 case 0: CASE_KSTEP(LX, 0); break; \
110 case 1: CASE_KSTEP(LX, 1); break; \
111 default: CASE_KSTEP(LX, 2); break; \
112 }
113
114#define CASE_DMMA(LX, C) \
115 dudxyz_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
116 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
117 ((real *) du, (real *) u, \
118 (real *) dr, (real *) ds, (real *) dt, \
119 (real *) dx, (real *) dy, (real *) dz, \
120 (real *) jacinv, *nel); \
121 CUDA_CHECK(cudaGetLastError());
122
123/* Runtime dispatch onto the tuned warps per block candidate */
124#define CASE_DMMA_SEL(LX, SEL) \
125 switch (SEL) { \
126 case 0: CASE_DMMA(LX, 0); break; \
127 case 1: CASE_DMMA(LX, 1); break; \
128 default: CASE_DMMA(LX, 2); break; \
129 }
130
131/*
132 * The TMA staged dmma variant. Same grid and the same warps per block
133 * candidates as CASE_DMMA -- at the one lx it supports NEKO_DMMA_PACK is 1,
134 * so the grid is nel blocks -- but it takes no nel: there is no packed tail
135 * to clamp when a block is exactly an element. See dmma_tma_kernel.h.
136 */
137#define CASE_DMMA_TMA(LX, C) \
138 dudxyz_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
139 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
140 ((real *) du, (real *) u, \
141 (real *) dr, (real *) ds, (real *) dt, \
142 (real *) dx, (real *) dy, (real *) dz, \
143 (real *) jacinv); \
144 CUDA_CHECK(cudaGetLastError());
145
146/* Runtime dispatch onto the tuned warps per block candidate */
147#define CASE_DMMA_TMA_SEL(LX, SEL) \
148 switch (SEL) { \
149 case 0: CASE_DMMA_TMA(LX, 0); break; \
150 case 1: CASE_DMMA_TMA(LX, 1); break; \
151 default: CASE_DMMA_TMA(LX, 2); break; \
152 }
153
154 #define CASE(LX) \
155 case LX: \
156 if(autotune[LX] == 0 ) { \
157 autotune[LX]=tune_dudxyz<LX>(du, u, \
158 dr, ds, dt, \
159 dx, dy, dz, \
160 jacinv, nel, lx, &autotune_eb[LX], \
161 &autotune_ch[LX], &autotune_nw[LX], \
162 &autotune_tw[LX]); \
163 } else if (autotune[LX] == 1 ) { \
164 CASE_1D_SEL(LX, autotune_ch[LX]); \
165 } else if (autotune[LX] == 2 ) { \
166 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
167 } else if (autotune[LX] == 3 ) { \
168 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
169 } else if (autotune[LX] == 4 ) { \
170 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
171 } \
172 break
173
174#define CASE_LARGE(LX) \
175 case LX: \
176 CASE_KSTEP(LX, 0); \
177 break
178
179
180 if ((*lx) < 11) {
181 switch(*lx) {
182 CASE(2);
183 CASE(3);
184 CASE(4);
185 CASE(5);
186 CASE(6);
187 CASE(7);
188 CASE(8);
189 CASE(9);
190 CASE(10);
191 default:
192 {
193 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
194 exit(1);
195 }
196 }
197 }
198 else {
199 switch(*lx) {
200 CASE_LARGE(11);
201 CASE_LARGE(12);
202 CASE_LARGE(13);
203 CASE_LARGE(14);
204 CASE_LARGE(15);
205 CASE_LARGE(16);
206 default:
207 {
208 fprintf(stderr, __FILE__ ": size not supported: %d\n", *lx);
209 exit(1);
210 }
211 }
212 }
213 }
214}
215
216template < const int LX >
217int tune_dudxyz(void *du, void *u,
218 void *dr, void *ds, void *dt,
219 void *dx, void *dy, void *dz,
220 void *jacinv, int *nel, int *lx, int *eb_sel, int *ch_sel,
221 int *nw_sel, int *tw_sel) {
222 cudaEvent_t start,stop;
223 float time1[NEKO_CHUNKS_CANDIDATES];
224 int best1 = 0;
225 float time2[NEKO_EB_CANDIDATES];
226 float time3[NEKO_DMMA_CANDIDATES];
227 int best3 = 0;
228 float time4[NEKO_DMMA_CANDIDATES];
229 int best4 = 0;
230 const int rounds = neko_tune_rounds();
231 const int iters = neko_tune_iters();
232 const int sweep = neko_eb_sweep();
233 const bool dmma = dmma_lx_supported<LX>() && cuda_have_dmma();
234 /* Whether the pointers are bulk copy aligned is a property of this call
235 rather than of the kernel, so it is checked rather than assumed, see
236 dmma_tma_dudxyz_aligned() in dmma_tma_kernel.h */
237 const bool tma = dmma_tma_dudxyz_lx_supported<LX>() && cuda_have_tma() &&
238 dmma_tma_dudxyz_aligned(du, u, dr, ds, dt, jacinv);
239 int best = 0;
240 int retval;
241
242 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) {
243 time2[c] = NEKO_TUNE_INIT;
244 }
245 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) {
246 time1[c] = NEKO_TUNE_INIT;
247 }
248 for (int c = 0; c < NEKO_DMMA_CANDIDATES; c++) {
249 time3[c] = NEKO_TUNE_INIT;
250 time4[c] = NEKO_TUNE_INIT;
251 }
252
253 const dim3 nthrds_1d(1024, 1, 1);
254 const dim3 nthrds_kstep((*lx), (*lx), 1);
255 const dim3 nblcks((*nel), 1, 1);
256 const cudaStream_t stream = (cudaStream_t) glb_cmd_queue;
257
258 char *env_value = NULL;
259 char neko_log_buf[80];
260
261 env_value=getenv("NEKO_AUTOTUNE");
262
263 sprintf(neko_log_buf, "Autotune dudxyz (lx: %d)", *lx);
264 log_section(neko_log_buf);
265
266 *eb_sel = 0;
267 *ch_sel = 0;
268 *nw_sel = 0;
269 *tw_sel = 0;
270
271 if(env_value) {
272 if( !strcmp(env_value,"1D") ) {
273 *ch_sel = neko_chunks_env();
274 CASE_1D_SEL(LX, *ch_sel);
275 sprintf(neko_log_buf,"Set by env : 1 (1D, %d chunk)",
276 NEKO_CHUNKS_SEL(LX, *ch_sel));
277 log_message(neko_log_buf);
278 log_end_section();
279 return 1;
280 } else if( !strcmp(env_value,"KSTEP") ) {
281 *eb_sel = neko_eb_env();
282 CASE_KSTEP_SEL(LX, *eb_sel);
283 sprintf(neko_log_buf,"Set by env : 2 (KSTEP, %d elem/block)",
284 NEKO_EB_SEL(LX, *eb_sel));
285 log_message(neko_log_buf);
286 log_end_section();
287 return 2;
288 } else if( !strcmp(env_value,"DMMA") ) {
289 if (dmma) {
290 const int c = neko_dmma_env();
291 *nw_sel = c;
292 CASE_DMMA_SEL(LX, c);
293 sprintf(neko_log_buf,"Set by env : 3 (DMMA, %d warps)",
294 NEKO_DMMA_NW(c));
295 log_message(neko_log_buf);
296 log_end_section();
297 return 3;
298 } else {
299 sprintf(neko_log_buf, "DMMA strategy not available for this config");
300 log_error(neko_log_buf);
301 }
302 } else if( !strcmp(env_value,"DMMA_TMA") ) {
303 if (tma) {
304 const int c = neko_dmma_tma_env();
305 *tw_sel = c;
306 CASE_DMMA_TMA_SEL(LX, c);
307 sprintf(neko_log_buf,"Set by env : 4 (DMMA_TMA, %d warps)",
308 NEKO_DMMA_NW(c));
309 log_message(neko_log_buf);
310 log_end_section();
311 return 4;
312 } else {
313 sprintf(neko_log_buf,
314 "DMMA_TMA strategy not available for this config");
315 log_error(neko_log_buf);
316 }
317 } else {
318 sprintf(neko_log_buf, "Invalid value set for NEKO_AUTOTUNE");
319 log_error(neko_log_buf);
320 }
321 }
322
323 cudaEventCreate(&start);
324 cudaEventCreate(&stop);
325
326 /* Warm every variant before timing anything: each specialisation has to be
327 resident and the clocks at steady state, or whichever is timed first is
328 measured on a colder part */
329 for (int i = 0; i < NEKO_TUNE_WARMUP; i++) {
330 CASE_1D(LX, 0);
331 CASE_1D(LX, 1);
332 CASE_1D(LX, 2);
333 CASE_1D(LX, 3);
334 CASE_KSTEP(LX, 0);
335 if (sweep) {
336 CASE_KSTEP(LX, 1);
337 CASE_KSTEP(LX, 2);
338 }
339 if (dmma) {
340 CASE_DMMA(LX, 0);
341 CASE_DMMA(LX, 1);
342 CASE_DMMA(LX, 2);
343 }
344 if (tma) {
345 CASE_DMMA_TMA(LX, 0);
346 CASE_DMMA_TMA(LX, 1);
347 CASE_DMMA_TMA(LX, 2);
348 }
349 }
350
351 /* Interleaved rounds, best time per variant */
352 for (int r = 0; r < rounds; r++) {
353 NEKO_TUNE_TIME(time1, CASE_1D, LX, 0, iters);
354 NEKO_TUNE_TIME(time1, CASE_1D, LX, 1, iters);
355 NEKO_TUNE_TIME(time1, CASE_1D, LX, 2, iters);
356 NEKO_TUNE_TIME(time1, CASE_1D, LX, 3, iters);
357 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 0, iters);
358 if (sweep) {
359 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 1, iters);
360 NEKO_TUNE_TIME(time2, CASE_KSTEP, LX, 2, iters);
361 }
362 if (dmma) {
363 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 0, iters);
364 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 1, iters);
365 NEKO_TUNE_TIME(time3, CASE_DMMA, LX, 2, iters);
366 }
367 if (tma) {
368 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 0, iters);
369 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 1, iters);
370 NEKO_TUNE_TIME(time4, CASE_DMMA_TMA, LX, 2, iters);
371 }
372 }
373
374 NEKO_TUNE_LOG(LX, time1, time2);
375 NEKO_TUNE_LOG_DMMA(LX, time3);
376 NEKO_TUNE_LOG_DMMA_TMA(LX, time4);
377
378 NEKO_TUNE_BEST(time1, best1, NEKO_CHUNKS_CANDIDATES);
379 NEKO_TUNE_BEST(time2, best, NEKO_EB_CANDIDATES);
380 NEKO_TUNE_BEST(time3, best3, NEKO_DMMA_CANDIDATES);
381 NEKO_TUNE_BEST(time4, best4, NEKO_DMMA_CANDIDATES);
382 *eb_sel = best;
383 *ch_sel = best1;
384 *nw_sel = best3;
385 *tw_sel = best4;
386
387 if (time1[best1] < time2[best]) {
388 retval = 1;
389 } else {
390 retval = 2;
391 }
392
393 /* The dmma variants join the comparison only where they exist, their
394 candidates are left at NEKO_TUNE_INIT otherwise */
395 float tbest = (retval == 1) ? time1[best1] : time2[best];
396
397 if (time3[best3] < tbest) {
398 retval = 3;
399 tbest = time3[best3];
400 }
401 if (time4[best4] < tbest) {
402 retval = 4;
403 }
404
405 /* Leave the chosen kernel's output in place: the tuner stands in for a real
406 evaluation and the variants do not sum in the same order */
407 if (retval == 1) {
408 CASE_1D_SEL(LX, best1);
409 } else if (retval == 2) {
410 CASE_KSTEP_SEL(LX, best);
411 } else if (retval == 3) {
412 CASE_DMMA_SEL(LX, best3);
413 } else {
414 CASE_DMMA_TMA_SEL(LX, best4);
415 }
416
417 if (retval == 1) {
418 sprintf(neko_log_buf, "Chose : 1 (1D, %d chunk)",
419 NEKO_CHUNKS_SEL(LX, best1));
420 } else if (retval == 2) {
421 sprintf(neko_log_buf, "Chose : 2 (KSTEP, %d elem/block)",
422 NEKO_EB_SEL(LX, best));
423 } else if (retval == 3) {
424 sprintf(neko_log_buf, "Chose : 3 (DMMA, %d warps, %d elem/blk)",
425 NEKO_DMMA_NW(best3), NEKO_DMMA_PACK(LX));
426 } else {
427 sprintf(neko_log_buf, "Chose : 4 (DMMA_TMA, %d warps)",
428 NEKO_DMMA_NW(best4));
429 }
430 log_message(neko_log_buf);
431 log_end_section();
432 return retval;
433}
ale_add_kinematics_kernel
__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)
Definition ale_kinematics_kernel.h:52
jacinv
__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
Definition ax_helm_full_kernel.h:62
i
const int i
Definition ax_helm_full_kernel.h:95
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__ dz
Definition ax_helm_full_kernel.h:51
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition ax_helm_full_kernel.h:49
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition ax_helm_full_kernel.h:50
ds
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ ds
Definition cdtp_kernel.h:112
dr
__global__ void const T *__restrict__ const T *__restrict__ dr
Definition cdtp_kernel.h:111
dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
Definition cdtp_kernel.h:113
NEKO_CHUNKS_CANDIDATES
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
NEKO_EB_CANDIDATES
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
NEKO_EB_SEL
#define NEKO_EB_SEL(LX, SEL)
Definition elem_block.h:108
NEKO_CHUNKS_SEL
#define NEKO_CHUNKS_SEL(LX, SEL)
Definition elem_block.h:147
NEKO_TUNE_TIME
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
Definition elem_block_tune.h:119
neko_eb_env
static int neko_eb_env()
Definition elem_block_tune.h:79
neko_tune_rounds
static int neko_tune_rounds()
Definition elem_block_tune.h:102
NEKO_TUNE_LOG
#define NEKO_TUNE_LOG(LX, T1, T2)
Definition elem_block_tune.h:149
NEKO_TUNE_INIT
#define NEKO_TUNE_INIT
Definition elem_block_tune.h:65
NEKO_TUNE_BEST
#define NEKO_TUNE_BEST(T, BEST, N)
Definition elem_block_tune.h:141
neko_tune_iters
static int neko_tune_iters()
Definition elem_block_tune.h:110
neko_chunks_env
static int neko_chunks_env()
Definition elem_block_tune.h:91
neko_eb_sweep
static int neko_eb_sweep()
Definition elem_block_tune.h:68
NEKO_TUNE_WARMUP
#define NEKO_TUNE_WARMUP
Definition elem_block_tune.h:66
device_config.h
NEKO_DMMA_CANDIDATES
#define NEKO_DMMA_CANDIDATES
Definition dmma_kernel.h:123
cuda_have_dmma
static bool cuda_have_dmma()
Definition dmma_kernel.h:224
neko_dmma_env
static int neko_dmma_env()
Definition dmma_kernel.h:154
NEKO_DMMA_PACK
#define NEKO_DMMA_PACK(LX)
Definition dmma_kernel.h:148
NEKO_TUNE_LOG_DMMA
#define NEKO_TUNE_LOG_DMMA(LX, T3)
Definition dmma_kernel.h:169
NEKO_DMMA_NW
#define NEKO_DMMA_NW(C)
Definition dmma_kernel.h:124
NEKO_TUNE_LOG_DMMA_TMA
#define NEKO_TUNE_LOG_DMMA_TMA(LX, T4)
Definition dmma_tma_kernel.h:580
neko_dmma_tma_env
static int neko_dmma_tma_env()
Definition dmma_tma_kernel.h:555
dmma_tma_dudxyz_aligned
static bool dmma_tma_dudxyz_aligned(const void *du, const void *u, const void *dr, const void *ds, const void *dt, const void *jacinv)
Definition dmma_tma_kernel.h:483
cuda_have_tma
static bool cuda_have_tma()
Definition dmma_tma_kernel.h:165
dudxyz_kernel.h
elem_block_tune.h
neko_log.h
log_error
void log_error(char *msg)
log_message
void log_message(char *msg)
log_end_section
void log_end_section()
log_section
void log_section(char *msg)
CASE_DMMA
#define CASE_DMMA(LX, C)
cuda_dudxyz
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:59
CASE_DMMA_SEL
#define CASE_DMMA_SEL(LX, SEL)
CASE_KSTEP_SEL
#define CASE_KSTEP_SEL(LX, SEL)
CASE
#define CASE(LX)
CASE_1D_SEL
#define CASE_1D_SEL(LX, SEL)
tune_dudxyz
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, int *nw_sel, int *tw_sel)
Definition opr_dudxyz.cu:217
CASE_KSTEP
#define CASE_KSTEP(LX, C)
CASE_LARGE
#define CASE_LARGE(LX)
CASE_DMMA_TMA_SEL
#define CASE_DMMA_TMA_SEL(LX, SEL)
CASE_DMMA_TMA
#define CASE_DMMA_TMA(LX, C)
CASE_1D
#define CASE_1D(LX, C)