#include <stdlib.h>
#include <stdio.h>
#include <math.h>
#include <hip/hip_runtime.h>
#include <device/device_config.h>
#include <device/hip/check.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.

Classes
struct	neko_tune_verify_t

struct	neko_tune_rej_t

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_1D(LX, T1)

#define	NEKO_TUNE_LOG_KSTEP(LX, T2)

#define	NEKO_TUNE_LOG(LX, T1, T2)

#define	NEKO_TUNE_VERIFY_DEFAULT 1

#define	NEKO_TUNE_VERIFY_BLOCKS 256

#define	NEKO_TUNE_VERIFY_THREADS 256

#define	NEKO_TUNE_VERIFY_MAXIN 8

#define	NEKO_TUNE_MAXCAND 16

#define	NEKO_TUNE_VERIFY_CAND(VF, REJ, LAUNCH, LX, C, NAME)

#define	NEKO_TUNE_APPLY_REJ(T, REJ, N)

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 ()

static int	neko_tune_verify_on ()

static double	neko_tune_verify_tol ()

static __global__ void	neko_tune_fill_kernel (real *__restrict__ x, const int n, const unsigned int seed)

static __global__ void	neko_tune_diff_kernel (const real __restrict__ a, const real __restrict__ b, const int n, double *__restrict__ part)

static int	neko_tune_verify_begin (neko_tune_verify_t v, real outs, const int nout, void **ins, const int nin, const int n, hipStream_t stream)

static void	neko_tune_verify_ref (neko_tune_verify_t *v, hipStream_t stream)

static int	neko_tune_verify_check (neko_tune_verify_t v, hipStream_t stream, double maxdiff, double *maxref)

static void	neko_tune_verify_end (neko_tune_verify_t *v, hipStream_t stream)

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 79 of file elem_block_tune.h.

◆ NEKO_TUNE_APPLY_REJ

#define NEKO_TUNE_APPLY_REJ	(	T,
		REJ,
		N
	)

Value:

  do {                                                                        \
    for (int c = 0; c < (N); c++) {                                           \
      if ((REJ).rej[c]) { (T)[c] = NEKO_TUNE_INIT; }                          \
    }                                                                         \
  } while (0)

Definition at line 498 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)

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

◆ NEKO_TUNE_INIT

#define NEKO_TUNE_INIT 1.0e30f

Definition at line 82 of file elem_block_tune.h.

◆ NEKO_TUNE_LOG

#define NEKO_TUNE_LOG	(	LX,
		T1,
		T2
	)

Value:

  do {                                                                        \
    NEKO_TUNE_LOG_1D(LX, T1);                                                 \
    NEKO_TUNE_LOG_KSTEP(LX, T2);                                              \
  } while (0)

Definition at line 226 of file elem_block_tune.h.

◆ NEKO_TUNE_LOG_1D

#define NEKO_TUNE_LOG_1D	(	LX,
		T1
	)

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), NEKO_TUNE_US((T1)[c], iters));          \
      log_message(neko_log_buf);                                              \
    }                                                                         \
  } while (0)

Definition at line 205 of file elem_block_tune.h.

◆ NEKO_TUNE_LOG_KSTEP

#define NEKO_TUNE_LOG_KSTEP	(	LX,
		T2
	)

Value:

  do {                                                                        \
    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), NEKO_TUNE_US((T2)[c], iters));              \
      log_message(neko_log_buf);                                              \
    }                                                                         \
  } while (0)

Definition at line 215 of file elem_block_tune.h.

◆ NEKO_TUNE_MAXCAND

#define NEKO_TUNE_MAXCAND 16

Definition at line 281 of file elem_block_tune.h.

◆ NEKO_TUNE_TIME

#define NEKO_TUNE_TIME	(	T,
		LAUNCH,
		LX,
		C,
		ITERS
	)

Value:

  do {                                                                        \
    float t_;                                                                 \
    HIP_CHECK(hipEventRecord(start, stream));                                           \
    for (int i_ = 0; i_ < (ITERS); i_++) { LAUNCH(LX, C); }                   \
    HIP_CHECK(hipEventRecord(stop, stream));                                            \
    HIP_CHECK(hipEventSynchronize(stop));                                               \
    HIP_CHECK(hipEventElapsedTime(&t_, start, stop));                                   \
    if (t_ < (T)[C]) { (T)[C] = t_; }                                         \
  } while (0)

Definition at line 173 of file elem_block_tune.h.

◆ NEKO_TUNE_US

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

Definition at line 193 of file elem_block_tune.h.

◆ NEKO_TUNE_VERIFY_BLOCKS

#define NEKO_TUNE_VERIFY_BLOCKS 256

Definition at line 277 of file elem_block_tune.h.

◆ NEKO_TUNE_VERIFY_CAND

#define NEKO_TUNE_VERIFY_CAND	(	VF,
		REJ,
		LAUNCH,
		LX,
		C,
		NAME
	)

