Neko
0.9.0
A portable framework for high-order spectral element flow simulations
|
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] |
T | ru [LX] |
T | rv [LX] |
T | rw [LX] |
T | ruw [LX] |
T | rvw [LX] |
T | rww [LX] |
T | rut |
T | rvt |
T | 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) |
__global__ void __launch_bounds__ | ( | LX * | LX, |
3 | |||
) |
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 | ( | ) |
__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 | ( | ) |
Definition at line 104 of file ax_helm_full_kernel.h.
__global__ void T *__restrict__ av |
Definition at line 44 of file ax_helm_full_kernel.h.
__global__ void T *__restrict__ T *__restrict__ aw |
Definition at line 45 of file ax_helm_full_kernel.h.
__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.
__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.
__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.
__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.
__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.
__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.
__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.
__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.
__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.
__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.
__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.
__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.
const int e = blockIdx.x |
Definition at line 93 of file ax_helm_full_kernel.h.
const int ele = e*LX*LX*LX |
Definition at line 97 of file ax_helm_full_kernel.h.
__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.
const int i = threadIdx.x |
Definition at line 95 of file ax_helm_full_kernel.h.
Definition at line 96 of file ax_helm_full_kernel.h.
Definition at line 343 of file ax_helm_full_kernel.h.
const int j = threadIdx.y |
Definition at line 94 of file ax_helm_full_kernel.h.
__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.
T ru |
Definition at line 81 of file ax_helm_full_kernel.h.
T rut |
Definition at line 89 of file ax_helm_full_kernel.h.
T ruw |
Definition at line 85 of file ax_helm_full_kernel.h.
T rv |
Definition at line 82 of file ax_helm_full_kernel.h.
T rvt |
Definition at line 90 of file ax_helm_full_kernel.h.
T rvw |
Definition at line 86 of file ax_helm_full_kernel.h.
T rw |
Definition at line 83 of file ax_helm_full_kernel.h.
T rwt |
Definition at line 91 of file ax_helm_full_kernel.h.
T rww |
Definition at line 87 of file ax_helm_full_kernel.h.
Definition at line 99 of file ax_helm_full_kernel.h.
Definition at line 66 of file ax_helm_full_kernel.h.
Definition at line 67 of file ax_helm_full_kernel.h.
__shared__ T shu |
Definition at line 69 of file ax_helm_full_kernel.h.
__shared__ T shur |
Definition at line 70 of file ax_helm_full_kernel.h.
__shared__ T shus |
Definition at line 71 of file ax_helm_full_kernel.h.
__shared__ T shv |
Definition at line 73 of file ax_helm_full_kernel.h.
__shared__ T shvr |
Definition at line 74 of file ax_helm_full_kernel.h.
__shared__ T shvs |
Definition at line 75 of file ax_helm_full_kernel.h.
__shared__ T shw |
Definition at line 77 of file ax_helm_full_kernel.h.
__shared__ T shwr |
Definition at line 78 of file ax_helm_full_kernel.h.
__shared__ T shws |
Definition at line 79 of file ax_helm_full_kernel.h.
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ u |
Definition at line 46 of file ax_helm_full_kernel.h.
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v |
Definition at line 47 of file ax_helm_full_kernel.h.
__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.
__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 |
Definition at line 63 of file ax_helm_full_kernel.h.