ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
ReactionKernelBulk_5_single_precision.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 ReactionKernelBulk_5_single_precision.cpp
17//! \\author pystencils
18//======================================================================================================================
19
20// kernel generated with pystencils v1.4+1.ge851f4e, lbmpy v1.4+1.ge9efe34, sympy v1.12.1, lbmpy_walberla/pystencils_walberla from waLBerla commit 3247aa7395049ca5bfb69d34d55e45db19fa439c
21
22#include <cmath>
23
24#include "ReactionKernelBulk_5_single_precision.h"
25#include "core/DataTypes.h"
26#include "core/Macros.h"
27
28#define FUNC_PREFIX
29
30#if (defined WALBERLA_CXX_COMPILER_IS_GNU) || (defined WALBERLA_CXX_COMPILER_IS_CLANG)
31#pragma GCC diagnostic push
32#pragma GCC diagnostic ignored "-Wfloat-equal"
33#pragma GCC diagnostic ignored "-Wshadow"
34#pragma GCC diagnostic ignored "-Wconversion"
35#pragma GCC diagnostic ignored "-Wunused-variable"
36#endif
37
38#if (defined WALBERLA_CXX_COMPILER_IS_INTEL)
39#pragma warning push
40#pragma warning(disable : 1599)
41#endif
42
43using namespace std;
44
45namespace walberla {
46namespace pystencils {
47
48namespace internal_737d6904f7c65dcfc320d619f189641b {
49static FUNC_PREFIX void reactionkernelbulk_5_single_precision_reactionkernelbulk_5_single_precision(float *RESTRICT _data_rho_0, float *RESTRICT _data_rho_1, float *RESTRICT _data_rho_2, float *RESTRICT _data_rho_3, float *RESTRICT _data_rho_4, int64_t const _size_rho_0_0, int64_t const _size_rho_0_1, int64_t const _size_rho_0_2, int64_t const _stride_rho_0_0, int64_t const _stride_rho_0_1, int64_t const _stride_rho_0_2, int64_t const _stride_rho_1_0, int64_t const _stride_rho_1_1, int64_t const _stride_rho_1_2, int64_t const _stride_rho_2_0, int64_t const _stride_rho_2_1, int64_t const _stride_rho_2_2, int64_t const _stride_rho_3_0, int64_t const _stride_rho_3_1, int64_t const _stride_rho_3_2, int64_t const _stride_rho_4_0, int64_t const _stride_rho_4_1, int64_t const _stride_rho_4_2, float order_0, float order_1, float order_2, float order_3, float order_4, float rate_coefficient, float stoech_0, float stoech_1, float stoech_2, float stoech_3, float stoech_4) {
50#pragma omp parallel
51 {
52#pragma omp for schedule(static)
53 for (int64_t ctr_2 = 0; ctr_2 < _size_rho_0_2; ctr_2 += 1) {
54 for (int64_t ctr_1 = 0; ctr_1 < _size_rho_0_1; ctr_1 += 1) {
55 for (int64_t ctr_0 = 0; ctr_0 < _size_rho_0_0; ctr_0 += 1) {
56 const float local_rho_0 = _data_rho_0[_stride_rho_0_0 * ctr_0 + _stride_rho_0_1 * ctr_1 + _stride_rho_0_2 * ctr_2];
57 const float local_rho_1 = _data_rho_1[_stride_rho_1_0 * ctr_0 + _stride_rho_1_1 * ctr_1 + _stride_rho_1_2 * ctr_2];
58 const float local_rho_2 = _data_rho_2[_stride_rho_2_0 * ctr_0 + _stride_rho_2_1 * ctr_1 + _stride_rho_2_2 * ctr_2];
59 const float local_rho_3 = _data_rho_3[_stride_rho_3_0 * ctr_0 + _stride_rho_3_1 * ctr_1 + _stride_rho_3_2 * ctr_2];
60 const float local_rho_4 = _data_rho_4[_stride_rho_4_0 * ctr_0 + _stride_rho_4_1 * ctr_1 + _stride_rho_4_2 * ctr_2];
61 const float rate_factor = rate_coefficient * powf(local_rho_0, order_0) * powf(local_rho_1, order_1) * powf(local_rho_2, order_2) * powf(local_rho_3, order_3) * powf(local_rho_4, order_4);
62 _data_rho_0[_stride_rho_0_0 * ctr_0 + _stride_rho_0_1 * ctr_1 + _stride_rho_0_2 * ctr_2] = local_rho_0 + rate_factor * stoech_0;
63 _data_rho_1[_stride_rho_1_0 * ctr_0 + _stride_rho_1_1 * ctr_1 + _stride_rho_1_2 * ctr_2] = local_rho_1 + rate_factor * stoech_1;
64 _data_rho_2[_stride_rho_2_0 * ctr_0 + _stride_rho_2_1 * ctr_1 + _stride_rho_2_2 * ctr_2] = local_rho_2 + rate_factor * stoech_2;
65 _data_rho_3[_stride_rho_3_0 * ctr_0 + _stride_rho_3_1 * ctr_1 + _stride_rho_3_2 * ctr_2] = local_rho_3 + rate_factor * stoech_3;
66 _data_rho_4[_stride_rho_4_0 * ctr_0 + _stride_rho_4_1 * ctr_1 + _stride_rho_4_2 * ctr_2] = local_rho_4 + rate_factor * stoech_4;
67 }
68 }
69 }
70 }
71}
72} // namespace internal_737d6904f7c65dcfc320d619f189641b
73
74void ReactionKernelBulk_5_single_precision::run(IBlock *block) {
75
76 auto rho_1 = block->getData<field::GhostLayerField<float, 1>>(rho_1ID);
77 auto rho_4 = block->getData<field::GhostLayerField<float, 1>>(rho_4ID);
78 auto rho_0 = block->getData<field::GhostLayerField<float, 1>>(rho_0ID);
79 auto rho_3 = block->getData<field::GhostLayerField<float, 1>>(rho_3ID);
80 auto rho_2 = block->getData<field::GhostLayerField<float, 1>>(rho_2ID);
81
82 auto &stoech_0 = this->stoech_0_;
83 auto &order_4 = this->order_4_;
84 auto &stoech_4 = this->stoech_4_;
85 auto &order_3 = this->order_3_;
86 auto &stoech_2 = this->stoech_2_;
87 auto &stoech_3 = this->stoech_3_;
88 auto &order_2 = this->order_2_;
89 auto &rate_coefficient = this->rate_coefficient_;
90 auto &order_0 = this->order_0_;
91 auto &order_1 = this->order_1_;
92 auto &stoech_1 = this->stoech_1_;
93 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_0->nrOfGhostLayers()))
94 float *RESTRICT _data_rho_0 = rho_0->dataAt(0, 0, 0, 0);
95 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_1->nrOfGhostLayers()))
96 float *RESTRICT _data_rho_1 = rho_1->dataAt(0, 0, 0, 0);
97 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_2->nrOfGhostLayers()))
98 float *RESTRICT _data_rho_2 = rho_2->dataAt(0, 0, 0, 0);
99 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_3->nrOfGhostLayers()))
100 float *RESTRICT _data_rho_3 = rho_3->dataAt(0, 0, 0, 0);
101 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_4->nrOfGhostLayers()))
102 float *RESTRICT _data_rho_4 = rho_4->dataAt(0, 0, 0, 0);
103 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->xSizeWithGhostLayer(), int64_t(int64_c(rho_0->xSize()) + 0))
104 const int64_t _size_rho_0_0 = int64_t(int64_c(rho_0->xSize()) + 0);
105 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->ySizeWithGhostLayer(), int64_t(int64_c(rho_0->ySize()) + 0))
106 const int64_t _size_rho_0_1 = int64_t(int64_c(rho_0->ySize()) + 0);
107 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->zSizeWithGhostLayer(), int64_t(int64_c(rho_0->zSize()) + 0))
108 const int64_t _size_rho_0_2 = int64_t(int64_c(rho_0->zSize()) + 0);
109 const int64_t _stride_rho_0_0 = int64_t(rho_0->xStride());
110 const int64_t _stride_rho_0_1 = int64_t(rho_0->yStride());
111 const int64_t _stride_rho_0_2 = int64_t(rho_0->zStride());
112 const int64_t _stride_rho_1_0 = int64_t(rho_1->xStride());
113 const int64_t _stride_rho_1_1 = int64_t(rho_1->yStride());
114 const int64_t _stride_rho_1_2 = int64_t(rho_1->zStride());
115 const int64_t _stride_rho_2_0 = int64_t(rho_2->xStride());
116 const int64_t _stride_rho_2_1 = int64_t(rho_2->yStride());
117 const int64_t _stride_rho_2_2 = int64_t(rho_2->zStride());
118 const int64_t _stride_rho_3_0 = int64_t(rho_3->xStride());
119 const int64_t _stride_rho_3_1 = int64_t(rho_3->yStride());
120 const int64_t _stride_rho_3_2 = int64_t(rho_3->zStride());
121 const int64_t _stride_rho_4_0 = int64_t(rho_4->xStride());
122 const int64_t _stride_rho_4_1 = int64_t(rho_4->yStride());
123 const int64_t _stride_rho_4_2 = int64_t(rho_4->zStride());
124 internal_737d6904f7c65dcfc320d619f189641b::reactionkernelbulk_5_single_precision_reactionkernelbulk_5_single_precision(_data_rho_0, _data_rho_1, _data_rho_2, _data_rho_3, _data_rho_4, _size_rho_0_0, _size_rho_0_1, _size_rho_0_2, _stride_rho_0_0, _stride_rho_0_1, _stride_rho_0_2, _stride_rho_1_0, _stride_rho_1_1, _stride_rho_1_2, _stride_rho_2_0, _stride_rho_2_1, _stride_rho_2_2, _stride_rho_3_0, _stride_rho_3_1, _stride_rho_3_2, _stride_rho_4_0, _stride_rho_4_1, _stride_rho_4_2, order_0, order_1, order_2, order_3, order_4, rate_coefficient, stoech_0, stoech_1, stoech_2, stoech_3, stoech_4);
125}
126
127void ReactionKernelBulk_5_single_precision::runOnCellInterval(const shared_ptr<StructuredBlockStorage> &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block) {
128
129 CellInterval ci = globalCellInterval;
130 CellInterval blockBB = blocks->getBlockCellBB(*block);
131 blockBB.expand(ghostLayers);
132 ci.intersect(blockBB);
133 blocks->transformGlobalToBlockLocalCellInterval(ci, *block);
134 if (ci.empty())
135 return;
136
137 auto rho_1 = block->getData<field::GhostLayerField<float, 1>>(rho_1ID);
138 auto rho_4 = block->getData<field::GhostLayerField<float, 1>>(rho_4ID);
139 auto rho_0 = block->getData<field::GhostLayerField<float, 1>>(rho_0ID);
140 auto rho_3 = block->getData<field::GhostLayerField<float, 1>>(rho_3ID);
141 auto rho_2 = block->getData<field::GhostLayerField<float, 1>>(rho_2ID);
142
143 auto &stoech_0 = this->stoech_0_;
144 auto &order_4 = this->order_4_;
145 auto &stoech_4 = this->stoech_4_;
146 auto &order_3 = this->order_3_;
147 auto &stoech_2 = this->stoech_2_;
148 auto &stoech_3 = this->stoech_3_;
149 auto &order_2 = this->order_2_;
150 auto &rate_coefficient = this->rate_coefficient_;
151 auto &order_0 = this->order_0_;
152 auto &order_1 = this->order_1_;
153 auto &stoech_1 = this->stoech_1_;
154 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_0->nrOfGhostLayers()))
155 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_0->nrOfGhostLayers()))
156 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_0->nrOfGhostLayers()))
157 float *RESTRICT _data_rho_0 = rho_0->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
158 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_1->nrOfGhostLayers()))
159 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_1->nrOfGhostLayers()))
160 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_1->nrOfGhostLayers()))
161 float *RESTRICT _data_rho_1 = rho_1->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
162 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_2->nrOfGhostLayers()))
163 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_2->nrOfGhostLayers()))
164 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_2->nrOfGhostLayers()))
165 float *RESTRICT _data_rho_2 = rho_2->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
166 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_3->nrOfGhostLayers()))
167 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_3->nrOfGhostLayers()))
168 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_3->nrOfGhostLayers()))
169 float *RESTRICT _data_rho_3 = rho_3->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
170 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_4->nrOfGhostLayers()))
171 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_4->nrOfGhostLayers()))
172 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_4->nrOfGhostLayers()))
173 float *RESTRICT _data_rho_4 = rho_4->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
174 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->xSizeWithGhostLayer(), int64_t(int64_c(ci.xSize()) + 0))
175 const int64_t _size_rho_0_0 = int64_t(int64_c(ci.xSize()) + 0);
176 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->ySizeWithGhostLayer(), int64_t(int64_c(ci.ySize()) + 0))
177 const int64_t _size_rho_0_1 = int64_t(int64_c(ci.ySize()) + 0);
178 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->zSizeWithGhostLayer(), int64_t(int64_c(ci.zSize()) + 0))
179 const int64_t _size_rho_0_2 = int64_t(int64_c(ci.zSize()) + 0);
180 const int64_t _stride_rho_0_0 = int64_t(rho_0->xStride());
181 const int64_t _stride_rho_0_1 = int64_t(rho_0->yStride());
182 const int64_t _stride_rho_0_2 = int64_t(rho_0->zStride());
183 const int64_t _stride_rho_1_0 = int64_t(rho_1->xStride());
184 const int64_t _stride_rho_1_1 = int64_t(rho_1->yStride());
185 const int64_t _stride_rho_1_2 = int64_t(rho_1->zStride());
186 const int64_t _stride_rho_2_0 = int64_t(rho_2->xStride());
187 const int64_t _stride_rho_2_1 = int64_t(rho_2->yStride());
188 const int64_t _stride_rho_2_2 = int64_t(rho_2->zStride());
189 const int64_t _stride_rho_3_0 = int64_t(rho_3->xStride());
190 const int64_t _stride_rho_3_1 = int64_t(rho_3->yStride());
191 const int64_t _stride_rho_3_2 = int64_t(rho_3->zStride());
192 const int64_t _stride_rho_4_0 = int64_t(rho_4->xStride());
193 const int64_t _stride_rho_4_1 = int64_t(rho_4->yStride());
194 const int64_t _stride_rho_4_2 = int64_t(rho_4->zStride());
195 internal_737d6904f7c65dcfc320d619f189641b::reactionkernelbulk_5_single_precision_reactionkernelbulk_5_single_precision(_data_rho_0, _data_rho_1, _data_rho_2, _data_rho_3, _data_rho_4, _size_rho_0_0, _size_rho_0_1, _size_rho_0_2, _stride_rho_0_0, _stride_rho_0_1, _stride_rho_0_2, _stride_rho_1_0, _stride_rho_1_1, _stride_rho_1_2, _stride_rho_2_0, _stride_rho_2_1, _stride_rho_2_2, _stride_rho_3_0, _stride_rho_3_1, _stride_rho_3_2, _stride_rho_4_0, _stride_rho_4_1, _stride_rho_4_2, order_0, order_1, order_2, order_3, order_4, rate_coefficient, stoech_0, stoech_1, stoech_2, stoech_3, stoech_4);
196}
197
198} // namespace pystencils
199} // namespace walberla
200
201#if (defined WALBERLA_CXX_COMPILER_IS_GNU) || (defined WALBERLA_CXX_COMPILER_IS_CLANG)
202#pragma GCC diagnostic pop
203#endif
204
205#if (defined WALBERLA_CXX_COMPILER_IS_INTEL)
206#pragma warning pop
207#endif
FUNC_PREFIX
#define FUNC_PREFIX
\file AdvectiveFluxKernel_double_precision.cpp \author pystencils
Definition AdvectiveFluxKernel_double_precision.cpp:28
RESTRICT
#define RESTRICT
\file AdvectiveFluxKernel_double_precision.h \author pystencils
Definition AdvectiveFluxKernel_double_precision.h:40
ReactionKernelBulk_5_single_precision.h
walberla::pystencils::ReactionKernelBulk_5_single_precision::runOnCellInterval
void runOnCellInterval(const shared_ptr< StructuredBlockStorage > &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block)
Definition ReactionKernelBulk_5_single_precision.cpp:127
walberla::pystencils::ReactionKernelBulk_5_single_precision::run
void run(IBlock *block)
Definition ReactionKernelBulk_5_single_precision.cpp:74
stream
cudaStream_t stream[1]
CUDA streams for parallel computing on CPU and GPU.
Definition common_cuda.cu:34
block
static double * block(double *p, std::size_t index, std::size_t size)
Definition elc.cpp:175
std
STL namespace.
walberla::pystencils::internal_737d6904f7c65dcfc320d619f189641b::reactionkernelbulk_5_single_precision_reactionkernelbulk_5_single_precision
static FUNC_PREFIX void reactionkernelbulk_5_single_precision_reactionkernelbulk_5_single_precision(float *RESTRICT _data_rho_0, float *RESTRICT _data_rho_1, float *RESTRICT _data_rho_2, float *RESTRICT _data_rho_3, float *RESTRICT _data_rho_4, int64_t const _size_rho_0_0, int64_t const _size_rho_0_1, int64_t const _size_rho_0_2, int64_t const _stride_rho_0_0, int64_t const _stride_rho_0_1, int64_t const _stride_rho_0_2, int64_t const _stride_rho_1_0, int64_t const _stride_rho_1_1, int64_t const _stride_rho_1_2, int64_t const _stride_rho_2_0, int64_t const _stride_rho_2_1, int64_t const _stride_rho_2_2, int64_t const _stride_rho_3_0, int64_t const _stride_rho_3_1, int64_t const _stride_rho_3_2, int64_t const _stride_rho_4_0, int64_t const _stride_rho_4_1, int64_t const _stride_rho_4_2, float order_0, float order_1, float order_2, float order_3, float order_4, float rate_coefficient, float stoech_0, float stoech_1, float stoech_2, float stoech_3, float stoech_4)
Definition ReactionKernelBulk_5_single_precision.cpp:49
walberla
\file PackInfoPdfDoublePrecision.cpp \author pystencils
Definition EKWalberla.hpp:38