1#ifndef __MATH_ELEM_BLOCK_TUNE_H__
2#define __MATH_ELEM_BLOCK_TUNE_H__
68#ifndef NEKO_EB_SWEEP_DEFAULT
69#define NEKO_EB_SWEEP_DEFAULT 1
72#define NEKO_TUNE_INIT 1.0e30f
73#define NEKO_TUNE_WARMUP 20
77 const char *
v =
getenv(
"NEKO_EB_TUNE");
80 return (
atoi(
v) != 0);
97 const char *
v =
getenv(
"NEKO_EB");
114 const char *
v =
getenv(
"NEKO_CHUNKS");
130 const char *
v =
getenv(
"NEKO_TUNE_ROUNDS");
133 return (n < 1) ? 1 : n;
138 const char *
v =
getenv(
"NEKO_TUNE_ITERS");
141 return (n < 1) ? 1 : n;
157#define NEKO_TUNE_FOR(C, ON, PIN, N) \
158 for (int C = (((PIN) >= 0) ? (PIN) : 0), \
159 C##_last_ = ((ON) ? (((PIN) >= 0) ? (PIN) + 1 : (N)) : 0); \
163#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
166 cudaEventRecord(start, stream); \
167 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
168 cudaEventRecord(stop, stream); \
169 cudaEventSynchronize(stop); \
170 cudaEventElapsedTime(&t_, start, stop); \
171 if (t_ < (T)[C]) { (T)[C] = t_; } \
183#define NEKO_TUNE_US(T, ITERS) ((T) * 1000.0 / (double) (ITERS))
185#define NEKO_TUNE_BEST(T, BEST, N) \
187 for (int c = 1; c < (N); c++) { \
188 if ((T)[c] < (T)[BEST]) { BEST = c; } \
193#define NEKO_TUNE_LOG(LX, T1, T2) \
195 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
196 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
197 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
198 NEKO_CHUNKS_SEL(LX, c), NEKO_TUNE_US((T1)[c], iters)); \
199 log_message(neko_log_buf); \
201 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
202 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
203 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
204 NEKO_EB_SEL(LX, c), NEKO_TUNE_US((T2)[c], iters)); \
205 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_chunks_pin()
static int neko_tune_rounds()
static int neko_tune_iters()
static int neko_eb_sweep()