Neko 0.9.99
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
conv1_kernel.h File Reference
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Functions

template<typename T , const int LX, const int CHUNKS>
__global__ void conv1_kernel_1d (T *__restrict__ du, const T *__restrict__ u, const T *__restrict__ vx, const T *__restrict__ vy, const T *__restrict__ vz, 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__ jacinv)
 
template<typename T , const int LX>
__global__ void __launch_bounds__ (LX *LX, 3) conv1_kernel_kstep(T *__restrict__ du
 
 for (int k=0;k< LX;++k)
 
 __syncthreads ()
 

Variables

__global__ void const T *__restrict__ u
 
__global__ void const T *__restrict__ const T *__restrict__ vx
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ vy
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ vz
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dz
 
__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__ drdx
 
__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__ const T *__restrict__ dsdx
 
__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__ const T *__restrict__ const T *__restrict__ dtdx
 
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdy
 
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdy
 
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dtdy
 
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ drdz
 
__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__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dsdz
 
__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__ 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
 
__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__ 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
 
__shared__ T shdx [LX *LX] = dx[ij]
 
__shared__ T shdy [LX *LX] = dy[ij]
 
__shared__ T shdz [LX *LX] = dz[ij]
 
const int e = blockIdx.x
 
const int j = threadIdx.y
 
const int i = threadIdx.x
 
const int ij = i + j * LX
 
const int ele = e*LX*LX*LX
 
T ru [LX]
 
T rvx [LX]
 
T rvy [LX]
 
T rvz [LX]
 
T rjacinv [LX]
 

Function Documentation

◆ __launch_bounds__()

template<typename T , const int LX>
__global__ void __launch_bounds__ ( LX LX,
 
)

◆ __syncthreads()

__syncthreads ( )

◆ conv1_kernel_1d()

template<typename T , const int LX, const int CHUNKS>
__global__ void conv1_kernel_1d ( T *__restrict__  du,
const T *__restrict__  u,
const T *__restrict__  vx,
const T *__restrict__  vy,
const T *__restrict__  vz,
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__  jacinv 
)

Device kernel for convective terms

Definition at line 42 of file conv1_kernel.h.

Here is the call graph for this function:

◆ for()

for ( )

Definition at line 173 of file conv1_kernel.h.

Here is the call graph for this function:

Variable Documentation

◆ drdx

◆ drdy

◆ drdz

◆ dsdx

◆ dsdy

◆ dsdz

◆ dtdx

◆ dtdy

◆ dtdz

◆ dx

◆ dy

◆ dz

◆ e

Definition at line 156 of file conv1_kernel.h.

◆ ele

const int ele = e*LX*LX*LX

Definition at line 160 of file conv1_kernel.h.

◆ i

Definition at line 158 of file conv1_kernel.h.

◆ ij

const int ij = i + j * LX

Definition at line 159 of file conv1_kernel.h.

◆ j

const int j = threadIdx.y

Definition at line 157 of file conv1_kernel.h.

◆ jacinv

Initial value:
{
__shared__ T shu[LX *LX]
__global__ void dirichlet_apply_scalar_kernel(const int *__restrict__ msk, T *__restrict__ x, const T g, const int m)

Definition at line 148 of file conv1_kernel.h.

◆ rjacinv

T rjacinv[LX]

Definition at line 170 of file conv1_kernel.h.

◆ ru

T ru[LX]

Definition at line 166 of file conv1_kernel.h.

◆ rvx

T rvx[LX]

Definition at line 167 of file conv1_kernel.h.

◆ rvy

T rvy[LX]

Definition at line 168 of file conv1_kernel.h.

◆ rvz

T rvz[LX]

Definition at line 169 of file conv1_kernel.h.

◆ shdx

shdx[ij] = dx[ij]

Definition at line 152 of file conv1_kernel.h.

◆ shdy

shdy[ij] = dy[ij]

Definition at line 153 of file conv1_kernel.h.

◆ shdz

shdz[ij] = dz[ij]

Definition at line 154 of file conv1_kernel.h.

◆ u

Definition at line 132 of file conv1_kernel.h.

◆ vx

◆ vy

◆ vz