Neko  0.9.99
A portable framework for high-order spectral element flow simulations
lambda2_kernel.h File Reference
Include dependency graph for lambda2_kernel.h:
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Functions

template<typename T >
__inline__ __device__ T eigen_val_calc (T grad11, T grad12, T grad13, T grad21, T grad22, T grad23, T grad31, T grad32, T grad33)
 
template<typename T , const int LX, const int CHUNKS>
__global__ void lambda2_kernel_1d (T *__restrict__ lambda2, const T *__restrict__ u, const T *__restrict__ v, const T *__restrict__ w, 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) lambda2_kernel_kstep(T *__restrict__ lambda2
 
 for (int k=0;k< LX;++k)
 
 __syncthreads ()
 

Variables

__global__ void const T *__restrict__ u
 
__global__ void const T *__restrict__ const T *__restrict__ v
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ w
 
__global__ void 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__ dy
 
__global__ void 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__ 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__ 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__ 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__ 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__ 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__ 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__ 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__ 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__ 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__ jacinv
 
__shared__ T shv [LX *LX]
 
__shared__ T shw [LX *LX]
 
__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
 
ru [LX]
 
rv [LX]
 
rw [LX]
 

Function Documentation

◆ __launch_bounds__()

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

◆ __syncthreads()

__syncthreads ( )
Here is the caller graph for this function:

◆ eigen_val_calc()

template<typename T >
__inline__ __device__ T eigen_val_calc ( grad11,
grad12,
grad13,
grad21,
grad22,
grad23,
grad31,
grad32,
grad33 
)

Device kernel for lambda2

Definition at line 44 of file lambda2_kernel.h.

◆ for()

for ( )

Definition at line 282 of file lambda2_kernel.h.

◆ lambda2_kernel_1d()

template<typename T , const int LX, const int CHUNKS>
__global__ void lambda2_kernel_1d ( T *__restrict__  lambda2,
const T *__restrict__  u,
const T *__restrict__  v,
const T *__restrict__  w,
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 
)

Definition at line 99 of file lambda2_kernel.h.

Here is the call graph for this function:

Variable Documentation

◆ drdx

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ drdx

Definition at line 248 of file lambda2_kernel.h.

◆ 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__ drdy

Definition at line 251 of file lambda2_kernel.h.

◆ 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__ drdz

Definition at line 254 of file lambda2_kernel.h.

◆ 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__ dsdx

Definition at line 249 of file lambda2_kernel.h.

◆ 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__ dsdy

Definition at line 252 of file lambda2_kernel.h.

◆ 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__ dsdz

Definition at line 255 of file lambda2_kernel.h.

◆ 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__ dtdx

Definition at line 250 of file lambda2_kernel.h.

◆ 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__ dtdy

Definition at line 253 of file lambda2_kernel.h.

◆ 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__ dtdz

Definition at line 256 of file lambda2_kernel.h.

◆ dx

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dx

Definition at line 245 of file lambda2_kernel.h.

◆ dy

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dy

Definition at line 246 of file lambda2_kernel.h.

◆ dz

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dz

Definition at line 247 of file lambda2_kernel.h.

◆ e

const int e = blockIdx.x

Definition at line 267 of file lambda2_kernel.h.

◆ ele

const int ele = e*LX*LX*LX

Definition at line 271 of file lambda2_kernel.h.

◆ i

const int i = threadIdx.x

Definition at line 269 of file lambda2_kernel.h.

◆ ij

const int ij = i + j * LX

Definition at line 270 of file lambda2_kernel.h.

◆ j

const int j = threadIdx.y

Definition at line 268 of file lambda2_kernel.h.

◆ jacinv

__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__ jacinv
Initial value:
{
__shared__ T shu[LX * LX]
__shared__ T shu[LX *LX]

Definition at line 257 of file lambda2_kernel.h.

◆ ru

T ru[LX]

Definition at line 277 of file lambda2_kernel.h.

◆ rv

T rv[LX]

Definition at line 278 of file lambda2_kernel.h.

◆ rw

T rw[LX]

Definition at line 279 of file lambda2_kernel.h.

◆ shdx

shdx[ij] = dx[ij]

Definition at line 263 of file lambda2_kernel.h.

◆ shdy

shdy[ij] = dy[ij]

Definition at line 264 of file lambda2_kernel.h.

◆ shdz

shdz[ij] = dz[ij]

Definition at line 265 of file lambda2_kernel.h.

◆ shv

__shared__ T shv[LX *LX]

Definition at line 260 of file lambda2_kernel.h.

◆ shw

__shared__ T shw[LX *LX]

Definition at line 261 of file lambda2_kernel.h.

◆ u

__global__ void const T* __restrict__ u

Definition at line 242 of file lambda2_kernel.h.

◆ v

__global__ void const T* __restrict__ const T* __restrict__ v

Definition at line 243 of file lambda2_kernel.h.

◆ w

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ w

Definition at line 244 of file lambda2_kernel.h.