Neko  0.9.0
A portable framework for high-order spectral element flow simulations
dudxyz_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 dudxyz_kernel_1d (T *__restrict__ du, const T *__restrict__ u, const T *__restrict__ dr, const T *__restrict__ ds, const T *__restrict__ dt, const T *__restrict__ dx, const T *__restrict__ dy, const T *__restrict__ dz, const T *__restrict__ jacinv)
 
template<typename T , const int LX>
__global__ void __launch_bounds__ (LX *LX, 3) dudxyz_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__ dr
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ ds
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
 
__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__ 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
 
ru [LX]
 
rdr [LX]
 
rds [LX]
 
rdt [LX]
 
rjacinv [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:

◆ dudxyz_kernel_1d()

template<typename T , const int LX, const int CHUNKS>
__global__ void dudxyz_kernel_1d ( T *__restrict__  du,
const T *__restrict__  u,
const T *__restrict__  dr,
const T *__restrict__  ds,
const T *__restrict__  dt,
const T *__restrict__  dx,
const T *__restrict__  dy,
const T *__restrict__  dz,
const T *__restrict__  jacinv 
)

Device kernel for derivative

Definition at line 42 of file dudxyz_kernel.h.

Here is the call graph for this function:

◆ for()

for ( )

Definition at line 144 of file dudxyz_kernel.h.

Variable Documentation

◆ dr

__global__ void const T* __restrict__ const T* __restrict__ dr

Definition at line 113 of file dudxyz_kernel.h.

◆ ds

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ ds

Definition at line 114 of file dudxyz_kernel.h.

◆ dt

__global__ void const T* __restrict__ const T* __restrict__ const T* __restrict__ const T* __restrict__ dt

Definition at line 115 of file dudxyz_kernel.h.

◆ dx

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

Definition at line 116 of file dudxyz_kernel.h.

◆ dy

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

Definition at line 117 of file dudxyz_kernel.h.

◆ dz

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

Definition at line 118 of file dudxyz_kernel.h.

◆ e

const int e = blockIdx.x

Definition at line 127 of file dudxyz_kernel.h.

◆ ele

const int ele = e*LX*LX*LX

Definition at line 131 of file dudxyz_kernel.h.

◆ i

const int i = threadIdx.x

Definition at line 129 of file dudxyz_kernel.h.

◆ ij

const int ij = i + j * LX

Definition at line 130 of file dudxyz_kernel.h.

◆ j

const int j = threadIdx.y

Definition at line 128 of file dudxyz_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__ jacinv
Initial value:
{
__shared__ T shu[LX * LX]
__shared__ T shu[LX *LX]

Definition at line 119 of file dudxyz_kernel.h.

◆ rdr

T rdr[LX]

Definition at line 138 of file dudxyz_kernel.h.

◆ rds

T rds[LX]

Definition at line 139 of file dudxyz_kernel.h.

◆ rdt

T rdt[LX]

Definition at line 140 of file dudxyz_kernel.h.

◆ rjacinv

T rjacinv[LX]

Definition at line 141 of file dudxyz_kernel.h.

◆ ru

T ru[LX]

Definition at line 137 of file dudxyz_kernel.h.

◆ shdx

shdx[ij] = dx[ij]

Definition at line 123 of file dudxyz_kernel.h.

◆ shdy

shdy[ij] = dy[ij]

Definition at line 124 of file dudxyz_kernel.h.

◆ shdz

shdz[ij] = dz[ij]

Definition at line 125 of file dudxyz_kernel.h.

◆ u

__global__ void const T* __restrict__ u

Definition at line 112 of file dudxyz_kernel.h.