1#ifndef __MATH_ELEM_BLOCK_TUNE_H__
2#define __MATH_ELEM_BLOCK_TUNE_H__
73#ifndef NEKO_EB_SWEEP_DEFAULT
74#define NEKO_EB_SWEEP_DEFAULT 1
77#define NEKO_TUNE_INIT 1.0e30f
78#define NEKO_TUNE_WARMUP 20
82 const char *
v =
getenv(
"NEKO_EB_TUNE");
85 return (
atoi(
v) != 0);
102 const char *
v =
getenv(
"NEKO_EB");
119 const char *
v =
getenv(
"NEKO_CHUNKS");
135 const char *
v =
getenv(
"NEKO_TUNE_ROUNDS");
138 return (n < 1) ? 1 : n;
143 const char *
v =
getenv(
"NEKO_TUNE_ITERS");
146 return (n < 1) ? 1 : n;
162#define NEKO_TUNE_FOR(C, ON, PIN, N) \
163 for (int C = (((PIN) >= 0) ? (PIN) : 0), \
164 C##_last_ = ((ON) ? (((PIN) >= 0) ? (PIN) + 1 : (N)) : 0); \
168#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
171 HIP_CHECK(hipEventRecord(start, stream)); \
172 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
173 HIP_CHECK(hipEventRecord(stop, stream)); \
174 HIP_CHECK(hipEventSynchronize(stop)); \
175 HIP_CHECK(hipEventElapsedTime(&t_, start, stop)); \
176 if (t_ < (T)[C]) { (T)[C] = t_; } \
188#define NEKO_TUNE_US(T, ITERS) ((T) * 1000.0 / (double) (ITERS))
190#define NEKO_TUNE_BEST(T, BEST, N) \
192 for (int c = 1; c < (N); c++) { \
193 if ((T)[c] < (T)[BEST]) { BEST = c; } \
200#define NEKO_TUNE_LOG_1D(LX, T1) \
202 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
203 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
204 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
205 NEKO_CHUNKS_SEL(LX, c), NEKO_TUNE_US((T1)[c], iters)); \
206 log_message(neko_log_buf); \
210#define NEKO_TUNE_LOG_KSTEP(LX, T2) \
212 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
213 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
214 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
215 NEKO_EB_SEL(LX, c), NEKO_TUNE_US((T2)[c], iters)); \
216 log_message(neko_log_buf); \
221#define NEKO_TUNE_LOG(LX, T1, T2) \
223 NEKO_TUNE_LOG_1D(LX, T1); \
224 NEKO_TUNE_LOG_KSTEP(LX, T2); \
__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__ v
#define NEKO_CHUNKS_CANDIDATES
#define NEKO_EB_CANDIDATES
#define NEKO_EB_SWEEP_DEFAULT
static int neko_chunks_pin()
static int neko_tune_rounds()
static int neko_tune_iters()
static int neko_eb_sweep()