49 void *
dr,
void *
ds,
void *
dt,
60 void *
dr,
void *
ds,
void *
dt,
62 void *
w3,
int *nel,
int *lx) {
79#define CASE_1D(LX, C) \
80 cdtp_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> \
81 <<<nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, stream>>> \
82 ((real *) dtx, (real *) x, \
83 (real *) dr, (real *) ds, (real *) dt, \
84 (real *) dxt, (real *) dyt, (real *) dzt, \
86 CUDA_CHECK(cudaGetLastError());
89#define CASE_1D_SEL(LX, 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; \
97#define CASE_KSTEP(LX, C) \
98 cdtp_kernel_kstep<real, LX, NEKO_EB(LX, C)> \
99 <<<NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0, stream>>> \
100 ((real *) dtx, (real *) x, \
101 (real *) dr, (real *) ds, (real *) dt, \
102 (real *) dxt, (real *) dyt, (real *) dzt, \
103 (real *) w3, *nel); \
104 CUDA_CHECK(cudaGetLastError());
107#define CASE_KSTEP_SEL(LX, SEL) \
109 case 0: CASE_KSTEP(LX, 0); break; \
110 case 1: CASE_KSTEP(LX, 1); break; \
111 default: CASE_KSTEP(LX, 2); break; \
114#define CASE_DMMA(LX, C) \
115 cdtp_kernel_dmma<real, LX, NEKO_DMMA_NW(C)> \
116 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
117 ((real *) dtx, (real *) x, \
118 (real *) dr, (real *) ds, (real *) dt, \
119 (real *) dxt, (real *) dyt, (real *) dzt, \
120 (real *) w3, *nel); \
121 CUDA_CHECK(cudaGetLastError());
124#define CASE_DMMA_SEL(LX, SEL) \
126 case 0: CASE_DMMA(LX, 0); break; \
127 case 1: CASE_DMMA(LX, 1); break; \
128 default: CASE_DMMA(LX, 2); break; \
139#define CASE_DMMA_TMA(LX, C) \
140 cdtp_kernel_dmma_tma<real, LX, NEKO_DMMA_NW(C)> \
141 <<<NEKO_DMMA_NBLCKS(*nel, LX), NEKO_DMMA_NTHRDS(C), 0, stream>>> \
142 ((real *) dtx, (real *) x, \
143 (real *) dr, (real *) ds, (real *) dt, \
144 (real *) dxt, (real *) dyt, (real *) dzt, \
146 CUDA_CHECK(cudaGetLastError());
149#define CASE_DMMA_TMA_SEL(LX, SEL) \
151 case 0: CASE_DMMA_TMA(LX, 0); break; \
152 case 1: CASE_DMMA_TMA(LX, 1); break; \
153 default: CASE_DMMA_TMA(LX, 2); break; \
158 if(autotune[LX] == 0 ) { \
159 autotune[LX]=tune_cdtp<LX>(dtx, x, \
162 w3, nel, lx, &autotune_eb[LX], \
163 &autotune_ch[LX], &autotune_nw[LX], \
165 } else if (autotune[LX] == 1 ) { \
166 CASE_1D_SEL(LX, autotune_ch[LX]); \
167 } else if (autotune[LX] == 2 ) { \
168 CASE_KSTEP_SEL(LX, autotune_eb[LX]); \
169 } else if (autotune[LX] == 3 ) { \
170 CASE_DMMA_SEL(LX, autotune_nw[LX]); \
171 } else if (autotune[LX] == 4 ) { \
172 CASE_DMMA_TMA_SEL(LX, autotune_tw[LX]); \
176#define CASE_LARGE(LX) \
218template < const
int LX >
220 void *
dr,
void *
ds,
void *
dt,
317 "DMMA_TMA strategy not available for this config");
355 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 const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
__global__ void const T *__restrict__ x
__global__ void const T *__restrict__ const T *__restrict__ dr
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dxt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w3
#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 dmma_tma_cdtp_aligned(const void *dtx, const void *x, const void *dr, const void *ds, const void *dt)
#define NEKO_TUNE_LOG_DMMA_TMA(LX, T4)
static int neko_dmma_tma_env()
static bool cuda_have_tma()
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)
int tune_cdtp(void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel, int *nw_sel, int *tw_sel)
#define CASE_1D_SEL(LX, SEL)
#define CASE_KSTEP(LX, C)
#define CASE_DMMA_TMA_SEL(LX, SEL)
void cuda_cdtp(void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx)
#define CASE_DMMA_TMA(LX, C)