ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
ReactionKernelBulk_3_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_3_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_3_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_54fb5dcfe8687c5ab8c7b22acb6d285f {
49static FUNC_PREFIX void reactionkernelbulk_3_single_precision_reactionkernelbulk_3_single_precision(float *RESTRICT _data_rho_0, float *RESTRICT _data_rho_1, float *RESTRICT _data_rho_2, 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, float order_0, float order_1, float order_2, float rate_coefficient, float stoech_0, float stoech_1, float stoech_2) {
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 rate_factor = rate_coefficient * powf(local_rho_0, order_0) * powf(local_rho_1, order_1) * powf(local_rho_2, order_2);
60 _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;
61 _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;
62 _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;
63 }
64 }
65 }
66 }
67}
68} // namespace internal_54fb5dcfe8687c5ab8c7b22acb6d285f
69
70void ReactionKernelBulk_3_single_precision::run(IBlock *block) {
71
72 auto rho_1 = block->getData<field::GhostLayerField<float, 1>>(rho_1ID);
73 auto rho_0 = block->getData<field::GhostLayerField<float, 1>>(rho_0ID);
74 auto rho_2 = block->getData<field::GhostLayerField<float, 1>>(rho_2ID);
75
76 auto &stoech_0 = this->stoech_0_;
77 auto &stoech_2 = this->stoech_2_;
78 auto &rate_coefficient = this->rate_coefficient_;
79 auto &order_2 = this->order_2_;
80 auto &order_0 = this->order_0_;
81 auto &order_1 = this->order_1_;
82 auto &stoech_1 = this->stoech_1_;
83 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_0->nrOfGhostLayers()))
84 float *RESTRICT _data_rho_0 = rho_0->dataAt(0, 0, 0, 0);
85 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_1->nrOfGhostLayers()))
86 float *RESTRICT _data_rho_1 = rho_1->dataAt(0, 0, 0, 0);
87 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_2->nrOfGhostLayers()))
88 float *RESTRICT _data_rho_2 = rho_2->dataAt(0, 0, 0, 0);
89 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->xSizeWithGhostLayer(), int64_t(int64_c(rho_0->xSize()) + 0))
90 const int64_t _size_rho_0_0 = int64_t(int64_c(rho_0->xSize()) + 0);
91 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->ySizeWithGhostLayer(), int64_t(int64_c(rho_0->ySize()) + 0))
92 const int64_t _size_rho_0_1 = int64_t(int64_c(rho_0->ySize()) + 0);
93 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->zSizeWithGhostLayer(), int64_t(int64_c(rho_0->zSize()) + 0))
94 const int64_t _size_rho_0_2 = int64_t(int64_c(rho_0->zSize()) + 0);
95 const int64_t _stride_rho_0_0 = int64_t(rho_0->xStride());
96 const int64_t _stride_rho_0_1 = int64_t(rho_0->yStride());
97 const int64_t _stride_rho_0_2 = int64_t(rho_0->zStride());
98 const int64_t _stride_rho_1_0 = int64_t(rho_1->xStride());
99 const int64_t _stride_rho_1_1 = int64_t(rho_1->yStride());
100 const int64_t _stride_rho_1_2 = int64_t(rho_1->zStride());
101 const int64_t _stride_rho_2_0 = int64_t(rho_2->xStride());
102 const int64_t _stride_rho_2_1 = int64_t(rho_2->yStride());
103 const int64_t _stride_rho_2_2 = int64_t(rho_2->zStride());
104 internal_54fb5dcfe8687c5ab8c7b22acb6d285f::reactionkernelbulk_3_single_precision_reactionkernelbulk_3_single_precision(_data_rho_0, _data_rho_1, _data_rho_2, _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, order_0, order_1, order_2, rate_coefficient, stoech_0, stoech_1, stoech_2);
105}
106
107void ReactionKernelBulk_3_single_precision::runOnCellInterval(const shared_ptr<StructuredBlockStorage> &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block) {
108
109 CellInterval ci = globalCellInterval;
110 CellInterval blockBB = blocks->getBlockCellBB(*block);
111 blockBB.expand(ghostLayers);
112 ci.intersect(blockBB);
113 blocks->transformGlobalToBlockLocalCellInterval(ci, *block);
114 if (ci.empty())
115 return;
116
117 auto rho_1 = block->getData<field::GhostLayerField<float, 1>>(rho_1ID);
118 auto rho_0 = block->getData<field::GhostLayerField<float, 1>>(rho_0ID);
119 auto rho_2 = block->getData<field::GhostLayerField<float, 1>>(rho_2ID);
120
121 auto &stoech_0 = this->stoech_0_;
122 auto &stoech_2 = this->stoech_2_;
123 auto &rate_coefficient = this->rate_coefficient_;
124 auto &order_2 = this->order_2_;
125 auto &order_0 = this->order_0_;
126 auto &order_1 = this->order_1_;
127 auto &stoech_1 = this->stoech_1_;
128 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_0->nrOfGhostLayers()))
129 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_0->nrOfGhostLayers()))
130 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_0->nrOfGhostLayers()))
131 float *RESTRICT _data_rho_0 = rho_0->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
132 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_1->nrOfGhostLayers()))
133 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_1->nrOfGhostLayers()))
134 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_1->nrOfGhostLayers()))
135 float *RESTRICT _data_rho_1 = rho_1->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
136 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_2->nrOfGhostLayers()))
137 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_2->nrOfGhostLayers()))
138 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_2->nrOfGhostLayers()))
139 float *RESTRICT _data_rho_2 = rho_2->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
140 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->xSizeWithGhostLayer(), int64_t(int64_c(ci.xSize()) + 0))
141 const int64_t _size_rho_0_0 = int64_t(int64_c(ci.xSize()) + 0);
142 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->ySizeWithGhostLayer(), int64_t(int64_c(ci.ySize()) + 0))
143 const int64_t _size_rho_0_1 = int64_t(int64_c(ci.ySize()) + 0);
144 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->zSizeWithGhostLayer(), int64_t(int64_c(ci.zSize()) + 0))
145 const int64_t _size_rho_0_2 = int64_t(int64_c(ci.zSize()) + 0);
146 const int64_t _stride_rho_0_0 = int64_t(rho_0->xStride());
147 const int64_t _stride_rho_0_1 = int64_t(rho_0->yStride());
148 const int64_t _stride_rho_0_2 = int64_t(rho_0->zStride());
149 const int64_t _stride_rho_1_0 = int64_t(rho_1->xStride());
150 const int64_t _stride_rho_1_1 = int64_t(rho_1->yStride());
151 const int64_t _stride_rho_1_2 = int64_t(rho_1->zStride());
152 const int64_t _stride_rho_2_0 = int64_t(rho_2->xStride());
153 const int64_t _stride_rho_2_1 = int64_t(rho_2->yStride());
154 const int64_t _stride_rho_2_2 = int64_t(rho_2->zStride());
155 internal_54fb5dcfe8687c5ab8c7b22acb6d285f::reactionkernelbulk_3_single_precision_reactionkernelbulk_3_single_precision(_data_rho_0, _data_rho_1, _data_rho_2, _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, order_0, order_1, order_2, rate_coefficient, stoech_0, stoech_1, stoech_2);
156}
157
158} // namespace pystencils
159} // namespace walberla
160
161#if (defined WALBERLA_CXX_COMPILER_IS_GNU) || (defined WALBERLA_CXX_COMPILER_IS_CLANG)
162#pragma GCC diagnostic pop
163#endif
164
165#if (defined WALBERLA_CXX_COMPILER_IS_INTEL)
166#pragma warning pop
167#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_3_single_precision.h
walberla::pystencils::ReactionKernelBulk_3_single_precision::run
void run(IBlock *block)
Definition ReactionKernelBulk_3_single_precision.cpp:70
walberla::pystencils::ReactionKernelBulk_3_single_precision::runOnCellInterval
void runOnCellInterval(const shared_ptr< StructuredBlockStorage > &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block)
Definition ReactionKernelBulk_3_single_precision.cpp:107
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_54fb5dcfe8687c5ab8c7b22acb6d285f::reactionkernelbulk_3_single_precision_reactionkernelbulk_3_single_precision
static FUNC_PREFIX void reactionkernelbulk_3_single_precision_reactionkernelbulk_3_single_precision(float *RESTRICT _data_rho_0, float *RESTRICT _data_rho_1, float *RESTRICT _data_rho_2, 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, float order_0, float order_1, float order_2, float rate_coefficient, float stoech_0, float stoech_1, float stoech_2)
Definition ReactionKernelBulk_3_single_precision.cpp:49
walberla
\file PackInfoPdfDoublePrecision.cpp \author pystencils
Definition EKWalberla.hpp:38