Commit 6245d049 authored by Dominik Thoennes's avatar Dominik Thoennes

Merge branch 'add-codegen-into-ci-rebase' into 'master'

Add codegen into CI

See merge request !259
parents 274e09cd 229a0ace
Pipeline #22683 passed with stages
in 504 minutes and 27 seconds
This diff is collapsed.
......@@ -560,7 +560,7 @@ endif ( )
#############################################################################################################################
if ( WALBERLA_BUILD_WITH_CODEGEN )
find_package( PythonInterp 3 QUIET REQUIRED)
execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "import pystencils_walberla; from pystencils.include import get_pystencils_include_path; print(get_pystencils_include_path())"
execute_process(COMMAND ${PYTHON_EXECUTABLE} -c "from pystencils.include import get_pystencils_include_path; print(get_pystencils_include_path())"
RESULT_VARIABLE PYTHON_RET_CODE
OUTPUT_VARIABLE PYSTENCILS_INCLUDE_PATH)
if(NOT PYTHON_RET_CODE EQUAL 0)
......
......@@ -155,7 +155,7 @@ int main( int argc, char **argv )
auto remainingTimeLoggerFrequency = parameters.getParameter< double >( "remainingTimeLoggerFrequency", -1.0 ); // in seconds
if (remainingTimeLoggerFrequency > 0) {
auto logger = timing::RemainingTimeLogger( timeLoop.getNrOfTimeSteps() * outerIterations, remainingTimeLoggerFrequency );
auto logger = timing::RemainingTimeLogger( timeLoop.getNrOfTimeSteps() * uint_c( outerIterations ), remainingTimeLoggerFrequency );
timeLoop.addFuncAfterTimeStep( logger, "remaining time logger" );
}
......
......@@ -152,23 +152,29 @@ with CodeGeneration() as ctx:
vec = { 'assume_aligned': True, 'assume_inner_stride_one': True}
# check if openmp is enabled in cmake
if ctx.openmp:
openmp_enabled = True
else:
openmp_enabled = False
# Sweeps
vec['nontemporal'] = opts['two_field_nt_stores']
generate_sweep(ctx, 'GenLbKernel', update_rule_two_field, field_swaps=[('pdfs', 'pdfs_tmp')],
cpu_vectorize_info=vec)
vec['nontemporal'] = opts['aa_even_nt_stores']
generate_sweep(ctx, 'GenLbKernelAAEven', update_rule_aa_even, cpu_vectorize_info=vec,
cpu_openmp=True, ghost_layers=1)
cpu_openmp=openmp_enabled, ghost_layers=1)
vec['nontemporal'] = opts['aa_odd_nt_stores']
generate_sweep(ctx, 'GenLbKernelAAOdd', update_rule_aa_odd, cpu_vectorize_info=vec,
cpu_openmp=True, ghost_layers=1)
cpu_openmp=openmp_enabled, ghost_layers=1)
setter_assignments = macroscopic_values_setter(update_rule_two_field.method, velocity=velocity_field.center_vector,
pdfs=pdfs.center_vector, density=1)
getter_assignments = macroscopic_values_getter(update_rule_two_field.method, velocity=velocity_field.center_vector,
pdfs=pdfs.center_vector, density=None)
generate_sweep(ctx, 'GenMacroSetter', setter_assignments, cpu_openmp=True)
generate_sweep(ctx, 'GenMacroGetter', getter_assignments, cpu_openmp=True)
generate_sweep(ctx, 'GenMacroSetter', setter_assignments, cpu_openmp=openmp_enabled)
generate_sweep(ctx, 'GenMacroGetter', getter_assignments, cpu_openmp=openmp_enabled)
# Communication
generate_pack_info_from_kernel(ctx, 'GenPackInfo', update_rule_two_field,
......
......@@ -64,10 +64,11 @@ function( handle_python_codegen sourceFilesOut generatedSourceFilesOut generator
string(REPLACE "\"" "\\\"" pythonParameters ${pythonParameters}) # even one more quoting level required
string(REPLACE "\n" "" pythonParameters ${pythonParameters}) # remove newline characters
set( WALBERLA_PYTHON_DIR ${walberla_SOURCE_DIR}/python)
file(MAKE_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/${codegenCfg}")
add_custom_command(OUTPUT ${generatedWithAbsolutePath}
DEPENDS ${sourceFile}
COMMAND ${PYTHON_EXECUTABLE} ${sourceFile} ${pythonParameters}
COMMAND ${CMAKE_COMMAND} -E env PYTHONPATH=${WALBERLA_PYTHON_DIR} ${PYTHON_EXECUTABLE} ${sourceFile} ${pythonParameters}
WORKING_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/${codegenCfg}")
endif()
else()
......
from .boundary import generate_boundary
from .walberla_lbm_generation import RefinementScaling, generate_lattice_model
__all__ = ['generate_lattice_model', 'RefinementScaling', 'generate_boundary']
import numpy as np
from jinja2 import Environment, PackageLoader, StrictUndefined
from lbmpy.boundaries.boundaryhandling import create_lattice_boltzmann_boundary_kernel
from lbmpy_walberla.walberla_lbm_generation import KernelInfo
from pystencils import Field, FieldType
from pystencils.boundaries.createindexlist import (
boundary_index_array_coordinate_names, direction_member_name,
numpy_data_type_for_boundary_object)
from pystencils.data_types import TypedSymbol, create_type
from pystencils_walberla.codegen import default_create_kernel_parameters
from pystencils_walberla.jinja_filters import add_pystencils_filters_to_jinja_env
def generate_boundary(generation_context, class_name, boundary_object, lb_method, **create_kernel_params):
struct_name = "IndexInfo"
boundary_object.name = class_name
create_kernel_params = default_create_kernel_parameters(generation_context, create_kernel_params)
target = create_kernel_params['target']
index_struct_dtype = numpy_data_type_for_boundary_object(boundary_object, lb_method.dim)
pdf_field = Field.create_generic('pdfs', lb_method.dim,
np.float64 if generation_context.double_accuracy else np.float32,
index_dimensions=1, layout='fzyx', index_shape=[len(lb_method.stencil)])
index_field = Field('indexVector', FieldType.INDEXED, index_struct_dtype, layout=[0],
shape=(TypedSymbol("indexVectorSize", create_type(np.int64)), 1), strides=(1, 1))
kernel = create_lattice_boltzmann_boundary_kernel(pdf_field, index_field, lb_method, boundary_object, target=target,
openmp=generation_context.openmp)
kernel.function_name = "boundary_" + boundary_object.name
# waLBerla is a 3D framework. Therefore, a zero for the z index has to be added if we work in 2D
if lb_method.dim == 2:
stencil = ()
for d in lb_method.stencil:
d = d + (0,)
stencil = stencil + (d,)
else:
stencil = lb_method.stencil
stencil_info = [(i, ", ".join([str(e) for e in d])) for i, d in enumerate(stencil)]
context = {
'class_name': boundary_object.name,
'StructName': struct_name,
'StructDeclaration': struct_from_numpy_dtype(struct_name, index_struct_dtype),
'kernel': KernelInfo(kernel),
'stencil_info': stencil_info,
'dim': lb_method.dim,
'target': target,
'namespace': 'lbm',
}
env = Environment(loader=PackageLoader('lbmpy_walberla'), undefined=StrictUndefined)
add_pystencils_filters_to_jinja_env(env)
header = env.get_template('Boundary.tmpl.h').render(**context)
source = env.get_template('Boundary.tmpl.cpp').render(**context)
source_extension = "cpp" if create_kernel_params.get("target", "cpu") == "cpu" else "cu"
generation_context.write_file("{}.h".format(class_name), header)
generation_context.write_file("{}.{}".format(class_name, source_extension), source)
def struct_from_numpy_dtype(struct_name, numpy_dtype):
result = "struct %s { \n" % (struct_name,)
equality_compare = []
constructor_params = []
constructor_initializer_list = []
for name, (sub_type, offset) in numpy_dtype.fields.items():
pystencils_type = create_type(sub_type)
result += " %s %s;\n" % (pystencils_type, name)
if name in boundary_index_array_coordinate_names or name == direction_member_name:
constructor_params.append("%s %s_" % (pystencils_type, name))
constructor_initializer_list.append("%s(%s_)" % (name, name))
else:
constructor_initializer_list.append("%s()" % name)
if pystencils_type.is_float():
equality_compare.append("floatIsEqual(%s, o.%s)" % (name, name))
else:
equality_compare.append("%s == o.%s" % (name, name))
result += " %s(%s) : %s {}\n" % \
(struct_name, ", ".join(constructor_params), ", ".join(constructor_initializer_list))
result += " bool operator==(const %s & o) const {\n return %s;\n }\n" % \
(struct_name, " && ".join(equality_compare))
result += "};\n"
return result
import numpy as np
from lbmpy.sparse import (
create_lb_update_rule_sparse, create_macroscopic_value_getter_sparse,
create_macroscopic_value_setter_sparse, create_symbolic_list)
from pystencils import TypedSymbol, create_kernel
class ListLbGenerator:
def __init__(self, collision_rule, pdf_type=np.float64, index_type=np.uint32, layout='SoA'):
self.collision_rule = collision_rule
num_cells = TypedSymbol('num_cells', np.uint64)
method = self.collision_rule.method
q = len(method.stencil)
self.num_cells = num_cells
self.src = create_symbolic_list('src', num_cells, q, pdf_type, layout)
self.dst = create_symbolic_list('dst', num_cells, q, pdf_type, layout)
self.idx = create_symbolic_list('idx', num_cells, q, index_type, layout)
def kernel(self, kernel_type='stream_pull_collide'):
ac = create_lb_update_rule_sparse(self.collision_rule, self.src, self.dst, self.idx, kernel_type)
return create_kernel(ac)
def getter_ast(self, density=True, velocity=True):
assert density or velocity
method = self.collision_rule.method
output_descriptor = {}
if density:
output_descriptor['density'] = create_symbolic_list('rho', self.src.spatial_shape[0], 1, self.src.dtype)
if velocity:
output_descriptor['velocity'] = create_symbolic_list('vel', self.src.spatial_shape[0],
method.dim, self.src.dtype)
ac = create_macroscopic_value_getter_sparse(method, self.src, output_descriptor)
return create_kernel(ac)
def setter_ast(self, density=True, velocity=True):
method = self.collision_rule.method
if density is True:
density = create_symbolic_list('rho', self.src.spatial_shape[0], 1, self.src.dtype).center
if velocity is True:
velocity = create_symbolic_list('vel', self.src.spatial_shape[0], method.dim, self.src.dtype).center_vector
ac = create_macroscopic_value_setter_sparse(method, self.src, density, velocity)
return create_kernel(ac)
//======================================================================================================================
//
// This file is part of waLBerla. waLBerla is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// waLBerla is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \\file {{class_name}}.cpp
//! \\ingroup lbm
//! \\author lbmpy
//======================================================================================================================
#include <cmath>
#include "core/DataTypes.h"
#include "core/Macros.h"
#include "{{class_name}}.h"
{% if target == 'gpu' -%}
#include "cuda/ErrorChecking.h"
{%- endif %}
{% if target == 'cpu' -%}
#define FUNC_PREFIX
{%- elif target == 'gpu' -%}
#define FUNC_PREFIX __global__
{%- endif %}
using namespace std;
namespace walberla {
namespace {{namespace}} {
#ifdef __GNUC__
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wstrict-aliasing"
#pragma GCC diagnostic ignored "-Wunused-variable"
#pragma GCC diagnostic ignored "-Wconversion"
#endif
#ifdef __CUDACC__
#pragma push
#pragma diag_suppress = declared_but_not_referenced
#endif
{{kernel|generate_definition}}
#ifdef __GNUC__
#pragma GCC diagnostic pop
#endif
#ifdef __CUDACC__
#pragma pop
#endif
void {{class_name}}::run( IBlock * block, IndexVectors::Type type {% if target == 'gpu'%}, cudaStream_t stream {%endif%})
{
auto * indexVectors = block->getData<IndexVectors>(indexVectorID);
int64_t indexVectorSize = int64_c( indexVectors->indexVector(type).size() );
if( indexVectorSize == 0)
return;
{% if target == 'gpu' -%}
auto pointer = indexVectors->pointerGpu(type);
{% else %}
auto pointer = indexVectors->pointerCpu(type);
{% endif %}
uint8_t * _data_indexVector = reinterpret_cast<uint8_t*>(pointer);
{{kernel|generate_block_data_to_field_extraction(['indexVector', 'indexVectorSize'])|indent(4)}}
{{kernel|generate_call(spatial_shape_symbols=['indexVectorSize'], stream='stream')|indent(4)}}
}
void {{class_name}}::operator() ( IBlock * block{% if target == 'gpu'%}, cudaStream_t stream {%endif%} )
{
run( block, IndexVectors::ALL{% if target == 'gpu'%}, stream {%endif%});
}
void {{class_name}}::inner( IBlock * block{% if target == 'gpu'%}, cudaStream_t stream {%endif%} )
{
run( block, IndexVectors::INNER{% if target == 'gpu'%}, stream {%endif%} );
}
void {{class_name}}::outer( IBlock * block{% if target == 'gpu'%}, cudaStream_t stream {%endif%} )
{
run( block, IndexVectors::OUTER{% if target == 'gpu'%}, stream {%endif%} );
}
} // namespace {{namespace}}
} // namespace walberla
//======================================================================================================================
//
// This file is part of waLBerla. waLBerla is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// waLBerla is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \\file {{class_name}}.h
//! \\author pystencils
//======================================================================================================================
#include "core/DataTypes.h"
{% if target is equalto 'cpu' -%}
#include "field/GhostLayerField.h"
{%- elif target is equalto 'gpu' -%}
#include "cuda/GPUField.h"
{%- endif %}
#include "domain_decomposition/BlockDataID.h"
#include "domain_decomposition/IBlock.h"
#include "blockforest/StructuredBlockForest.h"
#include "field/FlagField.h"
#include <set>
#include <vector>
#ifdef __GNUC__
#define RESTRICT __restrict__
#elif _MSC_VER
#define RESTRICT __restrict
#else
#define RESTRICT
#endif
namespace walberla {
namespace {{namespace}} {
class {{class_name}}
{
public:
{{StructDeclaration|indent(4)}}
class IndexVectors
{
public:
using CpuIndexVector = std::vector<{{StructName}}>;
enum Type {
ALL = 0,
INNER = 1,
OUTER = 2,
NUM_TYPES = 3
};
IndexVectors() : cpuVectors_(NUM_TYPES) {}
bool operator==(IndexVectors & other) { return other.cpuVectors_ == cpuVectors_; }
{% if target == 'gpu' -%}
~IndexVectors() {
for( auto & gpuVec: gpuVectors_)
cudaFree( gpuVec );
}
{% endif %}
CpuIndexVector & indexVector(Type t) { return cpuVectors_[t]; }
{{StructName}} * pointerCpu(Type t) { return &(cpuVectors_[t][0]); }
{% if target == 'gpu' -%}
{{StructName}} * pointerGpu(Type t) { return gpuVectors_[t]; }
{% endif %}
void syncGPU()
{
{% if target == 'gpu' -%}
gpuVectors_.resize( cpuVectors_.size() );
for(size_t i=0; i < size_t(NUM_TYPES); ++i )
{
auto & gpuVec = gpuVectors_[i];
auto & cpuVec = cpuVectors_[i];
cudaFree( gpuVec );
cudaMalloc( &gpuVec, sizeof({{StructName}}) * cpuVec.size() );
cudaMemcpy( gpuVec, &cpuVec[0], sizeof({{StructName}}) * cpuVec.size(), cudaMemcpyHostToDevice );
}
{%- endif %}
}
private:
std::vector<CpuIndexVector> cpuVectors_;
{% if target == 'gpu' -%}
using GpuIndexVector = {{StructName}} *;
std::vector<GpuIndexVector> gpuVectors_;
{% endif %}
};
{{class_name}}( const shared_ptr<StructuredBlockForest> & blocks,
{{kernel|generate_constructor_parameters(['indexVector', 'indexVectorSize'])}} )
: {{ kernel|generate_constructor_initializer_list(['indexVector', 'indexVectorSize']) }}
{
auto createIdxVector = []( IBlock * const , StructuredBlockStorage * const ) { return new IndexVectors(); };
indexVectorID = blocks->addStructuredBlockData< IndexVectors >( createIdxVector, "IndexField_{{class_name}}");
};
void operator() ( IBlock * block {% if target == 'gpu'%}, cudaStream_t stream = 0 {%endif%});
void inner( IBlock * block {% if target == 'gpu'%}, cudaStream_t stream = 0 {%endif%});
void outer( IBlock * block {% if target == 'gpu'%}, cudaStream_t stream = 0 {%endif%});
template<typename FlagField_T>
void fillFromFlagField( const shared_ptr<StructuredBlockForest> & blocks, ConstBlockDataID flagFieldID,
FlagUID boundaryFlagUID, FlagUID domainFlagUID)
{
for( auto blockIt = blocks->begin(); blockIt != blocks->end(); ++blockIt )
fillFromFlagField<FlagField_T>( &*blockIt, flagFieldID, boundaryFlagUID, domainFlagUID );
}
template<typename FlagField_T>
void fillFromFlagField( IBlock * block, ConstBlockDataID flagFieldID,
FlagUID boundaryFlagUID, FlagUID domainFlagUID )
{
auto * indexVectors = block->getData< IndexVectors > ( indexVectorID );
auto & indexVectorAll = indexVectors->indexVector(IndexVectors::ALL);
auto & indexVectorInner = indexVectors->indexVector(IndexVectors::INNER);
auto & indexVectorOuter = indexVectors->indexVector(IndexVectors::OUTER);
auto * flagField = block->getData< FlagField_T > ( flagFieldID );
if( !(flagField->flagExists(boundaryFlagUID) && flagField->flagExists(domainFlagUID) ))
return;
auto boundaryFlag = flagField->getFlag(boundaryFlagUID);
auto domainFlag = flagField->getFlag(domainFlagUID);
auto inner = flagField->xyzSize();
inner.expand( cell_idx_t(-1) );
indexVectorAll.clear();
indexVectorInner.clear();
indexVectorOuter.clear();
for( auto it = flagField->begin(); it != flagField->end(); ++it )
{
if( ! isFlagSet(it, domainFlag) )
continue;
{%- for dirIdx, offset in stencil_info %}
if ( isFlagSet( it.neighbor({{offset}} {%if dim == 3%}, 0 {%endif %}), boundaryFlag ) )
{
auto element = {{StructName}}(it.x(), it.y(), {%if dim == 3%} it.z(), {%endif %} {{dirIdx}} );
indexVectorAll.push_back( element );
if( inner.contains( it.x(), it.y(), it.z() ) )
indexVectorInner.push_back( element );
else
indexVectorOuter.push_back( element );
}
{% endfor %}
}
indexVectors->syncGPU();
}
private:
void run( IBlock * block, IndexVectors::Type type{% if target == 'gpu'%}, cudaStream_t stream = 0 {%endif%});
BlockDataID indexVectorID;
public:
{{kernel|generate_members(('indexVector', 'indexVectorSize'))|indent(4)}}
};
} // namespace {{namespace}}
} // namespace walberla
//======================================================================================================================
//
// This file is part of waLBerla. waLBerla is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// waLBerla is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \\author Martin Bauer <martin.bauer@fau.de>
//======================================================================================================================
#include <cmath>
#include "core/DataTypes.h"
#include "core/Macros.h"
#include "lbm/field/PdfField.h"
#include "lbm/sweeps/Streaming.h"
#include "{{class_name}}.h"
#ifdef _MSC_VER
# pragma warning( disable : 4458 )
#endif
{% if target is equalto 'cpu' -%}
#define FUNC_PREFIX
{%- elif target is equalto 'gpu' -%}
#define FUNC_PREFIX __global__
{%- endif %}
#if ( defined WALBERLA_CXX_COMPILER_IS_GNU ) || ( defined WALBERLA_CXX_COMPILER_IS_CLANG )
# pragma GCC diagnostic push
# pragma GCC diagnostic ignored "-Wfloat-equal"
# pragma GCC diagnostic ignored "-Wshadow"
# pragma GCC diagnostic ignored "-Wconversion"
# pragma GCC diagnostic ignored "-Wunused-variable"
# pragma GCC diagnostic ignored "-Wunused-parameter"
#endif
{% for header in headers %}
#include {{header}}
{% endfor %}
using namespace std;
namespace walberla {
namespace {{namespace}} {
{{stream_collide_kernel|generate_definition(target)}}
{{collide_kernel|generate_definition(target)}}
{{stream_kernel|generate_definition(target)}}
const real_t {{class_name}}::w[{{Q}}] = { {{weights}} };
const real_t {{class_name}}::wInv[{{Q}}] = { {{inverse_weights}} };
void {{class_name}}::Sweep::streamCollide( IBlock * block, const uint_t numberOfGhostLayersToInclude )
{
{{stream_collide_kernel|generate_block_data_to_field_extraction(parameters=['pdfs', 'pdfs_tmp'])|indent(4)}}
auto & lm = dynamic_cast< lbm::PdfField<{{class_name}}> * > (pdfs)->latticeModel();
WALBERLA_ASSERT_EQUAL( *(lm.blockId_), block->getId() );
{{stream_collide_kernel|generate_refs_for_kernel_parameters(prefix='lm.', parameters_to_ignore=['pdfs', 'pdfs_tmp'])|indent(4) }}
{{stream_collide_kernel|generate_call('cell_idx_c(numberOfGhostLayersToInclude)')|indent(4)}}
{{stream_collide_kernel|generate_swaps|indent(4)}}
}
void {{class_name}}::Sweep::collide( IBlock * block, const uint_t numberOfGhostLayersToInclude )
{
{{collide_kernel|generate_block_data_to_field_extraction(parameters=['pdfs'])|indent(4)}}
auto & lm = dynamic_cast< lbm::PdfField<{{class_name}}> * > (pdfs)->latticeModel();
WALBERLA_ASSERT_EQUAL( *(lm.blockId_), block->getId() );
{{collide_kernel|generate_refs_for_kernel_parameters(prefix='lm.', parameters_to_ignore=['pdfs', 'pdfs_tmp'])|indent(4) }}
{{collide_kernel|generate_call('cell_idx_c(numberOfGhostLayersToInclude)')|indent(4)}}
}
void {{class_name}}::Sweep::stream( IBlock * block, const uint_t numberOfGhostLayersToInclude )
{
{{stream_kernel|generate_block_data_to_field_extraction(parameters=['pdfs', 'pdfs_tmp'])|indent(4)}}
{{stream_kernel|generate_call('cell_idx_c(numberOfGhostLayersToInclude)')|indent(4)}}
{{stream_kernel|generate_swaps|indent(4)}}
}
} // namespace {{namespace}}
} // namespace walberla
// Buffer Packing
namespace walberla {
namespace mpi {
mpi::SendBuffer & operator<< (mpi::SendBuffer & buf, const ::walberla::{{namespace}}::{{class_name}} & lm)
{
buf << lm.currentLevel;
return buf;
}
mpi::RecvBuffer & operator>> (mpi::RecvBuffer & buf, ::walberla::{{namespace}}::{{class_name}} & lm)
{
buf >> lm.currentLevel;
return buf;
}
} // namespace mpi
} // namespace walberla
#if ( defined WALBERLA_CXX_COMPILER_IS_GNU ) || ( defined WALBERLA_CXX_COMPILER_IS_CLANG )
# pragma GCC diagnostic pop