ESPResSo
Extensible Simulation Package for Research on Soft Matter Systems
Loading...
Searching...
No Matches
FrictionCouplingKernel_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 FrictionCouplingKernel_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 "FrictionCouplingKernel_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_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda {
49static FUNC_PREFIX __launch_bounds__(256) void frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda(float D, float *RESTRICT _data_f, float *RESTRICT const _data_j, int64_t const _size_f_0, int64_t const _size_f_1, int64_t const _size_f_2, int64_t const _stride_f_0, int64_t const _stride_f_1, int64_t const _stride_f_2, int64_t const _stride_f_3, int64_t const _stride_j_0, int64_t const _stride_j_1, int64_t const _stride_j_2, int64_t const _stride_j_3, float kT, float rho_lb) {
50 if (blockDim.x * blockIdx.x + threadIdx.x + 1 < _size_f_0 - 1 && blockDim.y * blockIdx.y + threadIdx.y + 1 < _size_f_1 - 1 && blockDim.z * blockIdx.z + threadIdx.z + 1 < _size_f_2 - 1) {
51 const int64_t ctr_0 = blockDim.x * blockIdx.x + threadIdx.x + 1;
52 const int64_t ctr_1 = blockDim.y * blockIdx.y + threadIdx.y + 1;
53 const int64_t ctr_2 = blockDim.z * blockIdx.z + threadIdx.z + 1;
54 _data_f[_stride_f_0 * ctr_0 + _stride_f_1 * ctr_1 + _stride_f_2 * ctr_2] = kT * (-_data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + 3 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 9 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 10 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 5 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 6 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 + 4 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 11 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 12 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 10 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 11 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 12 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 3 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 4 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 5 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 6 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 9 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2]) * 0.5f * ((1.0f) / (D)) * ((1.0f) / (rho_lb));
55 _data_f[_stride_f_0 * ctr_0 + _stride_f_1 * ctr_1 + _stride_f_2 * ctr_2 + _stride_f_3] = kT * (-_data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + 3 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 9 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 10 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 + 4 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 11 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 12 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 7 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 8 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 10 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 11 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 12 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 3 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 4 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 7 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 8 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 9 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + _stride_j_3]) * 0.5f * ((1.0f) / (D)) * ((1.0f) / (rho_lb));
56 _data_f[_stride_f_0 * ctr_0 + _stride_f_1 * ctr_1 + _stride_f_2 * ctr_2 + 2 * _stride_f_3] = kT * (-_data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 9 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 10 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 5 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 6 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 11 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_0 + _stride_j_1 * ctr_1 - _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 12 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 7 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_1 + _stride_j_2 * ctr_2 - _stride_j_2 + 8 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 10 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 11 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 12 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 2 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 5 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 6 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 7 * _stride_j_3] + _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 8 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + 9 * _stride_j_3] - _data_j[_stride_j_0 * ctr_0 + _stride_j_1 * ctr_1 + _stride_j_2 * ctr_2 + _stride_j_2 + 2 * _stride_j_3]) * 0.5f * ((1.0f) / (D)) * ((1.0f) / (rho_lb));
57 }
58}
59} // namespace internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda
60
61void FrictionCouplingKernel_single_precision_CUDA::run(IBlock *block, gpuStream_t stream) {
62
63 auto f = block->getData<gpu::GPUField<float>>(fID);
64 auto j = block->getData<gpu::GPUField<float>>(jID);
65
66 auto &rho_lb = this->rho_lb_;
67 auto &D = this->D_;
68 auto &kT = this->kT_;
69 WALBERLA_ASSERT_GREATER_EQUAL(-1, -int_c(f->nrOfGhostLayers()))
70 float *RESTRICT _data_f = f->dataAt(-1, -1, -1, 0);
71 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
72 WALBERLA_ASSERT_GREATER_EQUAL(-1, -int_c(j->nrOfGhostLayers()))
73 float *RESTRICT const _data_j = j->dataAt(-1, -1, -1, 0);
74 WALBERLA_ASSERT_EQUAL(j->layout(), field::fzyx)
75 WALBERLA_ASSERT_GREATER_EQUAL(f->xSizeWithGhostLayer(), int64_t(int64_c(f->xSize()) + 2))
76 const int64_t _size_f_0 = int64_t(int64_c(f->xSize()) + 2);
77 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
78 WALBERLA_ASSERT_GREATER_EQUAL(f->ySizeWithGhostLayer(), int64_t(int64_c(f->ySize()) + 2))
79 const int64_t _size_f_1 = int64_t(int64_c(f->ySize()) + 2);
80 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
81 WALBERLA_ASSERT_GREATER_EQUAL(f->zSizeWithGhostLayer(), int64_t(int64_c(f->zSize()) + 2))
82 const int64_t _size_f_2 = int64_t(int64_c(f->zSize()) + 2);
83 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
84 const int64_t _stride_f_0 = int64_t(f->xStride());
85 const int64_t _stride_f_1 = int64_t(f->yStride());
86 const int64_t _stride_f_2 = int64_t(f->zStride());
87 const int64_t _stride_f_3 = int64_t(1 * int64_t(f->fStride()));
88 const int64_t _stride_j_0 = int64_t(j->xStride());
89 const int64_t _stride_j_1 = int64_t(j->yStride());
90 const int64_t _stride_j_2 = int64_t(j->zStride());
91 const int64_t _stride_j_3 = int64_t(1 * int64_t(j->fStride()));
92 dim3 _block(uint32_c(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)), uint32_c(((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))), uint32_c(((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))))));
93 dim3 _grid(uint32_c(((_size_f_0 - 2) % (((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)) == 0 ? (int64_t)(_size_f_0 - 2) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)) : ((int64_t)(_size_f_0 - 2) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))) + 1)), uint32_c(((_size_f_1 - 2) % (((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))) == 0 ? (int64_t)(_size_f_1 - 2) / (int64_t)(((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))) : ((int64_t)(_size_f_1 - 2) / (int64_t)(((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) + 1)), uint32_c(((_size_f_2 - 2) % (((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))))) == 0 ? (int64_t)(_size_f_2 - 2) / (int64_t)(((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))))) : ((int64_t)(_size_f_2 - 2) / (int64_t)(((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))))) + 1)));
94 internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda<<<_grid, _block, 0, stream>>>(D, _data_f, _data_j, _size_f_0, _size_f_1, _size_f_2, _stride_f_0, _stride_f_1, _stride_f_2, _stride_f_3, _stride_j_0, _stride_j_1, _stride_j_2, _stride_j_3, kT, rho_lb);
95}
96
97void FrictionCouplingKernel_single_precision_CUDA::runOnCellInterval(const shared_ptr<StructuredBlockStorage> &blocks, const CellInterval &globalCellInterval, cell_idx_t ghostLayers, IBlock *block, gpuStream_t stream) {
98
99 CellInterval ci = globalCellInterval;
100 CellInterval blockBB = blocks->getBlockCellBB(*block);
101 blockBB.expand(ghostLayers);
102 ci.intersect(blockBB);
103 blocks->transformGlobalToBlockLocalCellInterval(ci, *block);
104 if (ci.empty())
105 return;
106
107 auto f = block->getData<gpu::GPUField<float>>(fID);
108 auto j = block->getData<gpu::GPUField<float>>(jID);
109
110 auto &rho_lb = this->rho_lb_;
111 auto &D = this->D_;
112 auto &kT = this->kT_;
113 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin() - 1, -int_c(f->nrOfGhostLayers()))
114 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin() - 1, -int_c(f->nrOfGhostLayers()))
115 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin() - 1, -int_c(f->nrOfGhostLayers()))
116 float *RESTRICT _data_f = f->dataAt(ci.xMin() - 1, ci.yMin() - 1, ci.zMin() - 1, 0);
117 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
118 WALBERLA_ASSERT_GREATER_EQUAL(ci.xMin() - 1, -int_c(j->nrOfGhostLayers()))
119 WALBERLA_ASSERT_GREATER_EQUAL(ci.yMin() - 1, -int_c(j->nrOfGhostLayers()))
120 WALBERLA_ASSERT_GREATER_EQUAL(ci.zMin() - 1, -int_c(j->nrOfGhostLayers()))
121 float *RESTRICT const _data_j = j->dataAt(ci.xMin() - 1, ci.yMin() - 1, ci.zMin() - 1, 0);
122 WALBERLA_ASSERT_EQUAL(j->layout(), field::fzyx)
123 WALBERLA_ASSERT_GREATER_EQUAL(f->xSizeWithGhostLayer(), int64_t(int64_c(ci.xSize()) + 2))
124 const int64_t _size_f_0 = int64_t(int64_c(ci.xSize()) + 2);
125 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
126 WALBERLA_ASSERT_GREATER_EQUAL(f->ySizeWithGhostLayer(), int64_t(int64_c(ci.ySize()) + 2))
127 const int64_t _size_f_1 = int64_t(int64_c(ci.ySize()) + 2);
128 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
129 WALBERLA_ASSERT_GREATER_EQUAL(f->zSizeWithGhostLayer(), int64_t(int64_c(ci.zSize()) + 2))
130 const int64_t _size_f_2 = int64_t(int64_c(ci.zSize()) + 2);
131 WALBERLA_ASSERT_EQUAL(f->layout(), field::fzyx)
132 const int64_t _stride_f_0 = int64_t(f->xStride());
133 const int64_t _stride_f_1 = int64_t(f->yStride());
134 const int64_t _stride_f_2 = int64_t(f->zStride());
135 const int64_t _stride_f_3 = int64_t(1 * int64_t(f->fStride()));
136 const int64_t _stride_j_0 = int64_t(j->xStride());
137 const int64_t _stride_j_1 = int64_t(j->yStride());
138 const int64_t _stride_j_2 = int64_t(j->zStride());
139 const int64_t _stride_j_3 = int64_t(1 * int64_t(j->fStride()));
140 dim3 _block(uint32_c(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)), uint32_c(((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))), uint32_c(((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))))));
141 dim3 _grid(uint32_c(((_size_f_0 - 2) % (((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)) == 0 ? (int64_t)(_size_f_0 - 2) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)) : ((int64_t)(_size_f_0 - 2) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))) + 1)), uint32_c(((_size_f_1 - 2) % (((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))) == 0 ? (int64_t)(_size_f_1 - 2) / (int64_t)(((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))) : ((int64_t)(_size_f_1 - 2) / (int64_t)(((1024 < ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))) ? 1024 : ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) + 1)), uint32_c(((_size_f_2 - 2) % (((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))))) == 0 ? (int64_t)(_size_f_2 - 2) / (int64_t)(((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))))) : ((int64_t)(_size_f_2 - 2) / (int64_t)(((64 < ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))) ? 64 : ((_size_f_2 - 2 < ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2))))))) ? _size_f_2 - 2 : ((int64_t)(256) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2) * ((_size_f_1 - 2 < 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))) ? _size_f_1 - 2 : 2 * ((int64_t)(128) / (int64_t)(((128 < _size_f_0 - 2) ? 128 : _size_f_0 - 2)))))))))) + 1)));
142 internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda<<<_grid, _block, 0, stream>>>(D, _data_f, _data_j, _size_f_0, _size_f_1, _size_f_2, _stride_f_0, _stride_f_1, _stride_f_2, _stride_f_3, _stride_j_0, _stride_j_1, _stride_j_2, _stride_j_3, kT, rho_lb);
143}
144
145} // namespace pystencils
146} // namespace walberla
147
148#if (defined WALBERLA_CXX_COMPILER_IS_GNU) || (defined WALBERLA_CXX_COMPILER_IS_CLANG)
149#pragma GCC diagnostic pop
150#endif
151
152#if (defined WALBERLA_CXX_COMPILER_IS_INTEL)
153#pragma warning pop
154#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
FrictionCouplingKernel_single_precision_CUDA.h
walberla::pystencils::FrictionCouplingKernel_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 FrictionCouplingKernel_single_precision_CUDA.cu:97
walberla::pystencils::FrictionCouplingKernel_single_precision_CUDA::run
void run(IBlock *block, gpuStream_t stream=nullptr)
Definition FrictionCouplingKernel_single_precision_CUDA.cu:61
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_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_size_f_1
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const _size_f_1
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_j_1
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_j_1
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_f_0
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const _stride_f_0
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_j_3
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_j_3
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_j_0
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_j_0
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_f_3
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_f_3
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_size_f_0
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const _size_f_0
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_j_2
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_j_2
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_f_1
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_f_1
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_data_j
static FUNC_PREFIX float *RESTRICT float *RESTRICT const _data_j
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_stride_f_2
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const _stride_f_2
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_size_f_2
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const _size_f_2
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::kT
static FUNC_PREFIX float *RESTRICT float *RESTRICT const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const int64_t const float kT
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::_data_f
static FUNC_PREFIX float *RESTRICT _data_f
Definition FrictionCouplingKernel_single_precision_CUDA.cu:49
walberla::pystencils::internal_frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda::__launch_bounds__
static FUNC_PREFIX __launch_bounds__(256) void frictioncouplingkernel_single_precision_cuda_frictioncouplingkernel_single_precision_cuda(float D
walberla
\file PackInfoPdfDoublePrecision.cpp \author pystencils
Definition EKWalberla.hpp:36