Value:

  do {                                                                        \
    double md_, mr_;                                                          \
    LAUNCH(LX, C);                                                            \
    (REJ).rej[C] = !neko_tune_verify_check(&(VF), stream, &md_, &mr_);        \
    if ((REJ).rej[C]) {                                                       \
      sprintf(neko_log_buf,                                                   \
              "%-5s cand %2d REJECTED: |diff| %.1e, |ref| %.1e",              \
              NAME, C, md_, mr_);                                             \
      log_message(neko_log_buf);                                              \
    }                                                                         \
  } while (0)

Definition at line 484 of file elem_block_tune.h.

◆ NEKO_TUNE_VERIFY_DEFAULT

#define NEKO_TUNE_VERIFY_DEFAULT 1

Definition at line 275 of file elem_block_tune.h.

◆ NEKO_TUNE_VERIFY_MAXIN

#define NEKO_TUNE_VERIFY_MAXIN 8

Definition at line 279 of file elem_block_tune.h.

◆ NEKO_TUNE_VERIFY_THREADS

#define NEKO_TUNE_VERIFY_THREADS 256

Definition at line 278 of file elem_block_tune.h.

◆ NEKO_TUNE_WARMUP

#define NEKO_TUNE_WARMUP 20

Definition at line 83 of file elem_block_tune.h.

Function Documentation

◆ neko_chunks_pin()

static int neko_chunks_pin ( )

static

Definition at line 122 of file elem_block_tune.h.

Here is the call graph for this function:

◆ neko_eb_pin()

static int neko_eb_pin ( )

static

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

Here is the call graph for this function:

◆ neko_tune_diff_kernel()

static __global__ void neko_tune_diff_kernel	(	const real *__restrict__	a,
		const real *__restrict__	b,
		const int	n,
		double *__restrict__	part
	)

static

Definition at line 318 of file elem_block_tune.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_tune_fill_kernel()

static __global__ void neko_tune_fill_kernel	(	real *__restrict__	x,
		const int	n,
		const unsigned int	seed
	)

static

Definition at line 304 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 146 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 138 of file elem_block_tune.h.

Here is the call graph for this function:

◆ neko_tune_verify_begin()

static int neko_tune_verify_begin	(	neko_tune_verify_t *	v,
		real **	outs,
		const int	nout,
		void ***	ins,
		const int	nin,
		const int	n,
		hipStream_t	stream
	)

static

Definition at line 381 of file elem_block_tune.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_tune_verify_check()

static int neko_tune_verify_check	(	neko_tune_verify_t *	v,
		hipStream_t	stream,
		double *	maxdiff,
		double *	maxref
	)

static

Definition at line 442 of file elem_block_tune.h.

Here is the call graph for this function:

◆ neko_tune_verify_end()

static void neko_tune_verify_end	(	neko_tune_verify_t *	v,
		hipStream_t	stream
	)

static

Definition at line 471 of file elem_block_tune.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_tune_verify_on()

static int neko_tune_verify_on ( )

static

Definition at line 283 of file elem_block_tune.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_tune_verify_ref()

static void neko_tune_verify_ref	(	neko_tune_verify_t *	v,
		hipStream_t	stream
	)

static

Definition at line 423 of file elem_block_tune.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ neko_tune_verify_tol()

static double neko_tune_verify_tol ( )

static

Definition at line 293 of file elem_block_tune.h.

Here is the call graph for this function:

Here is the caller graph for this function:

Classes

Macros

Functions

Macro Definition Documentation

◆ NEKO_EB_SWEEP_DEFAULT

◆ NEKO_TUNE_APPLY_REJ

◆ NEKO_TUNE_BEST

◆ NEKO_TUNE_FOR

◆ NEKO_TUNE_INIT

◆ NEKO_TUNE_LOG

◆ NEKO_TUNE_LOG_1D

◆ NEKO_TUNE_LOG_KSTEP

◆ NEKO_TUNE_MAXCAND

◆ NEKO_TUNE_TIME

◆ NEKO_TUNE_US

◆ NEKO_TUNE_VERIFY_BLOCKS

◆ NEKO_TUNE_VERIFY_CAND

◆ NEKO_TUNE_VERIFY_DEFAULT

◆ NEKO_TUNE_VERIFY_MAXIN

◆ NEKO_TUNE_VERIFY_THREADS

◆ NEKO_TUNE_WARMUP

Function Documentation

◆ neko_chunks_pin()

◆ neko_eb_pin()

◆ neko_eb_sweep()

◆ neko_tune_diff_kernel()

◆ neko_tune_fill_kernel()

◆ neko_tune_iters()

◆ neko_tune_rounds()

◆ neko_tune_verify_begin()

◆ neko_tune_verify_check()

◆ neko_tune_verify_end()

◆ neko_tune_verify_on()

◆ neko_tune_verify_ref()

◆ neko_tune_verify_tol()