1#ifndef __MATH_ELEM_BLOCK_TUNE_H__
2#define __MATH_ELEM_BLOCK_TUNE_H__
62#ifndef NEKO_EB_SWEEP_DEFAULT
63#define NEKO_EB_SWEEP_DEFAULT 0
66#define NEKO_TUNE_INIT 1.0e30f
67#define NEKO_TUNE_WARMUP 20
71 const char *
v =
getenv(
"NEKO_EB_TUNE");
74 return (
atoi(
v) != 0);
82 const char *
v =
getenv(
"NEKO_EB");
94 const char *
v =
getenv(
"NEKO_CHUNKS");
105 const char *
v =
getenv(
"NEKO_TUNE_ROUNDS");
108 return (n < 1) ? 1 : n;
113 const char *
v =
getenv(
"NEKO_TUNE_ITERS");
116 return (n < 1) ? 1 : n;
120#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS) \
123 HIP_CHECK(hipEventRecord(start, stream)); \
124 for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
125 HIP_CHECK(hipEventRecord(stop, stream)); \
126 HIP_CHECK(hipEventSynchronize(stop)); \
127 HIP_CHECK(hipEventElapsedTime(&t_, start, stop)); \
128 if (t_ < (T)[C]) { (T)[C] = t_; } \
131#define NEKO_TUNE_BEST(T, BEST, N) \
133 for (int c = 1; c < (N); c++) { \
134 if ((T)[c] < (T)[BEST]) { BEST = c; } \
139#define NEKO_TUNE_LOG(LX, T1, T2) \
141 for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
142 if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
143 sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
144 NEKO_CHUNKS_SEL(LX, c), (T1)[c] * 10.0); \
145 log_message(neko_log_buf); \
147 for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
148 if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
149 sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
150 NEKO_EB_SEL(LX, c), (T2)[c] * 10.0); \
151 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()