35 #ifndef __COMMON_MAKEBDF_KERNEL__
36 #define __COMMON_MAKEBDF_KERNEL__
38 template<
typename T >
40 const T * __restrict__ ulag2,
41 const T * __restrict__ vlag1,
42 const T * __restrict__ vlag2,
43 const T * __restrict__ wlag1,
44 const T * __restrict__ wlag2,
48 const T * __restrict__
u,
49 const T * __restrict__
v,
50 const T * __restrict__
w,
51 const T * __restrict__ B,
60 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
61 const int str = blockDim.x * gridDim.x;
63 for (
int i = idx;
i < n;
i += str) {
64 T tb1_val =
u[
i] * B[
i] * bd2;
65 T tb2_val =
v[
i] * B[
i] * bd2;
66 T tb3_val =
w[
i] * B[
i] * bd2;
68 T ta1_val = ulag1[
i] * B[
i] * bd3;
69 T ta2_val = vlag1[
i] * B[
i] * bd3;
70 T ta3_val = wlag1[
i] * B[
i] * bd3;
77 tb1_val += ulag2[
i] * B[
i] * bd4;
78 tb2_val += vlag2[
i] * B[
i] * bd4;
79 tb3_val += wlag2[
i] * B[
i] * bd4;
82 bfx[
i] = bfx[
i] + tb1_val * (rho /
dt);
83 bfy[
i] = bfy[
i] + tb2_val * (rho /
dt);
84 bfz[
i] = bfz[
i] + tb3_val * (rho /
dt);
89 template<
typename T >
91 const T * __restrict__ s_laglag,
93 const T * __restrict__ s,
94 const T * __restrict__ B,
103 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
104 const int str = blockDim.x * gridDim.x;
106 for (
int i = idx;
i < n;
i += str) {
107 T tb1_val = s[
i] * B[
i] * bd2;
109 T ta1_val = s_lag[
i] * B[
i] * bd3;
114 tb1_val += s_laglag[
i] * B[
i] * bd4;
118 fs[
i] = fs[
i] + tb1_val * (rho /
dt);
__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__ u
__global__ void T *__restrict__ T *__restrict__ const T *__restrict__ const T *__restrict__ v
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
__global__ void scalar_makebdf_kernel(const T *__restrict__ s_lag, const T *__restrict__ s_laglag, T *__restrict__ fs, const T *__restrict__ s, const T *__restrict__ B, const T rho, const T dt, const T bd2, const T bd3, const T bd4, const int nbd, const int n)
__global__ void makebdf_kernel(const T *__restrict__ ulag1, const T *__restrict__ ulag2, const T *__restrict__ vlag1, const T *__restrict__ vlag2, const T *__restrict__ wlag1, const T *__restrict__ wlag2, T *__restrict__ bfx, T *__restrict__ bfy, T *__restrict__ bfz, const T *__restrict__ u, const T *__restrict__ v, const T *__restrict__ w, const T *__restrict__ B, const T rho, const T dt, const T bd2, const T bd3, const T bd4, const int nbd, const int n)