Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
cdtp_kernel.h File Reference
#include "elem_block.h"
Include dependency graph for cdtp_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 , const int LX, const int CHUNKS>
__global__ void cdtp_kernel_1d (T *__restrict__ dtx, const T *__restrict__ x, const T *__restrict__ dr, const T *__restrict__ ds, const T *__restrict__ dt, const T *__restrict__ dxt, const T *__restrict__ dyt, const T *__restrict__ dzt, const T *__restrict__ w3)
 
template<typename T , const int LX, const int EB>
__global__ void __launch_bounds__ ((LX *LX *EB)) cdtp_kernel_kstep(T *__restrict__ dtx
 
 if (eb==0)
 
 for (int k=0;k< LX;++k)
 
 __syncthreads ()
 

Variables

__global__ void const T *__restrict__ x
 
__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__ dxt
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dyt
 
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dzt
 
__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__ w3
 
__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 int nelv
 
__shared__ T shdyt [LX *LX]
 
__shared__ T shdzt [LX *LX]
 
__shared__ T shtar [EB *LX *LX]
 
__shared__ T shtas [EB *LX *LX]
 
T rtar [LX]
 
T rtas [LX]
 
T rtat [LX]
 
const int eb = (EB == 1) ? 0 : threadIdx.z
 
const int e_blk = blockIdx.x * EB + eb
 
const bool active = (EB == 1) ? true : (e_blk < nelv)
 
const int e = active ? e_blk : (nelv - 1)
 
const int sh = eb * LX * LX
 
const int j = threadIdx.y
 
const int i = threadIdx.x
 
const int ij = i + j * LX
 
const int ele = e*LX*LX*LX
 

Function Documentation

◆ __launch_bounds__()

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

◆ __syncthreads()

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

◆ cdtp_kernel_1d()

template<typename T , const int LX, const int CHUNKS>
__global__ void cdtp_kernel_1d ( T *__restrict__  dtx,
const T *__restrict__  x,
const T *__restrict__  dr,
const T *__restrict__  ds,
const T *__restrict__  dt,
const T *__restrict__  dxt,
const T *__restrict__  dyt,
const T *__restrict__  dzt,
const T *__restrict__  w3 
)

Device kernel for \( D^T x \)

Definition at line 44 of file cdtp_kernel.h.

Here is the call graph for this function:

◆ for()

for ( )

Definition at line 158 of file cdtp_kernel.h.

Here is the call graph for this function:

◆ if()

if ( eb  = = 0)

Definition at line 150 of file cdtp_kernel.h.

Here is the call graph for this function:

Variable Documentation

◆ active

const bool active = (EB == 1) ? true : (e_blk < nelv)

Definition at line 142 of file cdtp_kernel.h.

◆ dr

Definition at line 109 of file cdtp_kernel.h.

◆ ds

◆ dt

◆ dxt

◆ dyt

◆ dzt

◆ e

const int e = active ? e_blk : (nelv - 1)

Definition at line 143 of file cdtp_kernel.h.

◆ e_blk

const int e_blk = blockIdx.x * EB + eb

Definition at line 138 of file cdtp_kernel.h.

◆ eb

const int eb = (EB == 1) ? 0 : threadIdx.z

Definition at line 137 of file cdtp_kernel.h.

◆ ele

const int ele = e*LX*LX*LX

Definition at line 148 of file cdtp_kernel.h.

◆ i

Definition at line 146 of file cdtp_kernel.h.

◆ ij

const int ij = i + j * LX

Definition at line 147 of file cdtp_kernel.h.

◆ j

const int j = threadIdx.y

Definition at line 145 of file cdtp_kernel.h.

◆ nelv

Initial value:
{
__global__ void ale_add_kinematics_kernel(const int n, T *__restrict__ wx, T *__restrict__ wy, T *__restrict__ wz, const T *__restrict__ x_ref, const T *__restrict__ y_ref, const T *__restrict__ z_ref, const T *__restrict__ phi, const T *__restrict__ x, const T *__restrict__ y, const T *__restrict__ z, const kinematics_params_t kin_params)

Definition at line 116 of file cdtp_kernel.h.

◆ rtar

T rtar[LX]

Definition at line 133 of file cdtp_kernel.h.

◆ rtas

T rtas[LX]

Definition at line 134 of file cdtp_kernel.h.

◆ rtat

T rtat[LX]

Definition at line 135 of file cdtp_kernel.h.

◆ sh

const int sh = eb * LX * LX

Definition at line 144 of file cdtp_kernel.h.

◆ shdyt

__shared__ T shdyt[LX *LX]

Definition at line 119 of file cdtp_kernel.h.

◆ shdzt

__shared__ T shdzt[LX *LX]

Definition at line 120 of file cdtp_kernel.h.

◆ shtar

__shared__ T shtar[EB *LX *LX]

Definition at line 122 of file cdtp_kernel.h.

◆ shtas

__shared__ T shtas[EB *LX *LX]

Definition at line 123 of file cdtp_kernel.h.

◆ w3

◆ x

Definition at line 108 of file cdtp_kernel.h.