1#ifndef __MATH_ELEM_BLOCK_TUNE_H__
2#define __MATH_ELEM_BLOCK_TUNE_H__
66#ifndef NEKO_EB_SWEEP_DEFAULT
67#define NEKO_EB_SWEEP_DEFAULT 1
70#define NEKO_TUNE_INIT 1.0e30f
71#define NEKO_TUNE_WARMUP 20
75 const char *
v =
getenv(
"NEKO_EB_TUNE");
78 return (
atoi(
v) != 0);
86 const char *
v =
getenv(
"NEKO_EB");
98 const char *
v =
getenv(
"NEKO_CHUNKS");
109 const char *
v =
getenv(
"NEKO_TUNE_ROUNDS");
112 return (n < 1) ? 1 : n;
117 const char *
v =
getenv(
"NEKO_TUNE_ITERS");
120 return (n < 1) ? 1 : n;
124#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
127 HIP_CHECK(hipEventRecord(start, stream)); \
128 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
129 HIP_CHECK(hipEventRecord(stop, stream)); \
130 HIP_CHECK(hipEventSynchronize(stop)); \
131 HIP_CHECK(hipEventElapsedTime(&t_, start, stop)); \
132 if (t_ < (T)[C]) { (T)[C] = t_; } \
144#define NEKO_TUNE_US(T, ITERS) ((T) * 1000.0 / (double) (ITERS))
146#define NEKO_TUNE_BEST(T, BEST, N) \
148 for (int c = 1; c < (N); c++) { \
149 if ((T)[c] < (T)[BEST]) { BEST = c; } \
154#define NEKO_TUNE_LOG(LX, T1, T2) \
156 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
157 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
158 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
159 NEKO_CHUNKS_SEL(LX, c), NEKO_TUNE_US((T1)[c], iters)); \
160 log_message(neko_log_buf); \
162 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
163 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
164 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
165 NEKO_EB_SEL(LX, c), NEKO_TUNE_US((T2)[c], iters)); \
166 log_message(neko_log_buf); \
__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_tune_rounds()
static int neko_tune_iters()
static int neko_chunks_env()
static int neko_eb_sweep()