49 void *
vx,
void *
vy,
void *
vz,
50 void *
dx,
void *
dy,
void *
dz,
54 void *
jacinv,
int *nel,
int *gdim,
int *lx,
63 void *
vx,
void *
vy,
void *
vz,
64 void *
dx,
void *
dy,
void *
dz,
68 void *
jacinv,
int *nel,
int *gdim,
int *lx) {
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, \
95 CUDA_CHECK(cudaGetLastError());
98#define CASE_1D_SEL(LX, 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; \
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());
119#define CASE_KSTEP_SEL(LX, SEL) \
121 case 0: CASE_KSTEP(LX, 0); break; \
122 case 1: CASE_KSTEP(LX, 1); break; \
123 default: CASE_KSTEP(LX, 2); break; \
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());
139#define CASE_DMMA_SEL(LX, SEL) \
141 case 0: CASE_DMMA(LX, 0); break; \
142 case 1: CASE_DMMA(LX, 1); break; \
143 default: CASE_DMMA(LX, 2); break; \
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, \
166 CUDA_CHECK(cudaGetLastError());
169#define CASE_DMMA_TMA_SEL(LX, 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; \
178 if(autotune[LX] == 0 ) { \
179 autotune[LX]=tune_conv1<LX>(du, u, \
185 jacinv, nel, gdim, lx, &autotune_eb[LX], \
186 &autotune_ch[LX], &autotune_nw[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]); \
199#define CASE_LARGE(LX) \
241template < const
int LX >
243 void *
vx,
void *
vy,
void *
vz,
244 void *
dx,
void *
dy,
void *
dz,
248 void *
jacinv,
int *nel,
int *gdim,
int *lx,
347 "DMMA_TMA strategy not available for this config");
385 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__ 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
__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
#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()
#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)
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_section(char *msg)
#define CASE_DMMA_SEL(LX, SEL)
#define CASE_KSTEP_SEL(LX, SEL)
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)
#define CASE_1D_SEL(LX, SEL)
#define CASE_KSTEP(LX, C)
#define CASE_DMMA_TMA_SEL(LX, SEL)
#define CASE_DMMA_TMA(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)