ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
CollideSweepSinglePrecisionLeesEdwards.cpp
Go to the documentation of this file.
1//======================================================================================================================
2//
3// This file is part of waLBerla. waLBerla is free software: you can
4// redistribute it and/or modify it under the terms of the GNU General Public
5// License as published by the Free Software Foundation, either version 3 of
6// the License, or (at your option) any later version.
7//
8// waLBerla is distributed in the hope that it will be useful, but WITHOUT
9// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
10// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
11// for more details.
12//
13// You should have received a copy of the GNU General Public License along
14// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
15//
16//! \\file CollideSweepSinglePrecisionLeesEdwards.cpp
17//! \\ingroup lbm
18//! \\author lbmpy
19//======================================================================================================================
20
21// kernel generated with pystencils v1.2, lbmpy v1.2, lbmpy_walberla/pystencils_walberla from waLBerla commit 4d10e7f2358fc4a4f7e99195d0f67f0b759ecb6f
22
23#include <cmath>
24
26#include "core/DataTypes.h"
27#include "core/Macros.h"
28
29#define FUNC_PREFIX
30
31#if (defined WALBERLA_CXX_COMPILER_IS_GNU) || (defined WALBERLA_CXX_COMPILER_IS_CLANG)
32#pragma GCC diagnostic push
33#pragma GCC diagnostic ignored "-Wfloat-equal"
34#pragma GCC diagnostic ignored "-Wshadow"
35#pragma GCC diagnostic ignored "-Wconversion"
36#pragma GCC diagnostic ignored "-Wunused-variable"
37#endif
38
39#if (defined WALBERLA_CXX_COMPILER_IS_INTEL)
40#pragma warning push
41#pragma warning(disable : 1599)
42#endif
43
44using namespace std;
45
46namespace walberla {
47namespace pystencils {
48
49namespace internal_ab1f3bc3368574afb482da84ccb58898 {
50static FUNC_PREFIX void collidesweepsingleprecisionleesedwards_collidesweepsingleprecisionleesedwards(float *RESTRICT const _data_force, float *RESTRICT _data_pdfs, int64_t const _size_force_0, int64_t const _size_force_1, int64_t const _size_force_2, int64_t const _stride_force_0, int64_t const _stride_force_1, int64_t const _stride_force_2, int64_t const _stride_force_3, int64_t const _stride_pdfs_0, int64_t const _stride_pdfs_1, int64_t const _stride_pdfs_2, int64_t const _stride_pdfs_3, float grid_size, float omega_shear, float v_s) {
51 const float xi_0 = ((1.0f) / (omega_shear * -0.25f + 2.0f));
52 const float rr_0 = xi_0 * (omega_shear * -2.0f + 4.0f);
53 for (int64_t ctr_2 = 0; ctr_2 < _size_force_2; ctr_2 += 1) {
54 float *RESTRICT _data_pdfs_20_36 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 6 * _stride_pdfs_3;
55 float *RESTRICT _data_force_20_32 = _data_force + _stride_force_2 * ctr_2 + 2 * _stride_force_3;
56 float *RESTRICT _data_pdfs_20_31 = _data_pdfs + _stride_pdfs_2 * ctr_2 + _stride_pdfs_3;
57 float *RESTRICT _data_pdfs_20_32 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 2 * _stride_pdfs_3;
58 float *RESTRICT _data_pdfs_20_311 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 11 * _stride_pdfs_3;
59 float *RESTRICT _data_pdfs_20_318 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 18 * _stride_pdfs_3;
60 float *RESTRICT _data_pdfs_20_313 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 13 * _stride_pdfs_3;
61 float *RESTRICT _data_pdfs_20_34 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 4 * _stride_pdfs_3;
62 float *RESTRICT _data_pdfs_20_33 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 3 * _stride_pdfs_3;
63 float *RESTRICT _data_pdfs_20_317 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 17 * _stride_pdfs_3;
64 float *RESTRICT _data_force_20_30 = _data_force + _stride_force_2 * ctr_2;
65 float *RESTRICT _data_pdfs_20_35 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 5 * _stride_pdfs_3;
66 float *RESTRICT _data_pdfs_20_314 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 14 * _stride_pdfs_3;
67 float *RESTRICT _data_pdfs_20_38 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 8 * _stride_pdfs_3;
68 float *RESTRICT _data_pdfs_20_312 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 12 * _stride_pdfs_3;
69 float *RESTRICT _data_pdfs_20_316 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 16 * _stride_pdfs_3;
70 float *RESTRICT _data_pdfs_20_39 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 9 * _stride_pdfs_3;
71 float *RESTRICT _data_pdfs_20_315 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 15 * _stride_pdfs_3;
72 float *RESTRICT _data_force_20_31 = _data_force + _stride_force_2 * ctr_2 + _stride_force_3;
73 float *RESTRICT _data_pdfs_20_310 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 10 * _stride_pdfs_3;
74 float *RESTRICT _data_pdfs_20_30 = _data_pdfs + _stride_pdfs_2 * ctr_2;
75 float *RESTRICT _data_pdfs_20_37 = _data_pdfs + _stride_pdfs_2 * ctr_2 + 7 * _stride_pdfs_3;
76 for (int64_t ctr_1 = 0; ctr_1 < _size_force_1; ctr_1 += 1) {
77 float *RESTRICT _data_pdfs_20_36_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_36;
78 float *RESTRICT _data_force_20_32_10 = _stride_force_1 * ctr_1 + _data_force_20_32;
79 float *RESTRICT _data_pdfs_20_31_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_31;
80 float *RESTRICT _data_pdfs_20_32_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_32;
81 float *RESTRICT _data_pdfs_20_311_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_311;
82 float *RESTRICT _data_pdfs_20_318_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_318;
83 float *RESTRICT _data_pdfs_20_313_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_313;
84 float *RESTRICT _data_pdfs_20_34_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_34;
85 float *RESTRICT _data_pdfs_20_33_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_33;
86 float *RESTRICT _data_pdfs_20_317_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_317;
87 float *RESTRICT _data_force_20_30_10 = _stride_force_1 * ctr_1 + _data_force_20_30;
88 float *RESTRICT _data_pdfs_20_35_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_35;
89 float *RESTRICT _data_pdfs_20_314_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_314;
90 float *RESTRICT _data_pdfs_20_38_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_38;
91 float *RESTRICT _data_pdfs_20_312_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_312;
92 float *RESTRICT _data_pdfs_20_316_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_316;
93 float *RESTRICT _data_pdfs_20_39_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_39;
94 float *RESTRICT _data_pdfs_20_315_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_315;
95 float *RESTRICT _data_force_20_31_10 = _stride_force_1 * ctr_1 + _data_force_20_31;
96 float *RESTRICT _data_pdfs_20_310_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_310;
97 float *RESTRICT _data_pdfs_20_30_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_30;
98 float *RESTRICT _data_pdfs_20_37_10 = _stride_pdfs_1 * ctr_1 + _data_pdfs_20_37;
99 for (int64_t ctr_0 = 0; ctr_0 < _size_force_0; ctr_0 += 1) {
100 const float xi_25 = _data_pdfs_20_36_10[_stride_pdfs_0 * ctr_0];
101 const float xi_26 = _data_force_20_32_10[_stride_force_0 * ctr_0];
102 const float xi_27 = _data_pdfs_20_31_10[_stride_pdfs_0 * ctr_0];
103 const float xi_28 = _data_pdfs_20_32_10[_stride_pdfs_0 * ctr_0];
104 const float xi_29 = _data_pdfs_20_311_10[_stride_pdfs_0 * ctr_0];
105 const float xi_30 = _data_pdfs_20_318_10[_stride_pdfs_0 * ctr_0];
106 const float xi_31 = _data_pdfs_20_313_10[_stride_pdfs_0 * ctr_0];
107 const float xi_32 = _data_pdfs_20_34_10[_stride_pdfs_0 * ctr_0];
108 const float xi_33 = _data_pdfs_20_33_10[_stride_pdfs_0 * ctr_0];
109 const float xi_34 = _data_pdfs_20_317_10[_stride_pdfs_0 * ctr_0];
110 const float xi_35 = _data_force_20_30_10[_stride_force_0 * ctr_0];
111 const float xi_36 = _data_pdfs_20_35_10[_stride_pdfs_0 * ctr_0];
112 const float xi_37 = _data_pdfs_20_314_10[_stride_pdfs_0 * ctr_0];
113 const float xi_38 = _data_pdfs_20_38_10[_stride_pdfs_0 * ctr_0];
114 const float xi_39 = _data_pdfs_20_312_10[_stride_pdfs_0 * ctr_0];
115 const float xi_40 = _data_pdfs_20_316_10[_stride_pdfs_0 * ctr_0];
116 const float xi_41 = _data_pdfs_20_39_10[_stride_pdfs_0 * ctr_0];
117 const float xi_42 = _data_pdfs_20_315_10[_stride_pdfs_0 * ctr_0];
118 const float xi_43 = _data_force_20_31_10[_stride_force_0 * ctr_0];
119 const float xi_44 = _data_pdfs_20_310_10[_stride_pdfs_0 * ctr_0];
120 const float xi_45 = _data_pdfs_20_30_10[_stride_pdfs_0 * ctr_0];
121 const float xi_46 = _data_pdfs_20_37_10[_stride_pdfs_0 * ctr_0];
122 const float xi_3 = xi_25;
123 const float xi_4 = xi_26;
124 const float xi_5 = xi_27;
125 const float xi_6 = xi_28;
126 const float xi_7 = xi_29;
127 const float xi_8 = xi_30;
128 const float xi_9 = xi_31;
129 const float xi_10 = xi_45;
130 const float xi_11 = xi_32;
131 const float xi_12 = xi_33;
132 const float xi_13 = xi_34;
133 const float xi_14 = xi_35;
134 const float xi_15 = xi_36;
135 const float xi_16 = xi_37;
136 const float xi_17 = xi_38;
137 const float xi_18 = xi_39;
138 const float xi_19 = xi_40;
139 const float xi_20 = xi_42;
140 const float xi_21 = xi_43;
141 const float xi_22 = xi_44;
142 const float xi_23 = xi_41;
143 const float xi_24 = xi_46;
144 const float vel0Term = xi_11 + xi_16 + xi_17 + xi_22 + xi_8;
145 const float vel1Term = xi_20 + xi_24 + xi_5 + xi_7;
146 const float vel2Term = xi_15 + xi_18 + xi_9;
147 const float rho = vel0Term + vel1Term + vel2Term + xi_10 + xi_12 + xi_13 + xi_19 + xi_23 + xi_3 + xi_6;
148 const float xi_1 = ((1.0f) / (rho));
149 const float u_0 = xi_1 * xi_14 * 0.5f + xi_1 * (vel0Term + xi_12 * -1.0f + xi_13 * -1.0f + xi_23 * -1.0f + xi_24 * -1.0f + xi_9 * -1.0f);
150 const float u_1 = xi_1 * xi_21 * 0.5f + xi_1 * (vel1Term + xi_17 + xi_18 * -1.0f + xi_19 * -1.0f + xi_22 * -1.0f + xi_23 * -1.0f + xi_6 * -1.0f);
151 const float u_2 = xi_1 * xi_4 * 0.5f + xi_1 * (vel2Term + xi_13 * -1.0f + xi_16 + xi_19 * -1.0f + xi_20 * -1.0f + xi_3 * -1.0f + xi_7 + xi_8 * -1.0f);
152 const float forceTerm_0 = omega_shear * u_0 * xi_14 * 0.5f + omega_shear * u_1 * xi_21 * 0.5f + omega_shear * u_2 * xi_4 * 0.5f + u_0 * xi_14 * -1.0f + u_1 * xi_21 * -1.0f + u_2 * xi_4 * -1.0f;
153 const float forceTerm_1 = omega_shear * u_0 * xi_14 * 0.083333333333333329f + omega_shear * u_1 * xi_21 * -0.16666666666666666f + omega_shear * u_2 * xi_4 * 0.083333333333333329f + rr_0 * xi_21 * -0.083333333333333329f + u_0 * xi_14 * -0.16666666666666666f + u_1 * xi_21 * 0.33333333333333331f + u_2 * xi_4 * -0.16666666666666666f + xi_21 * 0.16666666666666666f;
154 const float forceTerm_2 = omega_shear * u_0 * xi_14 * 0.083333333333333329f + omega_shear * u_1 * xi_21 * -0.16666666666666666f + omega_shear * u_2 * xi_4 * 0.083333333333333329f + rr_0 * xi_21 * 0.083333333333333329f + u_0 * xi_14 * -0.16666666666666666f + u_1 * xi_21 * 0.33333333333333331f + u_2 * xi_4 * -0.16666666666666666f + xi_21 * -0.16666666666666666f;
155 const float forceTerm_3 = omega_shear * u_0 * xi_14 * -0.16666666666666666f + omega_shear * u_1 * xi_21 * 0.083333333333333329f + omega_shear * u_2 * xi_4 * 0.083333333333333329f + rr_0 * xi_14 * 0.083333333333333329f + u_0 * xi_14 * 0.33333333333333331f + u_1 * xi_21 * -0.16666666666666666f + u_2 * xi_4 * -0.16666666666666666f + xi_14 * -0.16666666666666666f;
156 const float forceTerm_4 = omega_shear * u_0 * xi_14 * -0.16666666666666666f + omega_shear * u_1 * xi_21 * 0.083333333333333329f + omega_shear * u_2 * xi_4 * 0.083333333333333329f + rr_0 * xi_14 * -0.083333333333333329f + u_0 * xi_14 * 0.33333333333333331f + u_1 * xi_21 * -0.16666666666666666f + u_2 * xi_4 * -0.16666666666666666f + xi_14 * 0.16666666666666666f;
157 const float forceTerm_5 = omega_shear * u_0 * xi_14 * 0.083333333333333329f + omega_shear * u_1 * xi_21 * 0.083333333333333329f + omega_shear * u_2 * xi_4 * -0.16666666666666666f + rr_0 * xi_4 * -0.083333333333333329f + u_0 * xi_14 * -0.16666666666666666f + u_1 * xi_21 * -0.16666666666666666f + u_2 * xi_4 * 0.33333333333333331f + xi_4 * 0.16666666666666666f;
158 const float forceTerm_6 = omega_shear * u_0 * xi_14 * 0.083333333333333329f + omega_shear * u_1 * xi_21 * 0.083333333333333329f + omega_shear * u_2 * xi_4 * -0.16666666666666666f + rr_0 * xi_4 * 0.083333333333333329f + u_0 * xi_14 * -0.16666666666666666f + u_1 * xi_21 * -0.16666666666666666f + u_2 * xi_4 * 0.33333333333333331f + xi_4 * -0.16666666666666666f;
159 const float forceTerm_7 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_21 * 0.125f + omega_shear * u_1 * xi_14 * 0.125f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_2 * xi_4 * 0.041666666666666664f + rr_0 * xi_14 * 0.041666666666666664f + rr_0 * xi_21 * -0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_21 * -0.25f + u_1 * xi_14 * -0.25f + u_1 * xi_21 * 0.16666666666666666f + u_2 * xi_4 * -0.083333333333333329f + xi_14 * -0.083333333333333329f + xi_21 * 0.083333333333333329f;
160 const float forceTerm_8 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_21 * -0.125f + omega_shear * u_1 * xi_14 * -0.125f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_2 * xi_4 * 0.041666666666666664f + rr_0 * xi_14 * -0.041666666666666664f + rr_0 * xi_21 * -0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_21 * 0.25f + u_1 * xi_14 * 0.25f + u_1 * xi_21 * 0.16666666666666666f + u_2 * xi_4 * -0.083333333333333329f + xi_14 * 0.083333333333333329f + xi_21 * 0.083333333333333329f;
161 const float forceTerm_9 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_21 * -0.125f + omega_shear * u_1 * xi_14 * -0.125f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_2 * xi_4 * 0.041666666666666664f + rr_0 * xi_14 * 0.041666666666666664f + rr_0 * xi_21 * 0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_21 * 0.25f + u_1 * xi_14 * 0.25f + u_1 * xi_21 * 0.16666666666666666f + u_2 * xi_4 * -0.083333333333333329f + xi_14 * -0.083333333333333329f + xi_21 * -0.083333333333333329f;
162 const float forceTerm_10 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_21 * 0.125f + omega_shear * u_1 * xi_14 * 0.125f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_2 * xi_4 * 0.041666666666666664f + rr_0 * xi_14 * -0.041666666666666664f + rr_0 * xi_21 * 0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_21 * -0.25f + u_1 * xi_14 * -0.25f + u_1 * xi_21 * 0.16666666666666666f + u_2 * xi_4 * -0.083333333333333329f + xi_14 * 0.083333333333333329f + xi_21 * -0.083333333333333329f;
163 const float forceTerm_11 = omega_shear * u_0 * xi_14 * 0.041666666666666664f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_1 * xi_4 * -0.125f + omega_shear * u_2 * xi_21 * -0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_21 * -0.041666666666666664f + rr_0 * xi_4 * -0.041666666666666664f + u_0 * xi_14 * -0.083333333333333329f + u_1 * xi_21 * 0.16666666666666666f + u_1 * xi_4 * 0.25f + u_2 * xi_21 * 0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_21 * 0.083333333333333329f + xi_4 * 0.083333333333333329f;
164 const float forceTerm_12 = omega_shear * u_0 * xi_14 * 0.041666666666666664f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_1 * xi_4 * 0.125f + omega_shear * u_2 * xi_21 * 0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_21 * 0.041666666666666664f + rr_0 * xi_4 * -0.041666666666666664f + u_0 * xi_14 * -0.083333333333333329f + u_1 * xi_21 * 0.16666666666666666f + u_1 * xi_4 * -0.25f + u_2 * xi_21 * -0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_21 * -0.083333333333333329f + xi_4 * 0.083333333333333329f;
165 const float forceTerm_13 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_4 * 0.125f + omega_shear * u_1 * xi_21 * 0.041666666666666664f + omega_shear * u_2 * xi_14 * 0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_14 * 0.041666666666666664f + rr_0 * xi_4 * -0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_4 * -0.25f + u_1 * xi_21 * -0.083333333333333329f + u_2 * xi_14 * -0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_14 * -0.083333333333333329f + xi_4 * 0.083333333333333329f;
166 const float forceTerm_14 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_4 * -0.125f + omega_shear * u_1 * xi_21 * 0.041666666666666664f + omega_shear * u_2 * xi_14 * -0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_14 * -0.041666666666666664f + rr_0 * xi_4 * -0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_4 * 0.25f + u_1 * xi_21 * -0.083333333333333329f + u_2 * xi_14 * 0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_14 * 0.083333333333333329f + xi_4 * 0.083333333333333329f;
167 const float forceTerm_15 = omega_shear * u_0 * xi_14 * 0.041666666666666664f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_1 * xi_4 * 0.125f + omega_shear * u_2 * xi_21 * 0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_21 * -0.041666666666666664f + rr_0 * xi_4 * 0.041666666666666664f + u_0 * xi_14 * -0.083333333333333329f + u_1 * xi_21 * 0.16666666666666666f + u_1 * xi_4 * -0.25f + u_2 * xi_21 * -0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_21 * 0.083333333333333329f + xi_4 * -0.083333333333333329f;
168 const float forceTerm_16 = omega_shear * u_0 * xi_14 * 0.041666666666666664f + omega_shear * u_1 * xi_21 * -0.083333333333333329f + omega_shear * u_1 * xi_4 * -0.125f + omega_shear * u_2 * xi_21 * -0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_21 * 0.041666666666666664f + rr_0 * xi_4 * 0.041666666666666664f + u_0 * xi_14 * -0.083333333333333329f + u_1 * xi_21 * 0.16666666666666666f + u_1 * xi_4 * 0.25f + u_2 * xi_21 * 0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_21 * -0.083333333333333329f + xi_4 * -0.083333333333333329f;
169 const float forceTerm_17 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_4 * -0.125f + omega_shear * u_1 * xi_21 * 0.041666666666666664f + omega_shear * u_2 * xi_14 * -0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_14 * 0.041666666666666664f + rr_0 * xi_4 * 0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_4 * 0.25f + u_1 * xi_21 * -0.083333333333333329f + u_2 * xi_14 * 0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_14 * -0.083333333333333329f + xi_4 * -0.083333333333333329f;
170 const float forceTerm_18 = omega_shear * u_0 * xi_14 * -0.083333333333333329f + omega_shear * u_0 * xi_4 * 0.125f + omega_shear * u_1 * xi_21 * 0.041666666666666664f + omega_shear * u_2 * xi_14 * 0.125f + omega_shear * u_2 * xi_4 * -0.083333333333333329f + rr_0 * xi_14 * -0.041666666666666664f + rr_0 * xi_4 * 0.041666666666666664f + u_0 * xi_14 * 0.16666666666666666f + u_0 * xi_4 * -0.25f + u_1 * xi_21 * -0.083333333333333329f + u_2 * xi_14 * -0.25f + u_2 * xi_4 * 0.16666666666666666f + xi_14 * 0.083333333333333329f + xi_4 * -0.083333333333333329f;
171 const float u0Mu1 = u_0 + u_1 * -1.0f;
172 const float u0Pu1 = u_0 + u_1;
173 const float u1Pu2 = u_1 + u_2;
174 const float u1Mu2 = u_1 + u_2 * -1.0f;
175 const float u0Mu2 = u_0 + u_2 * -1.0f;
176 const float u0Pu2 = u_0 + u_2;
177 const float f_eq_common = rho * -1.0f * (u_0 * u_0) + rho * -1.0f * (u_1 * u_1) + rho * -1.0f * (u_2 * u_2) + rho;
178 _data_pdfs_20_30_10[_stride_pdfs_0 * ctr_0] = forceTerm_0 + omega_shear * (f_eq_common * 0.33333333333333331f + xi_10 * -1.0f) + xi_10;
179 _data_pdfs_20_31_10[_stride_pdfs_0 * ctr_0] = forceTerm_1 + omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_1 * u_1)) + xi_5 * -0.5f + xi_6 * -0.5f) + rr_0 * (rho * u_1 * 0.16666666666666666f + xi_5 * -0.5f + xi_6 * 0.5f) + xi_5 + ((-1.0f <= grid_size * -1.0f + ((float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + v_s) * 0.16666666666666666f) : (0.0f));
180 _data_pdfs_20_32_10[_stride_pdfs_0 * ctr_0] = forceTerm_2 + omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_1 * u_1)) + xi_5 * -0.5f + xi_6 * -0.5f) + rr_0 * (rho * u_1 * -0.16666666666666666f + xi_5 * 0.5f + xi_6 * -0.5f) + xi_6 + ((0.0f >= ((float)(ctr_1))) ? (rho * v_s * (u_0 * -2.0f + v_s) * 0.16666666666666666f) : (0.0f));
181 _data_pdfs_20_33_10[_stride_pdfs_0 * ctr_0] = forceTerm_3 + omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_0 * u_0)) + xi_11 * -0.5f + xi_12 * -0.5f) + rr_0 * (rho * u_0 * -0.16666666666666666f + xi_11 * 0.5f + xi_12 * -0.5f) + xi_12;
182 _data_pdfs_20_34_10[_stride_pdfs_0 * ctr_0] = forceTerm_4 + omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_0 * u_0)) + xi_11 * -0.5f + xi_12 * -0.5f) + rr_0 * (rho * u_0 * 0.16666666666666666f + xi_11 * -0.5f + xi_12 * 0.5f) + xi_11;
183 _data_pdfs_20_35_10[_stride_pdfs_0 * ctr_0] = forceTerm_5 + omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_2 * u_2)) + xi_15 * -0.5f + xi_3 * -0.5f) + rr_0 * (rho * u_2 * 0.16666666666666666f + xi_15 * -0.5f + xi_3 * 0.5f) + xi_15;
184 _data_pdfs_20_36_10[_stride_pdfs_0 * ctr_0] = forceTerm_6 + omega_shear * (f_eq_common * 0.16666666666666666f + rho * (-0.1111111111111111f + 0.33333333333333331f * (u_2 * u_2)) + xi_15 * -0.5f + xi_3 * -0.5f) + rr_0 * (rho * u_2 * -0.16666666666666666f + xi_15 * 0.5f + xi_3 * -0.5f) + xi_3;
185 _data_pdfs_20_37_10[_stride_pdfs_0 * ctr_0] = forceTerm_7 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Mu1 * u0Mu1)) + xi_22 * -0.5f + xi_24 * -0.5f) + rr_0 * (rho * u0Mu1 * -0.083333333333333329f + xi_22 * 0.5f + xi_24 * -0.5f) + xi_24 + ((-1.0f <= grid_size * -1.0f + ((float)(ctr_1))) ? (rho * v_s * (u_0 * -2.0f + u_1 * 3.0f + v_s * -1.0f + 1.0f) * 0.083333333333333329f) : (0.0f));
186 _data_pdfs_20_38_10[_stride_pdfs_0 * ctr_0] = forceTerm_8 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Pu1 * u0Pu1)) + xi_17 * -0.5f + xi_23 * -0.5f) + rr_0 * (rho * u0Pu1 * 0.083333333333333329f + xi_17 * -0.5f + xi_23 * 0.5f) + xi_17 + ((-1.0f <= grid_size * -1.0f + ((float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + u_1 * 3.0f + v_s + 1.0f) * -0.083333333333333329f) : (0.0f));
187 _data_pdfs_20_39_10[_stride_pdfs_0 * ctr_0] = forceTerm_9 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Pu1 * u0Pu1)) + xi_17 * -0.5f + xi_23 * -0.5f) + rr_0 * (rho * u0Pu1 * -0.083333333333333329f + xi_17 * 0.5f + xi_23 * -0.5f) + xi_23 + ((0.0f >= ((float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + u_1 * 3.0f + v_s * -1.0f - 1.0f) * 0.083333333333333329f) : (0.0f));
188 _data_pdfs_20_310_10[_stride_pdfs_0 * ctr_0] = forceTerm_10 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_2 * u_2) + 0.125f * (u0Mu1 * u0Mu1)) + xi_22 * -0.5f + xi_24 * -0.5f) + rr_0 * (rho * u0Mu1 * 0.083333333333333329f + xi_22 * -0.5f + xi_24 * 0.5f) + xi_22 + ((0.0f >= ((float)(ctr_1))) ? (rho * v_s * (u_0 * 2.0f + u_1 * -3.0f + v_s * -1.0f + 1.0f) * 0.083333333333333329f) : (0.0f));
189 _data_pdfs_20_311_10[_stride_pdfs_0 * ctr_0] = forceTerm_11 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Pu2 * u1Pu2)) + xi_19 * -0.5f + xi_7 * -0.5f) + rr_0 * (rho * u1Pu2 * 0.083333333333333329f + xi_19 * 0.5f + xi_7 * -0.5f) + xi_7;
190 _data_pdfs_20_312_10[_stride_pdfs_0 * ctr_0] = forceTerm_12 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Mu2 * u1Mu2)) + xi_18 * -0.5f + xi_20 * -0.5f) + rr_0 * (rho * u1Mu2 * -0.083333333333333329f + xi_18 * -0.5f + xi_20 * 0.5f) + xi_18;
191 _data_pdfs_20_313_10[_stride_pdfs_0 * ctr_0] = forceTerm_13 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Mu2 * u0Mu2)) + xi_8 * -0.5f + xi_9 * -0.5f) + rr_0 * (rho * u0Mu2 * -0.083333333333333329f + xi_8 * 0.5f + xi_9 * -0.5f) + xi_9;
192 _data_pdfs_20_314_10[_stride_pdfs_0 * ctr_0] = forceTerm_14 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Pu2 * u0Pu2)) + xi_13 * -0.5f + xi_16 * -0.5f) + rr_0 * (rho * u0Pu2 * 0.083333333333333329f + xi_13 * 0.5f + xi_16 * -0.5f) + xi_16;
193 _data_pdfs_20_315_10[_stride_pdfs_0 * ctr_0] = forceTerm_15 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Mu2 * u1Mu2)) + xi_18 * -0.5f + xi_20 * -0.5f) + rr_0 * (rho * u1Mu2 * 0.083333333333333329f + xi_18 * 0.5f + xi_20 * -0.5f) + xi_20;
194 _data_pdfs_20_316_10[_stride_pdfs_0 * ctr_0] = forceTerm_16 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_0 * u_0) + 0.125f * (u1Pu2 * u1Pu2)) + xi_19 * -0.5f + xi_7 * -0.5f) + rr_0 * (rho * u1Pu2 * -0.083333333333333329f + xi_19 * -0.5f + xi_7 * 0.5f) + xi_19;
195 _data_pdfs_20_317_10[_stride_pdfs_0 * ctr_0] = forceTerm_17 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Pu2 * u0Pu2)) + xi_13 * -0.5f + xi_16 * -0.5f) + rr_0 * (rho * u0Pu2 * -0.083333333333333329f + xi_13 * -0.5f + xi_16 * 0.5f) + xi_13;
196 _data_pdfs_20_318_10[_stride_pdfs_0 * ctr_0] = forceTerm_18 + omega_shear * (f_eq_common * 0.041666666666666664f + rho * (-0.013888888888888888f + 0.041666666666666664f * (u_1 * u_1) + 0.125f * (u0Mu2 * u0Mu2)) + xi_8 * -0.5f + xi_9 * -0.5f) + rr_0 * (rho * u0Mu2 * 0.083333333333333329f + xi_8 * -0.5f + xi_9 * 0.5f) + xi_8;
197 }
198 }
199 }
200}
201} // namespace internal_ab1f3bc3368574afb482da84ccb58898
202
204 auto pdfs = block->getData<field::GhostLayerField<float, 19>>(pdfsID);
205 auto force = block->getData<field::GhostLayerField<float, 3>>(forceID);
206
207 auto &omega_shear = this->omega_shear_;
208 auto &v_s = this->v_s_;
209 auto &grid_size = this->grid_size_;
210 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(force->nrOfGhostLayers()));
211 float *RESTRICT const _data_force = force->dataAt(0, 0, 0, 0);
212 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
213 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(pdfs->nrOfGhostLayers()));
214 float *RESTRICT _data_pdfs = pdfs->dataAt(0, 0, 0, 0);
215 WALBERLA_ASSERT_EQUAL(pdfs->layout(), field::fzyx);
216 WALBERLA_ASSERT_GREATER_EQUAL(force->xSizeWithGhostLayer(), int64_t(cell_idx_c(force->xSize()) + 0));
217 const int64_t _size_force_0 = int64_t(cell_idx_c(force->xSize()) + 0);
218 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
219 WALBERLA_ASSERT_GREATER_EQUAL(force->ySizeWithGhostLayer(), int64_t(cell_idx_c(force->ySize()) + 0));
220 const int64_t _size_force_1 = int64_t(cell_idx_c(force->ySize()) + 0);
221 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
222 WALBERLA_ASSERT_GREATER_EQUAL(force->zSizeWithGhostLayer(), int64_t(cell_idx_c(force->zSize()) + 0));
223 const int64_t _size_force_2 = int64_t(cell_idx_c(force->zSize()) + 0);
224 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
225 const int64_t _stride_force_0 = int64_t(force->xStride());
226 const int64_t _stride_force_1 = int64_t(force->yStride());
227 const int64_t _stride_force_2 = int64_t(force->zStride());
228 const int64_t _stride_force_3 = int64_t(1 * int64_t(force->fStride()));
229 const int64_t _stride_pdfs_0 = int64_t(pdfs->xStride());
230 const int64_t _stride_pdfs_1 = int64_t(pdfs->yStride());
231 const int64_t _stride_pdfs_2 = int64_t(pdfs->zStride());
232 const int64_t _stride_pdfs_3 = int64_t(1 * int64_t(pdfs->fStride()));
233 internal_ab1f3bc3368574afb482da84ccb58898::collidesweepsingleprecisionleesedwards_collidesweepsingleprecisionleesedwards(_data_force, _data_pdfs, _size_force_0, _size_force_1, _size_force_2, _stride_force_0, _stride_force_1, _stride_force_2, _stride_force_3, _stride_pdfs_0, _stride_pdfs_1, _stride_pdfs_2, _stride_pdfs_3, grid_size, omega_shear, v_s);
234}
235
236void CollideSweepSinglePrecisionLeesEdwards::runOnCellInterval(const shared_ptr<StructuredBlockStorage> &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block) {
237 CellInterval ci = globalCellInterval;
238 CellInterval blockBB = blocks->getBlockCellBB(*block);
239 blockBB.expand(ghostLayers);
240 ci.intersect(blockBB);
241 blocks->transformGlobalToBlockLocalCellInterval(ci, *block);
242 if (ci.empty())
243 return;
244
245 auto pdfs = block->getData<field::GhostLayerField<float, 19>>(pdfsID);
246 auto force = block->getData<field::GhostLayerField<float, 3>>(forceID);
247
248 auto &omega_shear = this->omega_shear_;
249 auto &v_s = this->v_s_;
250 auto &grid_size = this->grid_size_;
251 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(force->nrOfGhostLayers()));
252 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(force->nrOfGhostLayers()));
253 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(force->nrOfGhostLayers()));
254 float *RESTRICT const _data_force = force->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
255 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
256 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(pdfs->nrOfGhostLayers()));
257 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(pdfs->nrOfGhostLayers()));
258 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(pdfs->nrOfGhostLayers()));
259 float *RESTRICT _data_pdfs = pdfs->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
260 WALBERLA_ASSERT_EQUAL(pdfs->layout(), field::fzyx);
261 WALBERLA_ASSERT_GREATER_EQUAL(force->xSizeWithGhostLayer(), int64_t(cell_idx_c(ci.xSize()) + 0));
262 const int64_t _size_force_0 = int64_t(cell_idx_c(ci.xSize()) + 0);
263 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
264 WALBERLA_ASSERT_GREATER_EQUAL(force->ySizeWithGhostLayer(), int64_t(cell_idx_c(ci.ySize()) + 0));
265 const int64_t _size_force_1 = int64_t(cell_idx_c(ci.ySize()) + 0);
266 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
267 WALBERLA_ASSERT_GREATER_EQUAL(force->zSizeWithGhostLayer(), int64_t(cell_idx_c(ci.zSize()) + 0));
268 const int64_t _size_force_2 = int64_t(cell_idx_c(ci.zSize()) + 0);
269 WALBERLA_ASSERT_EQUAL(force->layout(), field::fzyx);
270 const int64_t _stride_force_0 = int64_t(force->xStride());
271 const int64_t _stride_force_1 = int64_t(force->yStride());
272 const int64_t _stride_force_2 = int64_t(force->zStride());
273 const int64_t _stride_force_3 = int64_t(1 * int64_t(force->fStride()));
274 const int64_t _stride_pdfs_0 = int64_t(pdfs->xStride());
275 const int64_t _stride_pdfs_1 = int64_t(pdfs->yStride());
276 const int64_t _stride_pdfs_2 = int64_t(pdfs->zStride());
277 const int64_t _stride_pdfs_3 = int64_t(1 * int64_t(pdfs->fStride()));
278 internal_ab1f3bc3368574afb482da84ccb58898::collidesweepsingleprecisionleesedwards_collidesweepsingleprecisionleesedwards(_data_force, _data_pdfs, _size_force_0, _size_force_1, _size_force_2, _stride_force_0, _stride_force_1, _stride_force_2, _stride_force_3, _stride_pdfs_0, _stride_pdfs_1, _stride_pdfs_2, _stride_pdfs_3, grid_size, omega_shear, v_s);
279}
280
281} // namespace pystencils
282} // namespace walberla
283
284#if (defined WALBERLA_CXX_COMPILER_IS_GNU) || (defined WALBERLA_CXX_COMPILER_IS_CLANG)
285#pragma GCC diagnostic pop
286#endif
287
288#if (defined WALBERLA_CXX_COMPILER_IS_INTEL)
289#pragma warning pop
290#endif
#define FUNC_PREFIX
\file AdvectiveFluxKernel_double_precision.cpp \ingroup lbm \author lbmpy
#define RESTRICT
\file AdvectiveFluxKernel_double_precision.h \author pystencils
__global__ float * force
void runOnCellInterval(const shared_ptr< StructuredBlockStorage > &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block)
static double * block(double *p, std::size_t index, std::size_t size)
Definition elc.cpp:174
static FUNC_PREFIX void collidesweepsingleprecisionleesedwards_collidesweepsingleprecisionleesedwards(float *RESTRICT const _data_force, float *RESTRICT _data_pdfs, int64_t const _size_force_0, int64_t const _size_force_1, int64_t const _size_force_2, int64_t const _stride_force_0, int64_t const _stride_force_1, int64_t const _stride_force_2, int64_t const _stride_force_3, int64_t const _stride_pdfs_0, int64_t const _stride_pdfs_1, int64_t const _stride_pdfs_2, int64_t const _stride_pdfs_3, float grid_size, float omega_shear, float v_s)