1#ifndef __MATH_ELEM_BLOCK_TUNE_H__
2#define __MATH_ELEM_BLOCK_TUNE_H__
61#ifndef NEKO_EB_SWEEP_DEFAULT
62#define NEKO_EB_SWEEP_DEFAULT 1
65#define NEKO_TUNE_INIT 1.0e30f
66#define NEKO_TUNE_WARMUP 20
70 const char *
v =
getenv(
"NEKO_EB_TUNE");
73 return (
atoi(
v) != 0);
81 const char *
v =
getenv(
"NEKO_EB");
93 const char *
v =
getenv(
"NEKO_CHUNKS");
104 const char *
v =
getenv(
"NEKO_TUNE_ROUNDS");
107 return (n < 1) ? 1 : n;
112 const char *
v =
getenv(
"NEKO_TUNE_ITERS");
115 return (n < 1) ? 1 : n;
119#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
122 cudaEventRecord(start, stream); \
123 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
124 cudaEventRecord(stop, stream); \
125 cudaEventSynchronize(stop); \
126 cudaEventElapsedTime(&t_, start, stop); \
127 if (t_ < (T)[C]) { (T)[C] = t_; } \
130#define NEKO_TUNE_BEST(T, BEST, N) \
132 for (int c = 1; c < (N); c++) { \
133 if ((T)[c] < (T)[BEST]) { BEST = c; } \
138#define NEKO_TUNE_LOG(LX, T1, T2) \
140 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
141 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
142 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
143 NEKO_CHUNKS_SEL(LX, c), (T1)[c] * 10.0); \
144 log_message(neko_log_buf); \
146 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
147 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
148 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
149 NEKO_EB_SEL(LX, c), (T2)[c] * 10.0); \
150 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()