Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Macros | Functions
opr_opgrad.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 "opgrad_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_MFMA(LX, C)
 
#define CASE_MFMA_SEL(LX, SEL)
 
#define CASE(LX)
 

Functions

template<const int >
int tune_opgrad (void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx, int *eb_sel, int *ch_sel, int *nwf_sel)
 
void hip_opgrad (void *ux, void *uy, void *uz, void *u, void *dx, void *dy, void *dz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx)
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)
Value:
case LX: \
if(autotune[LX] == 0 ) { \
autotune[LX]=tune_opgrad<LX>(ux, uy, uz, u, \
dx, dy, dz, \
drdx, dsdx, dtdx, \
drdy, dsdy, dtdy, \
drdz, dsdz, dtdz, \
w3, nel, lx, &autotune_eb[LX], \
&autotune_ch[LX], \
&autotune_nwf[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]); \
} else if (autotune[LX] == 3 ) { \
CASE_MFMA_SEL(LX, autotune_nwf[LX]); \
} \
break
ale_add_kinematics_kernel
__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 ale_kinematics_kernel.h:52
dtdy
__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__ dtdy
Definition ax_helm_full_kernel.h:60
dtdx
__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__ dtdx
Definition ax_helm_full_kernel.h:59
dtdz
__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__ dtdz
Definition ax_helm_full_kernel.h:61
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
dsdz
__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__ dsdz
Definition ax_helm_full_kernel.h:58
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
u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
Definition ax_helm_full_kernel.h:46
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
drdz
__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__ drdz
Definition ax_helm_full_kernel.h:55
drdx
__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__ drdx
Definition ax_helm_full_kernel.h:53
dsdx
__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__ dsdx
Definition ax_helm_full_kernel.h:56
dsdy
__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__ dsdy
Definition ax_helm_full_kernel.h:57
drdy
__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__ drdy
Definition ax_helm_full_kernel.h:54
w3
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w3
Definition cdtp_kernel.h:117
uy
__global__ void T *__restrict__ uy
Definition opgrad_kernel.h:130
uz
__global__ void T *__restrict__ T *__restrict__ uz
Definition opgrad_kernel.h:131

◆ CASE_1D

#define CASE_1D (   LX,
  C 
)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME( \
opgrad_kernel_1d<real, LX, NEKO_CHUNKS(LX, C)> ), \
nblcks, NEKO_CHUNKS_NTHRDS(LX, C), 0, \
(hipStream_t) glb_cmd_queue, \
(real *) ux, (real *) uy, (real *) uz, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) drdx, (real *) dsdx, (real *) dtdx, \
(real *) drdy, (real *) dsdy, (real *) dtdy, \
(real *) drdz, (real *) dsdz, (real *) dtdz, \
(real *) w3); \
HIP_CHECK(hipGetLastError());
NEKO_CHUNKS
#define NEKO_CHUNKS(LX, C)
Definition elem_block.h:145
NEKO_CHUNKS_NTHRDS
#define NEKO_CHUNKS_NTHRDS(LX, C)
Definition elem_block.h:146
opgrad_kernel_1d
__global__ void opgrad_kernel_1d(T *__restrict__ ux, T *__restrict__ uy, T *__restrict__ uz, const T *__restrict__ u, const T *__restrict__ dx, const T *__restrict__ dy, const T *__restrict__ dz, const T *__restrict__ drdx, const T *__restrict__ dsdx, const T *__restrict__ dtdx, const T *__restrict__ drdy, const T *__restrict__ dsdy, const T *__restrict__ dtdy, const T *__restrict__ drdz, const T *__restrict__ dsdz, const T *__restrict__ dtdz, const T *__restrict__ w3)
Definition opgrad_kernel.h:46
real
double real
Definition device_config.h:12

◆ 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_1D
#define CASE_1D(LX, C)

◆ CASE_KSTEP

#define CASE_KSTEP (   LX,
  C 
)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME( \
opgrad_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 *) ux, (real *) uy, (real *) uz, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) drdx, (real *) dsdx, (real *) dtdx, \
(real *) drdy, (real *) dsdy, (real *) dtdy, \
(real *) drdz, (real *) dsdz, (real *) dtdz, \
(real *) w3, *nel); \
HIP_CHECK(hipGetLastError());
NEKO_EB
#define NEKO_EB(LX, C)
Definition elem_block.h:104
NEKO_EB_NBLCKS
#define NEKO_EB_NBLCKS(NELV, LX, C)
Definition elem_block.h:106
NEKO_EB_NTHRDS
#define NEKO_EB_NTHRDS(LX, C)
Definition elem_block.h:105

◆ 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_KSTEP
#define CASE_KSTEP(LX, C)

◆ CASE_MFMA

#define CASE_MFMA (   LX,
  C 
)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME( \
opgrad_kernel_mfma<real, LX, NEKO_MFMA_NWF(C)> ), \
NEKO_MFMA_NBLCKS(*nel, LX, C), NEKO_MFMA_NTHRDS(C), 0, \
(hipStream_t) glb_cmd_queue, \
(real *) ux, (real *) uy, (real *) uz, (real *) u, \
(real *) dx, (real *) dy, (real *) dz, \
(real *) drdx, (real *) dsdx, (real *) dtdx, \
(real *) drdy, (real *) dsdy, (real *) dtdy, \
(real *) drdz, (real *) dsdz, (real *) dtdz, \
(real *) w3, *nel); \
HIP_CHECK(hipGetLastError());
NEKO_MFMA_NWF
#define NEKO_MFMA_NWF(C)
Definition mfma_kernel.h:263
NEKO_MFMA_NTHRDS
#define NEKO_MFMA_NTHRDS(C)
Definition mfma_kernel.h:265
NEKO_MFMA_NBLCKS
#define NEKO_MFMA_NBLCKS(NELV, LX, C)
Definition mfma_kernel.h:495

◆ CASE_MFMA_SEL

#define CASE_MFMA_SEL (   LX,
  SEL 
)
Value:
switch (SEL) { \
case 0: CASE_MFMA(LX, 0); break; \
case 1: CASE_MFMA(LX, 1); break; \
case 2: CASE_MFMA(LX, 2); break; \
case 3: CASE_MFMA(LX, 3); break; \
default: CASE_MFMA(LX, 4); break; \
}
CASE_MFMA
#define CASE_MFMA(LX, C)

Function Documentation

◆ hip_opgrad()

void hip_opgrad ( void *  ux,
void *  uy,
void *  uz,
void *  u,
void *  dx,
void *  dy,
void *  dz,
void *  drdx,
void *  dsdx,
void *  dtdx,
void *  drdy,
void *  dsdy,
void *  dtdy,
void *  drdz,
void *  dsdz,
void *  dtdz,
void *  w3,
int *  nel,
int *  lx 
)

Fortran wrapper for device hip convective terms

Definition at line 62 of file opr_opgrad.hip.

Here is the call graph for this function:

◆ tune_opgrad()

template<const int >
int tune_opgrad ( void *  ux,
void *  uy,
void *  uz,
void *  u,
void *  dx,
void *  dy,
void *  dz,
void *  drdx,
void *  dsdx,
void *  dtdx,
void *  drdy,
void *  dsdy,
void *  dtdy,
void *  drdz,
void *  dsdz,
void *  dtdz,
void *  w3,
int *  nel,
int *  lx,
int *  eb_sel,
int *  ch_sel,
int *  nwf_sel 
)

Definition at line 202 of file opr_opgrad.hip.

Here is the call graph for this function: