Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
compressible_res_kernel.h
Go to the documentation of this file.
1/*
2 Copyright (c) 2025, The Neko Authors
3 All rights reserved.
4
5 Redistribution and use in source and binary forms, with or without
6 modification, are permitted provided that the following conditions
7 are met:
8
9 * Redistributions of source code must retain the above copyright
10 notice, this list of conditions and the following disclaimer.
11
12 * Redistributions in binary form must reproduce the above
13 copyright notice, this list of conditions and the following
14 disclaimer in the documentation and/or other materials provided
15 with the distribution.
16
17 * Neither the name of the authors nor the names of its
18 contributors may be used to endorse or promote products derived
19 from this software without specific prior written permission.
20
21 THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22 "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23 LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24 FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25 COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26 INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27 BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28 LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29 CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30 LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31 ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32 POSSIBILITY OF SUCH DAMAGE.
33*/
34#ifndef __FLUID_EULER_RES_KERNEL__
35#define __FLUID_EULER_RES_KERNEL__
36
37template< typename T >
38__global__ void compressible_res_part_visc_kernel(T * __restrict__ rhs_rho,
39 T * __restrict__ rhs_m_x,
40 T * __restrict__ rhs_m_y,
41 T * __restrict__ rhs_m_z,
42 T * __restrict__ rhs_E,
43 const T * __restrict__ visc_rho,
44 const T * __restrict__ visc_m_x,
45 const T * __restrict__ visc_m_y,
46 const T * __restrict__ visc_m_z,
47 const T * __restrict__ visc_E,
48 const T * __restrict__ Binv,
49 T * __restrict__ h1,
50 const int n) {
51
52 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
53 const int str = blockDim.x * gridDim.x;
54
55 for (int i = idx; i < n; i += str) {
56 const T Bi = Binv[i];
57 rhs_rho[i] = -rhs_rho[i] - Bi * visc_rho[i];
58 rhs_m_x[i] = -rhs_m_x[i] - Bi * visc_m_x[i];
59 rhs_m_y[i] = -rhs_m_y[i] - Bi * visc_m_y[i];
60 rhs_m_z[i] = -rhs_m_z[i] - Bi * visc_m_z[i];
61 rhs_E[i] = -rhs_E[i] - Bi * visc_E[i];
62 h1[i] = T(1);
63 }
64}
65
66template< typename T >
67__global__ void inviscid_res_part_mx_flux_kernel(T * __restrict__ f_x,
68 T * __restrict__ f_y,
69 T * __restrict__ f_z,
70 const T * __restrict__ m_x,
71 const T * __restrict__ m_y,
72 const T * __restrict__ m_z,
73 const T * __restrict__ rho_field,
74 const T * __restrict__ p,
75 const int n) {
76
77 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
78 const int str = blockDim.x * gridDim.x;
79
80 for (int i = idx; i < n; i += str) {
81 f_x[i] = m_x[i] * m_x[i] / rho_field[i] + p[i];
82 f_y[i] = m_x[i] * m_y[i] / rho_field[i];
83 f_z[i] = m_x[i] * m_z[i] / rho_field[i];
84 }
85}
86
87template< typename T >
88__global__ void inviscid_res_part_my_flux_kernel(T * __restrict__ f_x,
89 T * __restrict__ f_y,
90 T * __restrict__ f_z,
91 const T * __restrict__ m_x,
92 const T * __restrict__ m_y,
93 const T * __restrict__ m_z,
94 const T * __restrict__ rho_field,
95 const T * __restrict__ p,
96 const int n) {
97
98 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
99 const int str = blockDim.x * gridDim.x;
100
101 for (int i = idx; i < n; i += str) {
102 f_x[i] = m_y[i] * m_x[i] / rho_field[i];
103 f_y[i] = m_y[i] * m_y[i] / rho_field[i] + p[i];
104 f_z[i] = m_y[i] * m_z[i] / rho_field[i];
105 }
106}
107
108template< typename T >
109__global__ void inviscid_res_part_mz_flux_kernel(T * __restrict__ f_x,
110 T * __restrict__ f_y,
111 T * __restrict__ f_z,
112 const T * __restrict__ m_x,
113 const T * __restrict__ m_y,
114 const T * __restrict__ m_z,
115 const T * __restrict__ rho_field,
116 const T * __restrict__ p,
117 const int n) {
118
119 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
120 const int str = blockDim.x * gridDim.x;
121
122 for (int i = idx; i < n; i += str) {
123 f_x[i] = m_z[i] * m_x[i] / rho_field[i];
124 f_y[i] = m_z[i] * m_y[i] / rho_field[i];
125 f_z[i] = m_z[i] * m_z[i] / rho_field[i] + p[i];
126 }
127}
128
129template< typename T >
130__global__ void inviscid_res_part_E_flux_kernel(T * __restrict__ f_x,
131 T * __restrict__ f_y,
132 T * __restrict__ f_z,
133 const T * __restrict__ m_x,
134 const T * __restrict__ m_y,
135 const T * __restrict__ m_z,
136 const T * __restrict__ rho_field,
137 const T * __restrict__ p,
138 const T * __restrict__ E,
139 const int n) {
140
141 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
142 const int str = blockDim.x * gridDim.x;
143
144 for (int i = idx; i < n; i += str) {
145 f_x[i] = (E[i] + p[i]) * m_x[i] / rho_field[i];
146 f_y[i] = (E[i] + p[i]) * m_y[i] / rho_field[i];
147 f_z[i] = (E[i] + p[i]) * m_z[i] / rho_field[i];
148 }
149}
150
151template< typename T >
152__global__ void compressible_res_part_coef_mult_kernel(T * __restrict__ rhs_rho,
153 T * __restrict__ rhs_m_x,
154 T * __restrict__ rhs_m_y,
155 T * __restrict__ rhs_m_z,
156 T * __restrict__ rhs_E,
157 const T * __restrict__ mult,
158 const int n) {
159
160 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
161 const int str = blockDim.x * gridDim.x;
162
163 for (int i = idx; i < n; i += str) {
164 rhs_rho[i] = rhs_rho[i] * mult[i];
165 rhs_m_x[i] = rhs_m_x[i] * mult[i];
166 rhs_m_y[i] = rhs_m_y[i] * mult[i];
167 rhs_m_z[i] = rhs_m_z[i] * mult[i];
168 rhs_E[i] = rhs_E[i] * mult[i];
169 }
170}
171
172template< typename T >
173__global__ void compressible_res_part_rk_sum_kernel(T * __restrict__ rho,
174 T * __restrict__ m_x,
175 T * __restrict__ m_y,
176 T * __restrict__ m_z,
177 T * __restrict__ E,
178 const T * __restrict__ k_rho_i,
179 const T * __restrict__ k_m_x_i,
180 const T * __restrict__ k_m_y_i,
181 const T * __restrict__ k_m_z_i,
182 const T * __restrict__ k_E_i,
183 const T dt,
184 const T c,
185 const int n) {
186
187 const int idx = blockIdx.x * blockDim.x + threadIdx.x;
188 const int str = blockDim.x * gridDim.x;
189
190 for (int i = idx; i < n; i += str) {
191 rho[i] = rho[i] + dt * c * k_rho_i[i];
192 m_x[i] = m_x[i] + dt * c * k_m_x_i[i];
193 m_y[i] = m_y[i] + dt * c * k_m_y_i[i];
194 m_z[i] = m_z[i] + dt * c * k_m_z_i[i];
195 E[i] = E[i] + dt * c * k_E_i[i];
196 }
197}
198
199#endif // __FLUID_EULER_RES_KERNEL__
ale_add_kinematics_kernel
__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 ale_kinematics_kernel.h:52
i
const int i
Definition ax_helm_full_kernel.h:95
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 ax_helm_full_kernel.h:52
dt
__global__ void const T *__restrict__ const T *__restrict__ const T *__restrict__ const T *__restrict__ dt
Definition cdtp_kernel.h:113
compressible_res_part_coef_mult_kernel
__global__ void compressible_res_part_coef_mult_kernel(T *__restrict__ rhs_rho, T *__restrict__ rhs_m_x, T *__restrict__ rhs_m_y, T *__restrict__ rhs_m_z, T *__restrict__ rhs_E, const T *__restrict__ mult, const int n)
Definition compressible_res_kernel.h:152
inviscid_res_part_my_flux_kernel
__global__ void inviscid_res_part_my_flux_kernel(T *__restrict__ f_x, T *__restrict__ f_y, T *__restrict__ f_z, const T *__restrict__ m_x, const T *__restrict__ m_y, const T *__restrict__ m_z, const T *__restrict__ rho_field, const T *__restrict__ p, const int n)
Definition compressible_res_kernel.h:88
compressible_res_part_visc_kernel
__global__ void compressible_res_part_visc_kernel(T *__restrict__ rhs_rho, T *__restrict__ rhs_m_x, T *__restrict__ rhs_m_y, T *__restrict__ rhs_m_z, T *__restrict__ rhs_E, const T *__restrict__ visc_rho, const T *__restrict__ visc_m_x, const T *__restrict__ visc_m_y, const T *__restrict__ visc_m_z, const T *__restrict__ visc_E, const T *__restrict__ Binv, T *__restrict__ h1, const int n)
Definition compressible_res_kernel.h:38
inviscid_res_part_mz_flux_kernel
__global__ void inviscid_res_part_mz_flux_kernel(T *__restrict__ f_x, T *__restrict__ f_y, T *__restrict__ f_z, const T *__restrict__ m_x, const T *__restrict__ m_y, const T *__restrict__ m_z, const T *__restrict__ rho_field, const T *__restrict__ p, const int n)
Definition compressible_res_kernel.h:109
compressible_res_part_rk_sum_kernel
__global__ void compressible_res_part_rk_sum_kernel(T *__restrict__ rho, T *__restrict__ m_x, T *__restrict__ m_y, T *__restrict__ m_z, T *__restrict__ E, const T *__restrict__ k_rho_i, const T *__restrict__ k_m_x_i, const T *__restrict__ k_m_y_i, const T *__restrict__ k_m_z_i, const T *__restrict__ k_E_i, const T dt, const T c, const int n)
Definition compressible_res_kernel.h:173
inviscid_res_part_E_flux_kernel
__global__ void inviscid_res_part_E_flux_kernel(T *__restrict__ f_x, T *__restrict__ f_y, T *__restrict__ f_z, const T *__restrict__ m_x, const T *__restrict__ m_y, const T *__restrict__ m_z, const T *__restrict__ rho_field, const T *__restrict__ p, const T *__restrict__ E, const int n)
Definition compressible_res_kernel.h:130
inviscid_res_part_mx_flux_kernel
__global__ void inviscid_res_part_mx_flux_kernel(T *__restrict__ f_x, T *__restrict__ f_y, T *__restrict__ f_z, const T *__restrict__ m_x, const T *__restrict__ m_y, const T *__restrict__ m_z, const T *__restrict__ rho_field, const T *__restrict__ p, const int n)
Definition compressible_res_kernel.h:67