Neko  0.9.0
A portable framework for high-order spectral element flow simulations
Macros | Functions
opr_dudxyz.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 "dudxyz_kernel.h"
#include <common/neko_log.h>
Include dependency graph for opr_dudxyz.hip:

Go to the source code of this file.

Macros

#define CASE_1D(LX)
 
#define CASE_KSTEP(LX)
 
#define CASE(LX)
 
#define CASE_LARGE(LX)
 

Functions

template<const int >
int tune_dudxyz (void *du, void *u, void *dr, void *ds, void *dt, void *dx, void *dy, void *dz, void *jacinv, int *nel, int *lx)
 
void hip_dudxyz (void *du, void *u, void *dr, void *ds, void *dt, void *dx, void *dy, void *dz, void *jacinv, int *nel, int *lx)
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)
Value:
case LX: \
if(autotune[LX] == 0 ) { \
autotune[LX]=tune_dudxyz<LX>(du, u, \
dr, ds, dt, \
dx, dy, dz, \
jacinv, nel, lx); \
} else if (autotune[LX] == 1 ) { \
CASE_1D(LX); \
} else if (autotune[LX] == 2 ) { \
CASE_KSTEP(LX); \
} \
break
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition: ax_helm_full_kernel.h:46
dx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
Definition: ax_helm_full_kernel.h:49
dz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
Definition: ax_helm_full_kernel.h:51
dy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
Definition: ax_helm_full_kernel.h:50
jacinv
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ jacinv
Definition: ax_helm_full_kernel.h:62
dr
__global__ void const T *__restrict__ const T *__restrict__ dr
Definition: cdtp_kernel.h:107
ds
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ ds
Definition: cdtp_kernel.h:108
dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
Definition: cdtp_kernel.h:109

◆ CASE_1D

#define CASE_1D (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME(dudxyz_kernel_1d<real, LX, 1024> ), \
nblcks, nthrds_1d, 0, (hipStream_t) glb_cmd_queue, \
(real *) du, (real *) u, \
(real *) dr, (real *) ds, (real *) dt, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) jacinv); \
HIP_CHECK(hipGetLastError());
real
double real
Definition: device_config.h:12
glb_cmd_queue
void * glb_cmd_queue

◆ CASE_KSTEP

#define CASE_KSTEP (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME(dudxyz_kernel_kstep<real, LX> ), \
nblcks, nthrds_kstep, 0, (hipStream_t) glb_cmd_queue, \
(real *) du, (real *) u, \
(real *) dr, (real *) ds, (real *) dt, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) jacinv); \
HIP_CHECK(hipGetLastError());

◆ CASE_LARGE

#define CASE_LARGE (   LX)
Value:
case LX: \
CASE_KSTEP(LX); \
break

Function Documentation

◆ hip_dudxyz()

void hip_dudxyz ( void *  du,
void *  u,
void *  dr,
void *  ds,
void *  dt,
void *  dx,
void *  dy,
void *  dz,
void *  jacinv,
int *  nel,
int *  lx 
)

Fortran wrapper for device hip derivative kernels

Definition at line 58 of file opr_dudxyz.hip.

◆ tune_dudxyz()

template<const int >
int tune_dudxyz ( void *  du,
void *  u,
void *  dr,
void *  ds,
void *  dt,
void *  dx,
void *  dy,
void *  dz,
void *  jacinv,
int *  nel,
int *  lx 
)

Definition at line 144 of file opr_dudxyz.hip.

Here is the call graph for this function: