opr_cdtp.hip File Reference

#include <string.h>
#include <stdlib.h>
#include <stdio.h>
#include <hip/hip_runtime.h>
#include <device/device_config.h>
#include <device/hip/check.h>
#include "cdtp_kernel.h"
#include "elem_block_tune.h"
#include <common/neko_log.h>

Go to the source code of this file.

Macros
#define	CASE_1D(LX, C)
#define	CASE_1D_SEL(LX, SEL)
#define	CASE_KSTEP(LX, C)
#define	CASE_KSTEP_SEL(LX, SEL)
#define	CASE(LX)
#define	CASE_LARGE(LX)

Functions
template<const int >
int	tune_cdtp (void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel)
void	hip_cdtp (void *dtx, void *x, void *dr, void *ds, void *dt, void *dxt, void *dyt, void *dzt, void *w3, int *nel, int *lx)

Macro Definition Documentation

CASE

#define CASE

(

LX

)

Value:

    case LX:                                                                    \
      if(autotune[LX] == 0 ) {                                                  \
        autotune[LX]=tune_cdtp<LX>(dtx, x,                                      \
                                   dr, ds, dt,                                  \
                                   dxt, dyt, dzt,                               \
                                   w3, nel, lx, &autotune_eb[LX],               \
                                   &autotune_ch[LX]);                           \
      } else if (autotune[LX] == 1 ) {                                          \
        CASE_1D_SEL(LX, autotune_ch[LX]);                                       \
      } else if (autotune[LX] == 2 ) {                                          \
        CASE_KSTEP_SEL(LX, autotune_eb[LX]);                                    \
      }                                                                         \
      break

CASE_1D

#define CASE_1D	(		LX,
			C
	)

Value:

    hipLaunchKernelGGL( HIP_KERNEL_NAME(                                        \
                       cdtp_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> ),          \
                       nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0,                    \
                       (hipStream_t) glb_cmd_queue,                             \
                       (real *) dtx, (real *) x,                                \
                       (real *) dr, (real *) ds, (real *) dt,                   \
                       (real *) dxt, (real *) dyt, (real *) dzt,                \
                       (real *) w3);                                            \
    HIP_CHECK(hipGetLastError());

CASE_1D_SEL

#define CASE_1D_SEL	(		LX,
			SEL
	)

Value:

    switch (SEL) {                                                              \
    case 0:  CASE_1D(LX, 0); break;                                             \
    case 1:  CASE_1D(LX, 1); break;                                             \
    case 2:  CASE_1D(LX, 2); break;                                             \
    default: CASE_1D(LX, 3); break;                                             \
    }

CASE_KSTEP

#define CASE_KSTEP	(		LX,
			C
	)

Value:

    hipLaunchKernelGGL( HIP_KERNEL_NAME(                                        \
                        cdtp_kernel_kstep<real, LX, NEKO_EB(LX, C)> ),          \
                       NEKO_EB_NBLCKS(*nel, LX, C), NEKO_EB_NTHRDS(LX, C), 0,   \
                       (hipStream_t) glb_cmd_queue,                             \
                       (real *) dtx, (real *) x,                                \
                       (real *) dr, (real *) ds, (real *) dt,                   \
                       (real *) dxt, (real *) dyt, (real *) dzt,                \
                       (real *) w3, *nel);                                      \
    HIP_CHECK(hipGetLastError());

CASE_KSTEP_SEL

#define CASE_KSTEP_SEL	(		LX,
			SEL
	)

Value:

    switch (SEL) {                                                              \
    case 0:  CASE_KSTEP(LX, 0); break;                                          \
    case 1:  CASE_KSTEP(LX, 1); break;                                          \
    default: CASE_KSTEP(LX, 2); break;                                          \
    }

CASE_LARGE

#define CASE_LARGE

(

LX

)

Value:

    case LX:                                                                    \
      CASE_KSTEP(LX, 0);                                                        \
      break

Function Documentation

hip_cdtp()

void hip_cdtp	(	void *	dtx,
		void *	x,
		void *	dr,
		void *	ds,
		void *	dt,
		void *	dxt,
		void *	dyt,
		void *	dzt,
		void *	w3,
		int *	nel,
		int *	lx
	)

Fortran wrapper for device hip \( D^T X \)

Definition at line 59 of file opr_cdtp.hip.

Here is the call graph for this function:

tune_cdtp()

template<const int >

int tune_cdtp	(	void *	dtx,
		void *	x,
		void *	dr,
		void *	ds,
		void *	dt,
		void *	dxt,
		void *	dyt,
		void *	dzt,
		void *	w3,
		int *	nel,
		int *	lx,
		int *	eb_sel,
		int *	ch_sel
	)

Definition at line 171 of file opr_cdtp.hip.

Here is the call graph for this function: