Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
elem_block_tune.h File Reference
#include <stdlib.h>
Include dependency graph for elem_block_tune.h:
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Macros

#define NEKO_EB_SWEEP_DEFAULT   0
 
#define NEKO_TUNE_INIT   1.0e30f
 
#define NEKO_TUNE_WARMUP   20
 
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
 
#define NEKO_TUNE_BEST(T, BEST, N)
 
#define NEKO_TUNE_LOG(LX, T1, T2)
 

Functions

static int neko_eb_sweep ()
 
static int neko_eb_env ()
 
static int neko_chunks_env ()
 
static int neko_tune_rounds ()
 
static int neko_tune_iters ()
 

Macro Definition Documentation

◆ NEKO_EB_SWEEP_DEFAULT

#define NEKO_EB_SWEEP_DEFAULT   0

Shared autotuner scaffolding for the SEM operator kernels

Each operator's tune() picks between the 1d variant and the kstep variant, and – when sweeping is enabled – between the elements per block candidates in elem_block.h. The sampling is interleaved and min reduced: timing the variants one after another in a fixed order lets clock drift bias the comparison by candidate position, which no amount of extra iterations removes.

A tune function using these macros is expected to have start, stop and stream in scope, plus a CASE_1D(LX) macro and a kstep launch macro taking (LX, C).

Definition at line 63 of file elem_block_tune.h.

◆ NEKO_TUNE_BEST

#define NEKO_TUNE_BEST (   T,
  BEST,
  N 
)
Value:
do { \
for (int c = 1; c < (N); c++) { \
if ((T)[c] < (T)[BEST]) { BEST = c; } \
} \
} while (0)
__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)

Definition at line 131 of file elem_block_tune.h.

◆ NEKO_TUNE_INIT

#define NEKO_TUNE_INIT   1.0e30f

Definition at line 66 of file elem_block_tune.h.

◆ NEKO_TUNE_LOG

#define NEKO_TUNE_LOG (   LX,
  T1,
  T2 
)
Value:
do { \
for (int c = 0; c < NEKO_CHUNKS_CANDIDATES; c++) { \
if ((T1)[c] >= NEKO_TUNE_INIT) { continue; } \
sprintf(neko_log_buf, "1D ch=%-4d: %9.2f us/call", \
NEKO_CHUNKS_SEL(LX, c), (T1)[c] * 10.0); \
} \
for (int c = 0; c < NEKO_EB_CANDIDATES; c++) { \
if ((T2)[c] >= NEKO_TUNE_INIT) { continue; } \
sprintf(neko_log_buf, "KSTEP eb=%-4d: %9.2f us/call", \
NEKO_EB_SEL(LX, c), (T2)[c] * 10.0); \
} \
} while (0)
#define NEKO_CHUNKS_CANDIDATES
Definition elem_block.h:125
#define NEKO_EB_CANDIDATES
Definition elem_block.h:63
#define NEKO_EB_SEL(LX, SEL)
Definition elem_block.h:108
#define NEKO_CHUNKS_SEL(LX, SEL)
Definition elem_block.h:147
#define NEKO_TUNE_INIT

Definition at line 139 of file elem_block_tune.h.

◆ NEKO_TUNE_TIME

#define NEKO_TUNE_TIME (   T,
  LAUNCH,
  LX,
  C,
  ITERS 
)
Value:
do { \
float t_; \
for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); } \
if (t_ < (T)[C]) { (T)[C] = t_; } \
} while (0)

Definition at line 120 of file elem_block_tune.h.

◆ NEKO_TUNE_WARMUP

#define NEKO_TUNE_WARMUP   20

Definition at line 67 of file elem_block_tune.h.

Function Documentation

◆ neko_chunks_env()

static int neko_chunks_env ( )
static

Definition at line 92 of file elem_block_tune.h.

Here is the call graph for this function:

◆ neko_eb_env()

static int neko_eb_env ( )
static

Definition at line 80 of file elem_block_tune.h.

Here is the call graph for this function:

◆ neko_eb_sweep()

static int neko_eb_sweep ( )
static

Definition at line 69 of file elem_block_tune.h.

Here is the call graph for this function:

◆ neko_tune_iters()

static int neko_tune_iters ( )
static

Definition at line 111 of file elem_block_tune.h.

Here is the call graph for this function:

◆ neko_tune_rounds()

static int neko_tune_rounds ( )
static

Definition at line 103 of file elem_block_tune.h.

Here is the call graph for this function: