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

Go to the source code of this file.

Macros

#define CASE_VECTOR_KSTEP(LX)
 
#define CASE_VECTOR_KSTEP_PADDED(LX)
 
#define CASE_VECTOR(LX)
 
#define CASE_VECTOR_PADDED(LX)
 

Functions

void hip_ax_helm_stress_vector (void *au, void *av, void *aw, void *u, void *v, void *w, void *dx, void *dy, void *dz, void *dxt, void *dyt, void *dzt, void *h1, void *drdx, void *drdy, void *drdz, void *dsdx, void *dsdy, void *dsdz, void *dtdx, void *dtdy, void *dtdz, void *jacinv, void *w3, int *nelv, int *lx)
 
void hip_ax_helm_stress_vector_part2 (void *au, void *av, void *aw, void *u, void *v, void *w, void *h2, void *B, int *n)
 

Macro Definition Documentation

◆ CASE_VECTOR

#define CASE_VECTOR (   LX)
Value:
case LX: \
CASE_VECTOR_KSTEP(LX); \
break

◆ CASE_VECTOR_KSTEP

#define CASE_VECTOR_KSTEP (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_stress_kernel_vector_kstep<real, LX> ), \
nblcks, nthrds, 0, \
(hipStream_t) glb_cmd_queue, \
(real *) au, (real *) av, (real *) aw, \
(real *) u, (real *) v, (real *) w, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) drdx, (real *) drdy, (real *) drdz, \
(real *) dsdx, (real *) dsdy, (real *) dsdz, \
(real *) dtdx, (real *) dtdy, (real *) dtdz, \
(real *) jacinv, (real *) w3); \
HIP_CHECK(hipGetLastError());
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
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
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
aw
__global__ void T *__restrict__ T *__restrict__ aw
Definition: ax_helm_full_kernel.h:45
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
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition: ax_helm_full_kernel.h:48
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
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
av
__global__ void T *__restrict__ av
Definition: ax_helm_full_kernel.h:44
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
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
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
v
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
Definition: ax_helm_full_kernel.h:47
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
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
h1
__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__ h1
Definition: ax_helm_full_kernel.h:52
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
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:113
real
double real
Definition: device_config.h:12
glb_cmd_queue
void * glb_cmd_queue

◆ CASE_VECTOR_KSTEP_PADDED

#define CASE_VECTOR_KSTEP_PADDED (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME( ax_helm_stress_kernel_vector_kstep_padded<real, LX> ), \
nblcks, nthrds, 0, \
(hipStream_t) glb_cmd_queue, \
(real *) au, (real *) av, (real *) aw, \
(real *) u, (real *) v, (real *) w, \
(real *) dx, (real *) dy, (real *) dz, (real *) h1, \
(real *) drdx, (real *) drdy, (real *) drdz, \
(real *) dsdx, (real *) dsdy, (real *) dsdz, \
(real *) dtdx, (real *) dtdy, (real *) dtdz, \
(real *) jacinv, (real *) w3); \
HIP_CHECK(hipGetLastError());

◆ CASE_VECTOR_PADDED

#define CASE_VECTOR_PADDED (   LX)
Value:
case LX: \
CASE_VECTOR_KSTEP_PADDED(LX); \
break

Function Documentation

◆ hip_ax_helm_stress_vector()

void hip_ax_helm_stress_vector ( void *  au,
void *  av,
void *  aw,
void *  u,
void *  v,
void *  w,
void *  dx,
void *  dy,
void *  dz,
void *  dxt,
void *  dyt,
void *  dzt,
void *  h1,
void *  drdx,
void *  drdy,
void *  drdz,
void *  dsdx,
void *  dsdy,
void *  dsdz,
void *  dtdx,
void *  dtdy,
void *  dtdz,
void *  jacinv,
void *  w3,
int *  nelv,
int *  lx 
)

Fortran wrapper for device HIP Ax vector version

Definition at line 52 of file ax_helm_full.hip.

◆ hip_ax_helm_stress_vector_part2()

void hip_ax_helm_stress_vector_part2 ( void *  au,
void *  av,
void *  aw,
void *  u,
void *  v,
void *  w,
void *  h2,
void *  B,
int *  n 
)

Fortran wrapper for device HIP Ax vector version part2

Definition at line 128 of file ax_helm_full.hip.