Vlasiator ebf0dd394 on dev (v5.4.0 + 1054 commits)
Loading...
Searching...
No Matches
gpu_face_estimates.hpp
Go to the documentation of this file.
1/*
2 * This file is part of Vlasiator.
3 * Copyright 2010-2021 Finnish Meteorological Institute
4 *
5 * For details of usage, see the COPYING file and read the "Rules of the Road"
6 * at http://www.physics.helsinki.fi/vlasiator/
7 *
8 * This program is free software; you can redistribute it and/or modify
9 * it under the terms of the GNU General Public License as published by
10 * the Free Software Foundation; either version 2 of the License, or
11 * (at your option) any later version.
12 *
13 * This program is distributed in the hope that it will be useful,
14 * but WITHOUT ANY WARRANTY; without even the implied warranty of
15 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
16 * GNU General Public License for more details.
17 *
18 * You should have received a copy of the GNU General Public License along
19 * with this program; if not, write to the Free Software Foundation, Inc.,
20 * 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
21 */
22
23#ifndef GPU_FACE_ESTIMATES_H
24#define GPU_FACE_ESTIMATES_H
25
27#include "../definitions.h"
28
30
31/*enum for setting face value and derivative estimates. Implicit ones
32 not supported in the solver, so they are now not listed*/
37
38
48ARCH_DEV inline void compute_h8_left_face_value(const Realf* const values, int k, Realf &fv_l, const int index, const int stride)
49{
50 fv_l = (Realf)(1.0/840.0) * (
51 (Realf)(- 3.0) * values[(k-4)*stride+index]
52 + (Realf)(29.0) * values[(k-3)*stride+index]
53 - (Realf)(139.0) * values[(k-2)*stride+index]
54 + (Realf)(533.0) * values[(k-1)*stride+index]
55 + (Realf)(533.0) * values[k*stride+index]
56 - (Realf)(139.0) * values[(k+1)*stride+index]
57 + (Realf)(29.0) * values[(k+2)*stride+index]
58 - (Realf)(3.0) * values[(k+3)*stride+index]);
59}
60ARCH_DEV inline void compute_h7_left_face_derivative(const Realf* const values, int k, Realf &fd_l, const int index, const int stride){
61 fd_l = (Realf)(1.0/5040.0) * (
62 (Realf)(9.0) * values[(k-4)*stride+index]
63 - (Realf)(119.0) * values[(k-3)*stride+index]
64 + (Realf)(889.0) * values[(k-2)*stride+index]
65 - (Realf)(7175.0) * values[(k-1)*stride+index]
66 + (Realf)(7175.0) * values[k*stride+index]
67 - (Realf)(889.0) * values[(k+1)*stride+index]
68 + (Realf)(119.0) * values[(k+2)*stride+index]
69 - (Realf)(9.0) * values[(k+3)*stride+index]);
70}
71ARCH_DEV inline void compute_h6_left_face_value(const Realf* const values, int k, Realf &fv_l, const int index, const int stride)
72{
73 //compute left value
74 fv_l = (Realf)(1.0/60.0) * (values[(k-3)*stride+index]
75 - (Realf)(8.0) * values[(k-2)*stride+index]
76 + (Realf)(37.0) * values[(k-1)*stride+index]
77 + (Realf)(37.0) * values[k*stride+index]
78 - (Realf)(8.0) * values[(k+1)*stride+index]
79 + values[(k+2)*stride+index]);
80}
81ARCH_DEV inline void compute_h5_left_face_derivative(const Realf* const values, int k, Realf &fd_l, const int index, const int stride)
82{
83 fd_l = (Realf)(1.0/180.0) * ((Realf)(245.0) * (values[k*stride+index] - values[(k-1)*stride+index])
84 - (Realf)(25.0) * (values[(k+1)*stride+index] - values[(k-2)*stride+index])
85 + (Realf)(2.0) * (values[(k+2)*stride+index] - values[(k-3)*stride+index]));
86}
87ARCH_DEV inline void compute_h5_face_values(const Realf* const values, int k, Realf &fv_l, Realf &fv_r, const int index, const int stride)
88{
89 //compute left values
90 fv_l = (Realf)(1.0/60.0) * ((Realf)(- 3.0) * values[(k-2)*stride+index]
91 + (Realf)(27.0) * values[(k-1)*stride+index]
92 + (Realf)(47.0) * values[k*stride+index]
93 - (Realf)(13.0) * values[(k+1)*stride+index]
94 + (Realf)(2.0) * values[(k+2)*stride+index]);
95 fv_r = (Realf)(1.0/60.0) * ( (Realf)(2.0) * values[(k-2)*stride+index]
96 - (Realf)(13.0) * values[(k-1)*stride+index]
97 + (Realf)(47.0) * values[k*stride+index]
98 + (Realf)(27.0) * values[(k+1)*stride+index]
99 - (Realf)(3.0) * values[(k+2)*stride+index]);
100}
101ARCH_DEV inline void compute_h4_left_face_derivative(const Realf* const values, int k, Realf &fd_l, const int index, const int stride)
102{
103 fd_l = (Realf)(1.0/12.0) * ((Realf)(15.0) * (values[k*stride+index] - values[(k-1)*stride+index]) - (values[(k+1)*stride+index] - values[(k-2)*stride+index]));
104}
105ARCH_DEV inline void compute_h4_left_face_value(const Realf* const values, int k, Realf &fv_l, const int index, const int stride)
106{
107 //compute left value
108 fv_l = (Realf)(1.0/12.0) * ( (Realf)(- 1.0) * values[(k-2)*stride+index]
109 + (Realf)(7.0) * values[(k-1)*stride+index]
110 + (Realf)(7.0) * values[k*stride+index]
111 - (Realf)(1.0) * values[(k+1)*stride+index]);
112}
113
114// h is bin width (dv or dx)
115// u is values
116ARCH_DEV inline void compute_h4_left_face_value_nonuniform(const Realf* const h, const Realf* const u, int k, Realf &fv_l, const int index, const int stride) {
117 const Realf hkMinus2 = h[k-2];
118 const Realf hkMinus1 = h[k-1];
119 const Realf hk = h[k];
120 const Realf hkPlus1 = h[k+1];
121 fv_l = (
122 (Realf)(1.0) / ( hkMinus2 + hkMinus1 + hk + hkPlus1 )
123 * ( ( hkMinus2 + hkMinus1 ) * ( hk + hkPlus1 ) / ( hkMinus1 + hk )
124 * ( u[(k-1)*stride+index] * hk + u[k*stride+index] * hkMinus1 )
125 * ((Realf)(1.0) / ( hkMinus2 + hkMinus1 + hk ) + (Realf)(1.0) / ( hkMinus1 + hk + hkPlus1 ) )
126 + ( hk * ( hk + hkPlus1 ) ) / ( ( hkMinus2 + hkMinus1 + hk ) * (hkMinus2 + hkMinus1 ) )
127 * ( u[(k-1)*stride+index] * (hkMinus2 + (Realf)(2.0) * hkMinus1 ) - ( u[(k-2)*stride+index] * hkMinus1 ) )
128 + hkMinus1 * ( hkMinus2 + hkMinus1 ) / ( ( hkMinus1 + hk + hkPlus1 ) * ( hk + hkPlus1 ) )
129 * ( u[k*stride+index] * ( (Realf)(2.0) * hk + hkPlus1 ) - u[(k+1)*stride+index] * hk ) )
130 );
131}
132
133
134
135
136ARCH_DEV inline void compute_h3_left_face_derivative(const Realf* const values, int k, Realf &fv_l, const int index, const int stride)
137{
138 /*compute left value*/
139 fv_l = (Realf)(1.0/12.0) * ((Realf)(15.0) * (values[k*stride+index] - values[(k-1)*stride+index]) - (values[(k+1)*stride+index] - values[(k-2)*stride+index]));
140}
141
142ARCH_DEV inline void compute_filtered_face_values_derivatives(const Realf* const values, int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, Realf &fd_l, Realf &fd_r, const Realf threshold, const int index, const int stride)
143{
144 switch(order)
145 {
146 case h4:
147 compute_h4_left_face_value(values, k, fv_l, index, stride);
148 compute_h4_left_face_value(values, k + 1, fv_r, index, stride);
149 compute_h3_left_face_derivative(values, k, fd_l, index, stride);
150 compute_h3_left_face_derivative(values, k + 1, fd_r, index, stride);
151 break;
152 case h5:
153 compute_h5_face_values(values, k, fv_l, fv_r, index, stride);
154 compute_h4_left_face_derivative(values, k, fd_l, index, stride);
155 compute_h4_left_face_derivative(values, k + 1, fd_r, index, stride);
156 break;
157 default:
158 case h6:
159 compute_h6_left_face_value(values, k, fv_l, index, stride);
160 compute_h6_left_face_value(values, k + 1, fv_r, index, stride);
161 compute_h5_left_face_derivative(values, k, fd_l, index, stride);
162 compute_h5_left_face_derivative(values, k + 1, fd_r, index, stride);
163 break;
164 case h8:
165 compute_h8_left_face_value(values, k, fv_l, index, stride);
166 compute_h8_left_face_value(values, k + 1, fv_r, index, stride);
167 compute_h7_left_face_derivative(values, k, fd_l, index, stride);
168 compute_h7_left_face_derivative(values, k + 1, fd_r, index, stride);
169 break;
170 }
171 Realf slope_abs,slope_sign;
172 // scale values closer to 1 for more accurate slope limiter calculation
173 const Realf scale = (Realf)(1.0)/threshold;
174 slope_limiter(values[(k-1)*stride+index]*scale, values[k*stride+index]*scale, values[(k+1)*stride+index]*scale, slope_abs, slope_sign);
175 slope_abs = slope_abs*threshold;
176 //check for extrema, flatten if it is
177 bool is_extrema = (slope_abs == (Realf)(0.0));
178 if (is_extrema) {
179 fv_r = (is_extrema) ? values[k*stride+index] : fv_r;
180 fv_l = (is_extrema) ? values[k*stride+index] : fv_l;
181 fd_l = (is_extrema) ? (Realf)(0.0) : fd_l;
182 fd_r = (is_extrema) ? (Realf)(0.0) : fd_r;
183 }
184 //Fix left face if needed; boundary value is not bounded or slope is not consistent
185 bool filter = (values[(k-1)*stride+index] - fv_l) * (fv_l - values[k*stride+index]) < (Realf)(0.0) || slope_sign * fd_l < (Realf)(0.0);
186 if (filter) {
187 //Go to linear (PLM) estimates if not ok (this is always ok!)
188 fv_l= (filter) ? values[k*stride+index] - slope_sign * (Realf)(0.5) * slope_abs : fv_l;
189 fd_l= (filter) ? slope_sign * slope_abs : fd_l;
190 }
191 //Fix right face if needed; boundary value is not bounded or slope is not consistent
192 filter = (values[(k+1)*stride+index] - fv_r) * (fv_r - values[k*stride+index]) < (Realf)(0.0) || slope_sign * fd_r < (Realf)(0.0);
193 if (filter) {
194 //Go to linear (PLM) estimates if not ok (this is always ok!)
195 fv_r= (filter) ? values[k*stride+index] + slope_sign * (Realf)(0.5) * slope_abs : fv_r;
196 fd_r= (filter) ? slope_sign * slope_abs : fd_r;
197 }
198}
199
200/*Filters in section 2.6.1 of white et al. to be used for PPM
201 1) Checks for extrema and flattens them
202 2) Makes face values bounded
203 3) Makes sure face slopes are consistent with PLM slope
204*/
205ARCH_DEV inline void compute_filtered_face_values(const Realf* const values, int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, const Realf threshold, const int index, const int stride)
206{
207 switch(order)
208 {
209 case h4:
210 compute_h4_left_face_value(values, k, fv_l, index, stride);
211 compute_h4_left_face_value(values, k + 1, fv_r, index, stride);
212 break;
213 case h5:
214 compute_h5_face_values(values, k, fv_l, fv_r, index, stride);
215 break;
216 default:
217 case h6:
218 compute_h6_left_face_value(values, k, fv_l, index, stride);
219 compute_h6_left_face_value(values, k + 1, fv_r, index, stride);
220 break;
221 case h8:
222 compute_h8_left_face_value(values, k, fv_l, index, stride);
223 compute_h8_left_face_value(values, k + 1, fv_r, index, stride);
224 break;
225 }
226 Realf slope_abs, slope_sign;
227 // scale values closer to 1 for more accurate slope limiter calculation
228 const Realf scale = (Realf)(1.0)/threshold;
229 slope_limiter(values[(k-1)*stride+index]*scale, values[k*stride+index]*scale, values[(k+1)*stride+index]*scale, slope_abs, slope_sign);
230 slope_abs = slope_abs*threshold;
231
232 //check for extrema, flatten if it is
233 bool is_extrema = (slope_abs == (Realf)(0.0));
234 if (is_extrema) {
235 fv_r = (is_extrema) ? values[k*stride+index] : fv_r;
236 fv_l = (is_extrema) ? values[k*stride+index] : fv_l;
237 }
238 //Fix left face if needed; boundary value is not bounded
239 bool filter = (values[(k-1)*stride+index] - fv_l) * (fv_l - values[k*stride+index]) < (Realf)(0.0) ;
240 if (filter) {
241 //Go to linear (PLM) estimates if not ok (this is always ok!)
242 fv_l = (filter) ? values[k*stride+index] - slope_sign * (Realf)(0.5) * slope_abs : fv_l;
243 }
244 //Fix face if needed; boundary value is not bounded
245 filter = (values[(k+1)*stride+index] - fv_r) * (fv_r - values[k*stride+index]) < (Realf)(0.0);
246 if (filter) {
247 //Go to linear (PLM) estimates if not ok (this is always ok!)
248 fv_r = (filter) ? values[k*stride+index] + slope_sign * (Realf)(0.5) * slope_abs : fv_r;
249 }
250}
251
252ARCH_DEV inline void compute_filtered_face_values_nonuniform(const Realf * const dv, const Realf* const values,int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, const Realf threshold, const int index, const int stride){
253 switch(order){
254 case h4:
255 compute_h4_left_face_value_nonuniform(dv, values, k, fv_l, index, stride);
256 compute_h4_left_face_value_nonuniform(dv, values, k + 1, fv_r, index, stride);
257 break;
258 // case h5:
259 // compute_h5_face_values(dv, values, k, fv_l, fv_r, index, stride);
260 // break;
261 // case h6:
262 // compute_h6_left_face_value(dv, values, k, fv_l, index, stride);
263 // compute_h6_left_face_value(dv, values, k + 1, fv_r, index, stride);
264 // break;
265 // case h8:
266 // compute_h8_left_face_value(dv, values, k, fv_l, index, stride);
267 // compute_h8_left_face_value(dv, values, k + 1, fv_r, index, stride);
268 // break;
269 default:
270 printf("Order %d has not been implemented (yet)\n",order);
271 break;
272 }
273 Realf slope_abs,slope_sign;
274 if (threshold>(Realf)(0.0)) {
275 // scale values closer to 1 for more accurate slope limiter calculation
276 const Realf scale = (Realf)(1.0)/threshold;
277 slope_limiter(values[(k-1)*stride+index]*scale, values[k*stride+index]*scale, values[(k+1)*stride+index]*scale, slope_abs, slope_sign);
278 slope_abs = slope_abs*threshold;
279 } else {
280 slope_limiter(values[(k-1)*stride+index], values[k*stride+index], values[(k+1)*stride+index], slope_abs, slope_sign);
281 }
282
283 //check for extrema, flatten if it is
284 if (slope_abs == (Realf)(0.0)) {
285 fv_r = values[k*stride+index];
286 fv_l = values[k*stride+index];
287 }
288
289 //Fix left face if needed; boundary value is not bounded
290 if ((values[(k-1)*stride+index] - fv_l) * (fv_l - values[k*stride+index]) < (Realf)(0.0)) {
291 //Go to linear (PLM) estimates if not ok (this is always ok!)
292 fv_l=values[k*stride+index] - slope_sign * (Realf)(0.5) * slope_abs;
293 }
294
295 //Fix face if needed; boundary value is not bounded
296 if ((values[(k+1)*stride+index] - fv_r) * (fv_r - values[k*stride+index]) < (Realf)(0.0)) {
297 //Go to linear (PLM) estimates if not ok (this is always ok!)
298 fv_r=values[k*stride+index] + slope_sign * (Realf)(0.5) * slope_abs;
299 }
300}
301
302ARCH_DEV inline Realf get_D2aLim(const Realf* h, const Realf* values, int k, const Realf C, Realf & fv, const int index, const int stride) {
303
304 // Colella & Sekora, eq. 18
305 Realf invh2 = 1.0 / (h[k] * h[k]);
306 Realf d2a = invh2 * 3.0 * (values[k*stride +index] - 2.0 * fv + values[(k+1)*stride+index]);
307 Realf d2aL = invh2 * (values[(k-1)*stride+index] - 2.0 * values[k*stride +index] + values[(k+1)*stride+index]);
308 Realf d2aR = invh2 * (values[k*stride +index] - 2.0 * values[(k+1)*stride+index] + values[(k+2)*stride+index]);
309 Realf d2aLim;
310 if ( (d2a * d2aL >= 0) && (d2a * d2aR >= 0) && (d2a != 0) ) {
311 d2aLim = d2a / abs(d2a) * min(abs(d2a),min(C*abs(d2aL),C*abs(d2aR)));
312 } else {
313 d2aLim = 0.0;
314 }
315 return d2aLim;
316}
317
318ARCH_DEV inline void constrain_face_values(const Realf* h, const Realf* values,int k,Realf & fv_l, Realf & fv_r, const int index, const int stride) {
319
320 const Realf C = 1.25;
321 Realf invh2 = 1.0 / (h[k] * h[k]);
322
323 // Colella & Sekora, eq 19
324 Realf p_face = 0.5 * (values[k*stride+index] + values[(k+1)*stride+index])
325 - h[k] * h[k] / 3.0 * get_D2aLim(h,values,k ,C,fv_r, index, stride);
326 Realf m_face = 0.5 * (values[(k-1)*stride+index] + values[k*stride+index])
327 - h[k-1] * h[k-1] / 3.0 * get_D2aLim(h,values,k-1,C,fv_l, index, stride);
328
329 // Colella & Sekora, eq 21
330 Realf d2a = -2.0 * invh2 * 6.0 * (values[k*stride+index] - 3.0 * (m_face + p_face)); // a6,j from eq. 7
331 Realf d2aC = invh2 * (values[(k-1)*stride+index] - 2.0 * values[k*stride+index] + values[(k+1)*stride+index]);
332 // Note: Corrected the index of 2nd term in d2aL to k - 1.
333 // In the paper it is k but that is almost certainly an error.
334 Realf d2aL = invh2 * (values[(k-2)*stride+index] - 2.0 * values[(k-1)*stride+index] + values[k*stride+index]);
335 Realf d2aR = invh2 * (values[k*stride+index] - 2.0 * values[(k+1)*stride+index] + values[(k+2)*stride+index]);
336 Realf d2aLim;
337
338 // Colella & Sekora, eq 22
339 if ( (d2a * d2aL >= 0) && (d2a * d2aR >= 0) &&
340 (d2a * d2aC >= 0) && (d2a != 0) ) {
341
342 d2aLim = d2a / abs(d2a) * min(C * abs(d2aL), min(C * abs(d2aR), min(C * abs(d2aC), abs(d2a))));
343 } else {
344 d2aLim = 0.0;
345 if (d2a == 0.0) {
346 // Set a non-zero value for the denominator in eq. 23.
347 // According to the paper the ratio d2aLim/d2a should be 0
348 d2a = 1.0;
349 }
350 }
351
352 // Colella & Sekora, eq 23
353 fv_r = values[k*stride+index] + (p_face - values[k*stride+index]) * d2aLim / d2a;
354 fv_l = values[k*stride+index] + (m_face - values[k*stride+index]) * d2aLim / d2a;
355
356}
357
358ARCH_DEV inline void compute_filtered_face_values_nonuniform_conserving(const Realf * const dv, const Realf* const values,int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, const Realf threshold, const int index, const int stride){
359 switch(order){
360 case h4:
361 compute_h4_left_face_value_nonuniform(dv, values, k, fv_l, index, stride);
362 compute_h4_left_face_value_nonuniform(dv, values, k + 1, fv_r, index, stride);
363 break;
364 // case h5:
365 // compute_h5_face_values(dv, values, k, fv_l, fv_r, index, stride);
366 // break;
367 // case h6:
368 // compute_h6_left_face_value(dv, values, k, fv_l, index, stride);
369 // compute_h6_left_face_value(dv, values, k + 1, fv_r, index, stride);
370 // break;
371 // case h8:
372 // compute_h8_left_face_value(dv, values, k, fv_l, index, stride);
373 // compute_h8_left_face_value(dv, values, k + 1, fv_r, index, stride);
374 // break;
375 default:
376 printf("Order %d has not been implemented (yet)\n",order);
377 break;
378 }
379
380 Realf slope_abs,slope_sign;
381 if (threshold>0) {
382 // scale values closer to 1 for more accurate slope limiter calculation
383 const Realf scale = 1./threshold;
384 slope_limiter(values[(k-1)*stride+index]*scale, values[k*stride+index]*scale, values[(k+1)*stride+index]*scale, slope_abs, slope_sign);
385 slope_abs = slope_abs*threshold;
386 } else {
387 slope_limiter(values[(k-1)*stride+index], values[k*stride+index], values[(k+1)*stride+index], slope_abs, slope_sign);
388 }
389
390 //check for extrema
391 //bool is_extrema = (slope_abs == 0.0);
392 // bool filter_l = (values[(k-1)*stride+index] - fv_l) * (fv_l - values[k*stride+index]) < 0 ;
393 // bool filter_r = (values[(k+1)*stride+index] - fv_r) * (fv_r - values[k*stride+index]) < 0;
394 // if(horizontal_or(is_extrema) || horizontal_or(filter_l) || horizontal_or(filter_r)) {
395 // Colella & Sekora, eq. 20
396 if (((fv_r - values[k*stride+index]) * (values[k*stride+index] - fv_l) <= 0.0)
397 && ((values[(k-1)*stride+index] - values[k*stride+index]) * (values[k*stride+index] - values[(k+1)*stride+index]) <= 0.0)) {
398 constrain_face_values(dv, values, k, fv_l, fv_r, index, stride);
399
400 // fv_r = select(is_extrema, values[k], fv_r);
401 // fv_l = select(is_extrema, values[k], fv_l);
402 } else {
403
404 //Fix left face if needed; boundary value is not bounded
405 bool filter = (values[(k-1)*stride+index] - fv_l) * (fv_l - values[k*stride+index]) < 0 ;
406 if (filter) {
407 //Go to linear (PLM) estimates if not ok (this is always ok!)
408 fv_l=values[k*stride+index] - slope_sign * 0.5 * slope_abs;
409 }
410
411 //Fix face if needed; boundary value is not bounded
412 filter = (values[(k+1)*stride+index] - fv_r) * (fv_r - values[k*stride+index]) < 0;
413 if (filter) {
414 //Go to linear (PLM) estimates if not ok (this is always ok!)
415 fv_r=values[k*stride+index] + slope_sign * 0.5 * slope_abs;
416 }
417 }
418}
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436// ARCH_DEV inline void compute_h4_left_face_value_nonuniform(const Realf * const h, const Realf * const u, int k, Realf &fv_l, const int index, const int stride) {
437// fv_l = (
438// 1.0 / ( h[k - 2] + h[k - 1] + h[k] + h[k + 1] )
439// * ( ( h[k - 2] + h[k - 1] ) * ( h[k] + h[k + 1] ) / ( h[k - 1] + h[k] )
440// * ( u[(k-1)*stride+index] * h[k] + u[k*stride+index] * h[k - 1] )
441// * (1.0 / ( h[k - 2] + h[k - 1] + h[k] ) + 1.0 / ( h[k - 1] + h[k] + h[k + 1] ) )
442// + ( h[k] * ( h[k] + h[k + 1] ) ) / ( ( h[k - 2] + h[k - 1] + h[k] ) * (h[k - 2] + h[k - 1] ) )
443// * ( u[(k-1)*stride+index] * (h[k - 2] + 2.0 * h[k - 1] ) - ( u[(k-2)*stride+index] * h[k - 1] ) )
444// + h[k - 1] * ( h[k - 2] + h[k - 1] ) / ( ( h[k - 1] + h[k] + h[k + 1] ) * ( h[k] + h[k + 1] ) )
445// * ( u[k*stride+index] * ( 2.0 * h[k] + h[k + 1] ) - u[(k+1)*stride+index] * h[k] ) )
446// );
447// }
448
449
450// ARCH_DEV inline void compute_filtered_face_values_nonuniform(const Realf * const dv, const Realf * const values,int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, const Realf threshold, const int index, const int stride){
451// switch(order){
452// case h4:
453// compute_h4_left_face_value_nonuniform(dv, values, k, fv_l, index, stride);
454// compute_h4_left_face_value_nonuniform(dv, values, k + 1, fv_r, index, stride);
455// break;
456// // case h5:
457// // compute_h5_face_values(dv, values, k, fv_l, fv_r, index, stride);
458// // break;
459// // case h6:
460// // compute_h6_left_face_value(dv, values, k, fv_l, index, stride);
461// // compute_h6_left_face_value(dv, values, k + 1, fv_r, index, stride);
462// // break;
463// // case h8:
464// // compute_h8_left_face_value(dv, values, k, fv_l, index, stride);
465// // compute_h8_left_face_value(dv, values, k + 1, fv_r, index, stride);
466// // break;
467// default:
468// printf("Order %d has not been implemented (yet)\n",order);
469// break;
470// }
471// Realf slope_abs,slope_sign;
472// if (threshold>0) {
473// // scale values closer to 1 for more accurate slope limiter calculation
474// const Realf scale = 1./threshold;
475// slope_limiter(values[(k-1)*stride+index]*scale, values[k*stride+index]*scale, values[(k+1)*stride+index]*scale, slope_abs, slope_sign);
476// slope_abs = slope_abs*threshold;
477// } else {
478// slope_limiter(values[(k-1)*stride+index], values[k*stride+index], values[(k+1)*stride+index], slope_abs, slope_sign);
479// }
480
481// //check for extrema, flatten if it is
482// if (slope_abs == 0) {
483// fv_r = values[k*stride+index];
484// fv_l = values[k*stride+index];
485// }
486
487// //Fix left face if needed; boundary value is not bounded
488// if ((values[(k-1)*stride+index] - fv_l) * (fv_l - values[k*stride+index]) < 0) {
489// //Go to linear (PLM) estimates if not ok (this is always ok!)
490// fv_l=values[k*stride+index] - slope_sign * 0.5 * slope_abs;
491// }
492
493// //Fix face if needed; boundary value is not bounded
494// if ((values[(k+1)*stride+index] - fv_r) * (fv_r - values[k*stride+index]) < 0) {
495// //Go to linear (PLM) estimates if not ok (this is always ok!)
496// fv_r=values[k*stride+index] + slope_sign * 0.5 * slope_abs;
497// }
498// }
499
500
501
502#endif
#define ARCH_DEV
face_estimate_order
float Realf
Definition definitions.h:33
__global__ void vmesh::VelocityMesh **__restrict__ ColumnOffsets split::SplitVector< vmesh::GlobalID > Hashinator::Hashmap< vmesh::GlobalID, vmesh::LocalID > const uint *__restrict__ const Realf const int const int const Realf const Realf dv
ARCH_DEV void compute_h6_left_face_value(const Realf *const values, int k, Realf &fv_l, const int index, const int stride)
ARCH_DEV void compute_h5_face_values(const Realf *const values, int k, Realf &fv_l, Realf &fv_r, const int index, const int stride)
ARCH_DEV void compute_filtered_face_values_nonuniform(const Realf *const dv, const Realf *const values, int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, const Realf threshold, const int index, const int stride)
ARCH_DEV Realf get_D2aLim(const Realf *h, const Realf *values, int k, const Realf C, Realf &fv, const int index, const int stride)
ARCH_DEV void compute_h7_left_face_derivative(const Realf *const values, int k, Realf &fd_l, const int index, const int stride)
ARCH_DEV void compute_h5_left_face_derivative(const Realf *const values, int k, Realf &fd_l, const int index, const int stride)
ARCH_DEV void compute_filtered_face_values_nonuniform_conserving(const Realf *const dv, const Realf *const values, int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, const Realf threshold, const int index, const int stride)
ARCH_DEV void compute_h4_left_face_value_nonuniform(const Realf *const h, const Realf *const u, int k, Realf &fv_l, const int index, const int stride)
ARCH_DEV void compute_h4_left_face_derivative(const Realf *const values, int k, Realf &fd_l, const int index, const int stride)
ARCH_DEV void constrain_face_values(const Realf *h, const Realf *values, int k, Realf &fv_l, Realf &fv_r, const int index, const int stride)
ARCH_DEV void compute_filtered_face_values_derivatives(const Realf *const values, int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, Realf &fd_l, Realf &fd_r, const Realf threshold, const int index, const int stride)
ARCH_DEV void compute_h3_left_face_derivative(const Realf *const values, int k, Realf &fv_l, const int index, const int stride)
ARCH_DEV void compute_h8_left_face_value(const Realf *const values, int k, Realf &fv_l, const int index, const int stride)
ARCH_DEV void compute_h4_left_face_value(const Realf *const values, int k, Realf &fv_l, const int index, const int stride)
ARCH_DEV void compute_filtered_face_values(const Realf *const values, int k, face_estimate_order order, Realf &fv_l, Realf &fv_r, const Realf threshold, const int index, const int stride)
const int k
__global__ void const Realf const uint *__restrict__ const uint *__restrict__ const vmesh::GlobalID *__restrict__ const uint const uint const uint const Realf threshold
#define index(i, j, k)
Real slope_limiter(const Real &l, const Real &m, const Real &r)
static ARCH_HOSTDEV VecSimple< T > min(VecSimple< T > const &l, VecSimple< T > const &r)
static ARCH_HOSTDEV VecSimple< T > abs(const VecSimple< T > &l)