ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
ReactionKernelBulk_1_single_precision_CUDA.cu
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_1_single_precision_CUDA.cpp
17//! \\author pystencils
18//======================================================================================================================
19
20// kernel generated with pystencils v1.3.7+13.gdfd203a, lbmpy v1.3.7+10.gd3f6236, sympy v1.12.1, lbmpy_walberla/pystencils_walberla from waLBerla commit c69cb11d6a95d32b2280544d3d9abde1fe5fdbb5
21
22#include <cmath>
23
24#include "ReactionKernelBulk_1_single_precision_CUDA.h"
25#include "core/DataTypes.h"
26#include "core/Macros.h"
27
28#define FUNC_PREFIX __global__
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_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda {
49static FUNC_PREFIX __launch_bounds__(256) void reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda(float *RESTRICT _data_rho_0, 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, float order_0, float rate_coefficient, float stoech_0) {
50 if (blockDim.x * blockIdx.x + threadIdx.x < _size_rho_0_0 && blockDim.y * blockIdx.y + threadIdx.y < _size_rho_0_1 && blockDim.z * blockIdx.z + threadIdx.z < _size_rho_0_2) {
51 const int64_t ctr_0 = blockDim.x * blockIdx.x + threadIdx.x;
52 const int64_t ctr_1 = blockDim.y * blockIdx.y + threadIdx.y;
53 const int64_t ctr_2 = blockDim.z * blockIdx.z + threadIdx.z;
54 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];
55 const float rate_factor = rate_coefficient * powf(local_rho_0, order_0);
56 _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;
57 }
58}
59} // namespace internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda
60
61void ReactionKernelBulk_1_single_precision_CUDA::run(IBlock *block, gpuStream_t stream) {
62
63 auto rho_0 = block->getData<gpu::GPUField<float>>(rho_0ID);
64
65 auto &rate_coefficient = this->rate_coefficient_;
66 auto &stoech_0 = this->stoech_0_;
67 auto &order_0 = this->order_0_;
68 WALBERLA_ASSERT_GREATER_EQUAL(0, -int_c(rho_0->nrOfGhostLayers()))
69 float *RESTRICT _data_rho_0 = rho_0->dataAt(0, 0, 0, 0);
70 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->xSizeWithGhostLayer(), int64_t(int64_c(rho_0->xSize()) + 0))
71 const int64_t _size_rho_0_0 = int64_t(int64_c(rho_0->xSize()) + 0);
72 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->ySizeWithGhostLayer(), int64_t(int64_c(rho_0->ySize()) + 0))
73 const int64_t _size_rho_0_1 = int64_t(int64_c(rho_0->ySize()) + 0);
74 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->zSizeWithGhostLayer(), int64_t(int64_c(rho_0->zSize()) + 0))
75 const int64_t _size_rho_0_2 = int64_t(int64_c(rho_0->zSize()) + 0);
76 const int64_t _stride_rho_0_0 = int64_t(rho_0->xStride());
77 const int64_t _stride_rho_0_1 = int64_t(rho_0->yStride());
78 const int64_t _stride_rho_0_2 = int64_t(rho_0->zStride());
79 dim3 _block(uint32_c(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)), uint32_c(((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))), uint32_c(((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))))));
80 dim3 _grid(uint32_c(((_size_rho_0_0) % (((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)) == 0 ? (int64_t)(_size_rho_0_0) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)) : ((int64_t)(_size_rho_0_0) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))) + 1)), uint32_c(((_size_rho_0_1) % (((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))) == 0 ? (int64_t)(_size_rho_0_1) / (int64_t)(((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))) : ((int64_t)(_size_rho_0_1) / (int64_t)(((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) + 1)), uint32_c(((_size_rho_0_2) % (((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))))) == 0 ? (int64_t)(_size_rho_0_2) / (int64_t)(((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))))) : ((int64_t)(_size_rho_0_2) / (int64_t)(((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))))) + 1)));
81 internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda<<<_grid, _block, 0, stream>>>(_data_rho_0, _size_rho_0_0, _size_rho_0_1, _size_rho_0_2, _stride_rho_0_0, _stride_rho_0_1, _stride_rho_0_2, order_0, rate_coefficient, stoech_0);
82}
83
84void ReactionKernelBulk_1_single_precision_CUDA::runOnCellInterval(const shared_ptr<StructuredBlockStorage> &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block, gpuStream_t stream) {
85
86 CellInterval ci = globalCellInterval;
87 CellInterval blockBB = blocks->getBlockCellBB(*block);
88 blockBB.expand(ghostLayers);
89 ci.intersect(blockBB);
90 blocks->transformGlobalToBlockLocalCellInterval(ci, *block);
91 if (ci.empty())
92 return;
93
94 auto rho_0 = block->getData<gpu::GPUField<float>>(rho_0ID);
95
96 auto &rate_coefficient = this->rate_coefficient_;
97 auto &stoech_0 = this->stoech_0_;
98 auto &order_0 = this->order_0_;
99 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin(), -int_c(rho_0->nrOfGhostLayers()))
100 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin(), -int_c(rho_0->nrOfGhostLayers()))
101 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin(), -int_c(rho_0->nrOfGhostLayers()))
102 float *RESTRICT _data_rho_0 = rho_0->dataAt(ci.xMin(), ci.yMin(), ci.zMin(), 0);
103 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->xSizeWithGhostLayer(), int64_t(int64_c(ci.xSize()) + 0))
104 const int64_t _size_rho_0_0 = int64_t(int64_c(ci.xSize()) + 0);
105 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->ySizeWithGhostLayer(), int64_t(int64_c(ci.ySize()) + 0))
106 const int64_t _size_rho_0_1 = int64_t(int64_c(ci.ySize()) + 0);
107 WALBERLA_ASSERT_GREATER_EQUAL(rho_0->zSizeWithGhostLayer(), int64_t(int64_c(ci.zSize()) + 0))
108 const int64_t _size_rho_0_2 = int64_t(int64_c(ci.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 dim3 _block(uint32_c(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)), uint32_c(((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))), uint32_c(((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))))));
113 dim3 _grid(uint32_c(((_size_rho_0_0) % (((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)) == 0 ? (int64_t)(_size_rho_0_0) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)) : ((int64_t)(_size_rho_0_0) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))) + 1)), uint32_c(((_size_rho_0_1) % (((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))) == 0 ? (int64_t)(_size_rho_0_1) / (int64_t)(((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))) : ((int64_t)(_size_rho_0_1) / (int64_t)(((1024 < ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))) ? 1024 : ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) + 1)), uint32_c(((_size_rho_0_2) % (((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))))) == 0 ? (int64_t)(_size_rho_0_2) / (int64_t)(((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))))) : ((int64_t)(_size_rho_0_2) / (int64_t)(((64 < ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))) ? 64 : ((_size_rho_0_2 < ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0))))))) ? _size_rho_0_2 : ((int64_t)(256) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0) * ((_size_rho_0_1 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))) ? _size_rho_0_1 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_rho_0_0) ? 128 : _size_rho_0_0)))))))))) + 1)));
114 internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda<<<_grid, _block, 0, stream>>>(_data_rho_0, _size_rho_0_0, _size_rho_0_1, _size_rho_0_2, _stride_rho_0_0, _stride_rho_0_1, _stride_rho_0_2, order_0, rate_coefficient, stoech_0);
115}
116
117} // namespace pystencils
118} // namespace walberla
119
120#if (defined WALBERLA_CXX_COMPILER_IS_GNU) || (defined WALBERLA_CXX_COMPILER_IS_CLANG)
121#pragma GCC diagnostic pop
122#endif
123
124#if (defined WALBERLA_CXX_COMPILER_IS_INTEL)
125#pragma warning pop
126#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:42
ReactionKernelBulk_1_single_precision_CUDA.h
walberla::pystencils::ReactionKernelBulk_1_single_precision_CUDA::run
void run(IBlock *block, gpuStream_t stream=nullptr)
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:61
walberla::pystencils::ReactionKernelBulk_1_single_precision_CUDA::runOnCellInterval
void runOnCellInterval(const shared_ptr< StructuredBlockStorage > &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block, gpuStream_t stream=nullptr)
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:84
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:176
std
STL namespace.
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::_size_rho_0_1
static FUNC_PREFIX int64_t const int64_t const _size_rho_0_1
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::rate_coefficient
static FUNC_PREFIX int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const float float rate_coefficient
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::_size_rho_0_0
static FUNC_PREFIX int64_t const _size_rho_0_0
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::_stride_rho_0_2
static FUNC_PREFIX int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_rho_0_2
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::order_0
static FUNC_PREFIX int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const float order_0
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::__launch_bounds__
static FUNC_PREFIX __launch_bounds__(256) void reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda(float *RESTRICT _data_rho_0
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::_size_rho_0_2
static FUNC_PREFIX int64_t const int64_t const int64_t const _size_rho_0_2
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::_stride_rho_0_0
static FUNC_PREFIX int64_t const int64_t const int64_t const int64_t const _stride_rho_0_0
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla::pystencils::internal_reactionkernelbulk_1_single_precision_cuda_reactionkernelbulk_1_single_precision_cuda::_stride_rho_0_1
static FUNC_PREFIX int64_t const int64_t const int64_t const int64_t const int64_t const _stride_rho_0_1
Definition ReactionKernelBulk_1_single_precision_CUDA.cu:49
walberla
\file PackInfoPdfDoublePrecision.cpp \author pystencils
Definition EKWalberla.hpp:36