Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
Macros | Functions
ax_helm_svv_one_sided.cu File Reference
#include <stdio.h>
#include <stdlib.h>
#include <device/device_config.h>
#include <device/cuda/check.h>
#include "ax_helm_svv_one_sided_kernel.h"

Go to the source code of this file.

Macros

#define LAUNCH(LX)
 
#define CASE(LX)
 
#define CASE_LARGE(LX)
 

Functions

void cuda_ax_helm_svv_one_sided (void *w, void *u, void *dx, void *dy, void *dz, void *h1, void *drdx, void *drdy, void *drdz, void *dsdx, void *dsdy, void *dsdz, void *dtdx, void *dtdy, void *dtdz, void *jacinv, void *w3, void *h1_svv, void *filter_r, void *filter_s, void *filter_t, int *nelv, int *lx)
 

Macro Definition Documentation

◆ CASE

#define CASE (   LX)
Value:
case LX: \
LAUNCH(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

◆ CASE_LARGE

#define CASE_LARGE (   LX)
Value:
case LX: \
if (!shared_configured[LX]) { \
CUDA_CHECK(cudaFuncSetAttribute( \
ax_helm_svv_one_sided_kernel<real, LX>, \
cudaFuncAttributeMaxDynamicSharedMemorySize, shared_size)); \
shared_configured[LX] = true; \
} \
LAUNCH(LX); \
break

◆ LAUNCH

#define LAUNCH (   LX)
Value:
ax_helm_svv_one_sided_kernel<real, LX> \
<<<blocks, threads, shared_size, stream>>>( \
(real *) w, (real *) u, \
(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, (real *) h1_svv, \
(real *) filter_r, (real *) filter_s, (real *) filter_t); \
CUDA_CHECK(cudaGetLastError())
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
w
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
Definition ax_helm_full_kernel.h:48
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
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
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
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
real
double real
Definition device_config.h:12

Function Documentation

◆ cuda_ax_helm_svv_one_sided()

void cuda_ax_helm_svv_one_sided ( void *  w,
void *  u,
void *  dx,
void *  dy,
void *  dz,
void *  h1,
void *  drdx,
void *  drdy,
void *  drdz,
void *  dsdx,
void *  dsdy,
void *  dsdz,
void *  dtdx,
void *  dtdy,
void *  dtdz,
void *  jacinv,
void *  w3,
void *  h1_svv,
void *  filter_r,
void *  filter_s,
void *  filter_t,
int *  nelv,
int *  lx 
)

Fortran wrapper for the fused CUDA SVV Helmholtz operator.

Definition at line 46 of file ax_helm_svv_one_sided.cu.

Here is the call graph for this function:
Here is the caller graph for this function: