49 void *
dr,
void *
ds,
void *
dt,
50 void *
dx,
void *
dy,
void *
dz,
59 void *
dr,
void *
ds,
void *
dt,
60 void *
dx,
void *
dy,
void *
dz,
61 void *
jacinv,
int *nel,
int *lx) {
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, \
81 CUDA_CHECK(cudaGetLastError());
84#define CASE_1D_SEL(LX, 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; \
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());
102#define CASE_KSTEP_SEL(LX, SEL) \
104 case 0: CASE_KSTEP(LX, 0); break; \
105 case 1: CASE_KSTEP(LX, 1); break; \
106 default: CASE_KSTEP(LX, 2); break; \
111 if(autotune[LX] == 0 ) { \
112 autotune[LX]=tune_dudxyz<LX>(du, u, \
115 jacinv, nel, lx, &autotune_eb[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]); \
124#define CASE_LARGE(LX) \
166template < const
int LX >
168 void *
dr,
void *
ds,
void *
dt,
169 void *
dx,
void *
dy,
void *
dz,
246 for (
int r = 0; r <
rounds; r++) {
__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__ 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
#define NEKO_EB_CANDIDATES
#define NEKO_EB_SEL(LX, SEL)
#define NEKO_CHUNKS_SEL(LX, SEL)
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
static int neko_tune_rounds()
#define NEKO_TUNE_LOG(LX, T1, T2)
#define NEKO_TUNE_BEST(T, BEST, N)
static int neko_tune_iters()
static int neko_chunks_env()
static int neko_eb_sweep()
void log_error(char *msg)
void log_message(char *msg)
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)
#define CASE_KSTEP_SEL(LX, SEL)
#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)