Neko  0.9.0
A portable framework for high-order spectral element flow simulations
ax_helm_full_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>
__global__ void __launch_bounds__ (LX *LX, 3) ax_helm_stress_kernel_vector_kstep(T *__restrict__ au
 
 for (int k=0;k< LX;++k)
 
 __syncthreads ()
 
template<typename T >
__global__ void ax_helm_stress_kernel_vector_part2 (T *__restrict__ au, T *__restrict__ av, T *__restrict__ aw, const T *__restrict__ u, const T *__restrict__ v, const T *__restrict__ w, const T *__restrict__ h2, const T *__restrict__ B, const int n)
 

Variables

__global__ void T *__restrict__ av
 
__global__ void T *__restrict__ T *__restrict__ aw
 
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u
 
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
 
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ w
 
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dx
 
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dy
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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
 
__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__ const T *__restrict__ weight3
 
__shared__ T shdy [LX *LX] = dy[ij]
 
__shared__ T shdz [LX *LX] = dz[ij]
 
__shared__ T shu [LX *LX]
 
__shared__ T shur [LX *LX]
 
__shared__ T shus [LX *LX]
 
__shared__ T shv [LX *LX]
 
__shared__ T shvr [LX *LX]
 
__shared__ T shvs [LX *LX]
 
__shared__ T shw [LX *LX]
 
__shared__ T shwr [LX *LX]
 
__shared__ T shws [LX *LX]
 
ru [LX]
 
rv [LX]
 
rw [LX]
 
ruw [LX]
 
rvw [LX]
 
rww [LX]
 
rut
 
rvt
 
rwt
 
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
 
 shdx [ij] = dx[ij]
 
const int ij_p = i + j*(LX+1)
 

Function Documentation

◆ __launch_bounds__()

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

Device kernels for Ax helm full

Device kernel for axhelm with padding in shared memory to remove bank conflicts when LX is a power of 2

◆ __syncthreads()

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

◆ ax_helm_stress_kernel_vector_part2()

template<typename T >
__global__ void ax_helm_stress_kernel_vector_part2 ( T *__restrict__  au,
T *__restrict__  av,
T *__restrict__  aw,
const T *__restrict__  u,
const T *__restrict__  v,
const T *__restrict__  w,
const T *__restrict__  h2,
const T *__restrict__  B,
const int  n 
)

Definition at line 535 of file ax_helm_full_kernel.h.

◆ for()

for ( )

Definition at line 104 of file ax_helm_full_kernel.h.

Variable Documentation

◆ av

__global__ void T *__restrict__ av

Definition at line 44 of file ax_helm_full_kernel.h.

◆ aw

__global__ void T *__restrict__ T *__restrict__ aw

Definition at line 45 of file ax_helm_full_kernel.h.

◆ 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 at line 53 of file ax_helm_full_kernel.h.

◆ 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 at line 54 of file ax_helm_full_kernel.h.

◆ 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 at line 55 of file ax_helm_full_kernel.h.

◆ 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 at line 56 of file ax_helm_full_kernel.h.

◆ 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 at line 57 of file ax_helm_full_kernel.h.

◆ 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 at line 58 of file ax_helm_full_kernel.h.

◆ 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 at line 59 of file ax_helm_full_kernel.h.

◆ 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 at line 60 of file ax_helm_full_kernel.h.

◆ 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 at line 61 of file ax_helm_full_kernel.h.

◆ dx

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

Definition at line 49 of file ax_helm_full_kernel.h.

◆ dy

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

Definition at line 50 of file ax_helm_full_kernel.h.

◆ 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 at line 51 of file ax_helm_full_kernel.h.

◆ e

const int e = blockIdx.x

Definition at line 93 of file ax_helm_full_kernel.h.

◆ ele

const int ele = e*LX*LX*LX

Definition at line 97 of file ax_helm_full_kernel.h.

◆ 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 at line 52 of file ax_helm_full_kernel.h.

◆ i

const int i = threadIdx.x

Definition at line 95 of file ax_helm_full_kernel.h.

◆ ij

const int ij = i + j*LX

Definition at line 96 of file ax_helm_full_kernel.h.

◆ ij_p

const int ij_p = i + j*(LX+1)

Definition at line 343 of file ax_helm_full_kernel.h.

◆ j

const int j = threadIdx.y

Definition at line 94 of file ax_helm_full_kernel.h.

◆ 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 at line 62 of file ax_helm_full_kernel.h.

◆ ru

T ru

Definition at line 81 of file ax_helm_full_kernel.h.

◆ rut

T rut

Definition at line 89 of file ax_helm_full_kernel.h.

◆ ruw

T ruw

Definition at line 85 of file ax_helm_full_kernel.h.

◆ rv

T rv

Definition at line 82 of file ax_helm_full_kernel.h.

◆ rvt

T rvt

Definition at line 90 of file ax_helm_full_kernel.h.

◆ rvw

T rvw

Definition at line 86 of file ax_helm_full_kernel.h.

◆ rw

T rw

Definition at line 83 of file ax_helm_full_kernel.h.

◆ rwt

T rwt

Definition at line 91 of file ax_helm_full_kernel.h.

◆ rww

T rww

Definition at line 87 of file ax_helm_full_kernel.h.

◆ shdx

shdx[ij_p] = dx[ij]

Definition at line 99 of file ax_helm_full_kernel.h.

◆ shdy

__shared__ T shdy = dy[ij]

Definition at line 66 of file ax_helm_full_kernel.h.

◆ shdz

__shared__ T shdz = dz[ij]

Definition at line 67 of file ax_helm_full_kernel.h.

◆ shu

__shared__ T shu

Definition at line 69 of file ax_helm_full_kernel.h.

◆ shur

__shared__ T shur

Definition at line 70 of file ax_helm_full_kernel.h.

◆ shus

__shared__ T shus

Definition at line 71 of file ax_helm_full_kernel.h.

◆ shv

__shared__ T shv

Definition at line 73 of file ax_helm_full_kernel.h.

◆ shvr

__shared__ T shvr

Definition at line 74 of file ax_helm_full_kernel.h.

◆ shvs

__shared__ T shvs

Definition at line 75 of file ax_helm_full_kernel.h.

◆ shw

__shared__ T shw

Definition at line 77 of file ax_helm_full_kernel.h.

◆ shwr

__shared__ T shwr

Definition at line 78 of file ax_helm_full_kernel.h.

◆ shws

__shared__ T shws

Definition at line 79 of file ax_helm_full_kernel.h.

◆ u

__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u

Definition at line 46 of file ax_helm_full_kernel.h.

◆ v

__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v

Definition at line 47 of file ax_helm_full_kernel.h.

◆ w

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

Definition at line 48 of file ax_helm_full_kernel.h.

◆ weight3

__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__ const T *__restrict__ weight3
Initial value:
{
__shared__ T shdx[LX * LX]
shdx[ij]

Definition at line 63 of file ax_helm_full_kernel.h.