Neko 1.99.1
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Macros | Functions
set_convect_rst.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 "set_convect_rst_kernel.h"
#include <common/neko_log.h>

Go to the source code of this file.

Macros

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

Functions

template<const int >
int tune_set_convect_rst (void *cr, void *cs, void *ct, void *cx, void *cy, void *cz, void *drdx, void *dsdx, void *dtdx, void *drdy, void *dsdy, void *dtdy, void *drdz, void *dsdz, void *dtdz, void *w3, int *nel, int *lx)
 
void hip_set_convect_rst (void *cr, void *cs, void *ct, void *cx, void *cy, void *cz, 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_set_convect_rst<LX>(cr, cs, ct, \
cx, cy, cz, \
drdx, dsdx, dtdx, \
drdy, dsdy, dtdy, \
drdz, dsdz, dtdz, \
w3, nel, lx); \
} else if (autotune[LX] == 1 ) { \
CASE_1D(LX); \
} else if (autotune[LX] == 2 ) { \
CASE_KSTEP(LX); \
} \
break
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
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
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
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
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
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
cr
__global__ void const T *__restrict__ const T *__restrict__ cr
Definition convect_scalar_kernel.h:112
cs
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ cs
Definition convect_scalar_kernel.h:113
ct
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ ct
Definition convect_scalar_kernel.h:114
dirichlet_apply_scalar_kernel
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)
Definition dirichlet_kernel.h:42
cz
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ cz
Definition set_convect_rst_kernel.h:100
cx
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ cx
Definition set_convect_rst_kernel.h:98
cy
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ cy
Definition set_convect_rst_kernel.h:99

◆ CASE_1D

#define CASE_1D (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME(set_convect_rst_kernel_1d<real, LX, 1024> ), \
nblcks, nthrds_1d, 0, (hipStream_t) glb_cmd_queue, \
(real *) cr, (real *) cs, (real *) ct, \
(real *) cx, (real *) cy, (real *) cz, \
(real *) drdx, (real *) dsdx, (real *) dtdx, \
(real *) drdy, (real *) dsdy, (real *) dtdy, \
(real *) drdz, (real *) dsdz, (real *) dtdz, \
(real *) w3); \
HIP_CHECK(hipGetLastError());
real
double real
Definition device_config.h:12

◆ CASE_KSTEP

#define CASE_KSTEP (   LX)
Value:
hipLaunchKernelGGL( HIP_KERNEL_NAME(set_convect_rst_kernel_kstep<real, LX> ), \
nblcks, nthrds_kstep, 0, (hipStream_t) glb_cmd_queue, \
(real *) cr, (real *) cs, (real *) ct, \
(real *) cx, (real *) cy, (real *) cz, \
(real *) drdx, (real *) dsdx, (real *) dtdx, \
(real *) drdy, (real *) dsdy, (real *) dtdy, \
(real *) drdz, (real *) dsdz, (real *) dtdz, \
(real *) w3); \
HIP_CHECK(hipGetLastError());

Function Documentation

◆ hip_set_convect_rst()

void hip_set_convect_rst ( void *  cr,
void *  cs,
void *  ct,
void *  cx,
void *  cy,
void *  cz,
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 60 of file set_convect_rst.hip.

Here is the call graph for this function:

◆ tune_set_convect_rst()

template<const int >
int tune_set_convect_rst ( void *  cr,
void *  cs,
void *  ct,
void *  cx,
void *  cy,
void *  cz,
void *  drdx,
void *  dsdx,
void *  dtdx,
void *  drdy,
void *  dsdy,
void *  dtdy,
void *  drdz,
void *  dsdz,
void *  dtdz,
void *  w3,
int *  nel,
int *  lx 
)

Definition at line 140 of file set_convect_rst.hip.

Here is the call graph for this function: