Neko 1.99.9
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   1
 
#define NEKO_TUNE_INIT   1.0e30f
 
#define NEKO_TUNE_WARMUP   20
 
#define NEKO_TUNE_FOR(C, ON, PIN, N)
 
#define NEKO_TUNE_TIME(T, LAUNCH, LX, C, ITERS)
 
#define NEKO_TUNE_US(T, ITERS)   ((T) * 1000.0 / (double) (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_pin ()
 
static int neko_chunks_pin ()
 
static int neko_tune_rounds ()
 
static int neko_tune_iters ()
 

Macro Definition Documentation

◆ NEKO_EB_SWEEP_DEFAULT

#define NEKO_EB_SWEEP_DEFAULT   1

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.

NEKO_AUTOTUNE takes a formulation out of the tuner's hands, but only the formulation: the geometry candidates of the one it names are still measured and reported, and it takes that formulation's own variable (NEKO_EB, NEKO_CHUNKS, ...) to fix the geometry too. Pinning the formulation used to imply candidate 0, which silently answered a different question than the one an A/B run of two formulations is asking. See NEKO_TUNE_FOR() below.

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

Definition at line 69 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 185 of file elem_block_tune.h.

◆ NEKO_TUNE_FOR

#define NEKO_TUNE_FOR (   C,
  ON,
  PIN,
  N 
)
Value:
for (int C = (((PIN) >= 0) ? (PIN) : 0), \
C##_last_ = ((ON) ? (((PIN) >= 0) ? (PIN) + 1 : (N)) : 0); \
C < C##_last_; C++)

Definition at line 157 of file elem_block_tune.h.

◆ NEKO_TUNE_INIT

#define NEKO_TUNE_INIT   1.0e30f

Definition at line 72 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", \
} \
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", \
} \
} 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_US(T, ITERS)
#define NEKO_TUNE_INIT

Definition at line 193 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)
#define LAUNCH(LX)

Definition at line 163 of file elem_block_tune.h.

◆ NEKO_TUNE_US

#define NEKO_TUNE_US (   T,
  ITERS 
)    ((T) * 1000.0 / (double) (ITERS))

Definition at line 183 of file elem_block_tune.h.

◆ NEKO_TUNE_WARMUP

#define NEKO_TUNE_WARMUP   20

Definition at line 73 of file elem_block_tune.h.

Function Documentation

◆ neko_chunks_pin()

static int neko_chunks_pin ( )
static

Definition at line 112 of file elem_block_tune.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ neko_eb_pin()

static int neko_eb_pin ( )
static

Definition at line 95 of file elem_block_tune.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ neko_eb_sweep()

static int neko_eb_sweep ( )
static

Definition at line 75 of file elem_block_tune.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ neko_tune_iters()

static int neko_tune_iters ( )
static

Definition at line 136 of file elem_block_tune.h.

Here is the call graph for this function:
Here is the caller graph for this function:

◆ neko_tune_rounds()

static int neko_tune_rounds ( )
static

Definition at line 128 of file elem_block_tune.h.

Here is the call graph for this function:
Here is the caller graph for this function: