Skip to content
Snippets Groups Projects

Compare revisions

Changes are shown as if the source revision was being merged into the target revision. Learn more about comparing revisions.

Source

Select target project
No results found

Target

Select target project
No results found
Show changes
Expand all Hide whitespace changes
Inline Side-by-side
Showing
with 1139 additions and 43 deletions
src/field/communication/MPIDatatypes.impl.h
View file @ 7da42ec8
...@@ -28,11 +28,12 @@ namespace field { ...@@ -28,11 +28,12 @@ namespace field {
namespace communication { namespace communication {
template<typename T, uint_t fSize> template<typename Field_T>
MPI_Datatype mpiDatatypeSlice( const Field< T, fSize > & field, MPI_Datatype mpiDatatypeSlice( const Field_T & field,
const cell_idx_t xBeg, const cell_idx_t yBeg, const cell_idx_t zBeg, const cell_idx_t fBeg, const cell_idx_t xBeg, const cell_idx_t yBeg, const cell_idx_t zBeg, const cell_idx_t fBeg,
const cell_idx_t xEnd, const cell_idx_t yEnd, const cell_idx_t zEnd, const cell_idx_t fEnd ) const cell_idx_t xEnd, const cell_idx_t yEnd, const cell_idx_t zEnd, const cell_idx_t fEnd )
{ {
typedef typename Field_T::value_type T;
int sizes[4]; int sizes[4];
int subsizes[4]; int subsizes[4];
int starts[4]; int starts[4];
...@@ -101,8 +102,8 @@ MPI_Datatype mpiDatatypeSlice( const Field< T, fSize > & field, ...@@ -101,8 +102,8 @@ MPI_Datatype mpiDatatypeSlice( const Field< T, fSize > & field,
template<typename T, uint_t fSize> template<typename Field_T>
MPI_Datatype mpiDatatype( const Field< T, fSize > & field ) MPI_Datatype mpiDatatype( const Field_T & field )
{ {
return mpiDatatypeSlice( field, return mpiDatatypeSlice( field,
cell_idx_t( 0 ), cell_idx_t( 0 ), cell_idx_t( 0 ), cell_idx_t( 0 ), cell_idx_t( 0 ), cell_idx_t( 0 ), cell_idx_t( 0 ), cell_idx_t( 0 ),
...@@ -111,8 +112,8 @@ MPI_Datatype mpiDatatype( const Field< T, fSize > & field ) ...@@ -111,8 +112,8 @@ MPI_Datatype mpiDatatype( const Field< T, fSize > & field )
} }
template<typename T, uint_t fSize> template<typename Field_T>
MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInterval & interval, cell_idx_t f /*= 0*/ ) MPI_Datatype mpiDatatypeSliceXYZ( const Field_T & field, const CellInterval & interval, cell_idx_t f /*= 0*/ )
{ {
return mpiDatatypeSlice( field, return mpiDatatypeSlice( field,
interval.xMin(), interval.yMin(), interval.zMin(), f, interval.xMin(), interval.yMin(), interval.zMin(), f,
...@@ -120,8 +121,8 @@ MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInt ...@@ -120,8 +121,8 @@ MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInt
} }
template<typename T, uint_t fSize> template<typename Field_T>
MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInterval & interval, const cell_idx_t fBeg, const cell_idx_t fEnd ) MPI_Datatype mpiDatatypeSliceXYZ( const Field_T & field, const CellInterval & interval, const cell_idx_t fBeg, const cell_idx_t fEnd )
{ {
return mpiDatatypeSlice( field, return mpiDatatypeSlice( field,
interval.xMin(), interval.yMin(), interval.zMin(), fBeg, interval.xMin(), interval.yMin(), interval.zMin(), fBeg,
...@@ -129,9 +130,11 @@ MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInt ...@@ -129,9 +130,11 @@ MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInt
} }
template<typename T, uint_t fSize> template<typename Field_T>
MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInterval & interval, const std::set<cell_idx_t> & fs ) MPI_Datatype mpiDatatypeSliceXYZ( const Field_T & field, const CellInterval & interval, const std::set<cell_idx_t> & fs )
{ {
typedef typename Field_T::value_type T;
MPI_Datatype newType = MPI_DATATYPE_NULL; MPI_Datatype newType = MPI_DATATYPE_NULL;
int sizes[3]; int sizes[3];
...@@ -206,14 +209,14 @@ MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInt ...@@ -206,14 +209,14 @@ MPI_Datatype mpiDatatypeSliceXYZ( const Field< T, fSize > & field, const CellInt
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeWithGhostLayer( const GhostLayerField< T, fSize > & field ) MPI_Datatype mpiDatatypeWithGhostLayer( const GhostLayerField_T & field )
{ {
return mpiDatatypeWithGhostLayer( field, field.nrOfGhostLayers() ); return mpiDatatypeWithGhostLayer( field, field.nrOfGhostLayers() );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeWithGhostLayer( const GhostLayerField< T, fSize > & field, const uint_t numGhostLayers ) MPI_Datatype mpiDatatypeWithGhostLayer( const GhostLayerField_T & field, const uint_t numGhostLayers )
{ {
const cell_idx_t xBeg = - cell_idx_c( numGhostLayers ); const cell_idx_t xBeg = - cell_idx_c( numGhostLayers );
const cell_idx_t yBeg = - cell_idx_c( numGhostLayers ); const cell_idx_t yBeg = - cell_idx_c( numGhostLayers );
...@@ -231,14 +234,14 @@ MPI_Datatype mpiDatatypeWithGhostLayer( const GhostLayerField< T, fSize > & fiel ...@@ -231,14 +234,14 @@ MPI_Datatype mpiDatatypeWithGhostLayer( const GhostLayerField< T, fSize > & fiel
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeGhostLayerOnly( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const bool fullSlice /*= false*/ ) MPI_Datatype mpiDatatypeGhostLayerOnly( const GhostLayerField_T & field, const stencil::Direction dir, const bool fullSlice /*= false*/ )
{ {
return mpiDatatypeGhostLayerOnly( field, field.nrOfGhostLayers(), dir, fullSlice ); return mpiDatatypeGhostLayerOnly( field, field.nrOfGhostLayers(), dir, fullSlice );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeGhostLayerOnly( const GhostLayerField< T, fSize > & field, const uint_t thickness, const stencil::Direction dir, const bool fullSlice /*= false*/ ) MPI_Datatype mpiDatatypeGhostLayerOnly( const GhostLayerField_T & field, const uint_t thickness, const stencil::Direction dir, const bool fullSlice /*= false*/ )
{ {
CellInterval ci; CellInterval ci;
field.getGhostRegion( dir, ci, cell_idx_c( thickness ), fullSlice ); field.getGhostRegion( dir, ci, cell_idx_c( thickness ), fullSlice );
...@@ -250,8 +253,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnly( const GhostLayerField< T, fSize > & fiel ...@@ -250,8 +253,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnly( const GhostLayerField< T, fSize > & fiel
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const bool fullSlice /*= false*/, const cell_idx_t f /*= 0*/ ) MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField_T & field, const stencil::Direction dir, const bool fullSlice /*= false*/, const cell_idx_t f /*= 0*/ )
{ {
CellInterval ci; CellInterval ci;
field.getGhostRegion( dir, ci, cell_idx_c( field.nrOfGhostLayers() ), fullSlice ); field.getGhostRegion( dir, ci, cell_idx_c( field.nrOfGhostLayers() ), fullSlice );
...@@ -259,8 +262,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & f ...@@ -259,8 +262,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & f
return mpiDatatypeSliceXYZ( field, ci, f ); return mpiDatatypeSliceXYZ( field, ci, f );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const bool fullSlice, const cell_idx_t fBeg, const cell_idx_t fEnd ) MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField_T & field, const stencil::Direction dir, const bool fullSlice, const cell_idx_t fBeg, const cell_idx_t fEnd )
{ {
CellInterval ci; CellInterval ci;
field.getGhostRegion( dir, ci, cell_idx_c( field.nrOfGhostLayers() ), fullSlice ); field.getGhostRegion( dir, ci, cell_idx_c( field.nrOfGhostLayers() ), fullSlice );
...@@ -268,8 +271,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & f ...@@ -268,8 +271,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & f
return mpiDatatypeSliceXYZ( field, ci, fBeg, fEnd ); return mpiDatatypeSliceXYZ( field, ci, fBeg, fEnd );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const bool fullSlice, const std::set<cell_idx_t> & fs ) MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField_T & field, const stencil::Direction dir, const bool fullSlice, const std::set<cell_idx_t> & fs )
{ {
CellInterval ci; CellInterval ci;
field.getGhostRegion( dir, ci, cell_idx_c( field.nrOfGhostLayers() ), fullSlice ); field.getGhostRegion( dir, ci, cell_idx_c( field.nrOfGhostLayers() ), fullSlice );
...@@ -277,8 +280,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & f ...@@ -277,8 +280,8 @@ MPI_Datatype mpiDatatypeGhostLayerOnlyXYZ( const GhostLayerField< T, fSize > & f
return mpiDatatypeSliceXYZ( field, ci, fs ); return mpiDatatypeSliceXYZ( field, ci, fs );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeSliceBeforeGhostlayer( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const uint_t thickness /*= 1*/, const bool fullSlice /*= false*/ ) MPI_Datatype mpiDatatypeSliceBeforeGhostlayer( const GhostLayerField_T & field, const stencil::Direction dir, const uint_t thickness /*= 1*/, const bool fullSlice /*= false*/ )
{ {
CellInterval ci; CellInterval ci;
field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice ); field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice );
...@@ -289,8 +292,8 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayer( const GhostLayerField< T, fSize > ...@@ -289,8 +292,8 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayer( const GhostLayerField< T, fSize >
return mpiDatatypeSliceXYZ( field, ci, fBeg, fEnd ); return mpiDatatypeSliceXYZ( field, ci, fBeg, fEnd );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const uint_t thickness /*= 1*/, const cell_idx_t f /*= 0*/, const bool fullSlice /*= false*/ ) MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField_T & field, const stencil::Direction dir, const uint_t thickness /*= 1*/, const cell_idx_t f /*= 0*/, const bool fullSlice /*= false*/ )
{ {
CellInterval ci; CellInterval ci;
field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice ); field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice );
...@@ -298,8 +301,8 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSiz ...@@ -298,8 +301,8 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSiz
return mpiDatatypeSliceXYZ( field, ci, f ); return mpiDatatypeSliceXYZ( field, ci, f );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const uint_t thickness, const cell_idx_t fBeg, const cell_idx_t fEnd, const bool fullSlice ) MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField_T & field, const stencil::Direction dir, const uint_t thickness, const cell_idx_t fBeg, const cell_idx_t fEnd, const bool fullSlice )
{ {
CellInterval ci; CellInterval ci;
field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice ); field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice );
...@@ -307,8 +310,8 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSiz ...@@ -307,8 +310,8 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSiz
return mpiDatatypeSliceXYZ( field, ci, fBeg, fEnd ); return mpiDatatypeSliceXYZ( field, ci, fBeg, fEnd );
} }
template<typename T, uint_t fSize> template<typename GhostLayerField_T>
MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSize > & field, const stencil::Direction dir, const uint_t thickness, const std::set<cell_idx_t> & fs, const bool fullSlice ) MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField_T & field, const stencil::Direction dir, const uint_t thickness, const std::set<cell_idx_t> & fs, const bool fullSlice )
{ {
CellInterval ci; CellInterval ci;
field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice ); field.getSliceBeforeGhostLayer( dir, ci, cell_idx_c( thickness ), fullSlice );
...@@ -319,4 +322,4 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSiz ...@@ -319,4 +322,4 @@ MPI_Datatype mpiDatatypeSliceBeforeGhostlayerXYZ( const GhostLayerField< T, fSiz
} // namespace communication } // namespace communication
} // namespace field } // namespace field
} // namespace walberla } // namespace walberla
\ No newline at end of file
src/field/iterators/FieldIterator.h
View file @ 7da42ec8
...@@ -26,6 +26,7 @@ ...@@ -26,6 +26,7 @@
#include "core/DataTypes.h" #include "core/DataTypes.h"
#include "core/cell/Cell.h" #include "core/cell/Cell.h"
#include "core/debug/Debug.h" #include "core/debug/Debug.h"
#include "field/Layout.h"
#include "stencil/Directions.h" #include "stencil/Directions.h"
...@@ -39,16 +40,6 @@ namespace walberla { ...@@ -39,16 +40,6 @@ namespace walberla {
namespace field { namespace field {
/**
* \brief Layout for field (
* \ingroup field
*/
enum Layout {
fzyx = 0, //!< Value-sorted data layout (f should be outermost loop)
zyxf = 1 //!< Cell-sorted data layout, (f should be innermost loop)
};
template<typename T, uint_t fSize_> class Field; // forward for friend declaration template<typename T, uint_t fSize_> class Field; // forward for friend declaration
......
tests/CMakeLists.txt
View file @ 7da42ec8
...@@ -5,6 +5,7 @@ include_directories( ${walberla_BINARY_DIR}/src ) # for generated headers ...@@ -5,6 +5,7 @@ include_directories( ${walberla_BINARY_DIR}/src ) # for generated headers
add_subdirectory( blockforest ) add_subdirectory( blockforest )
add_subdirectory( boundary ) add_subdirectory( boundary )
add_subdirectory( core ) add_subdirectory( core )
add_subdirectory( cuda )
add_subdirectory( domain_decomposition ) add_subdirectory( domain_decomposition )
add_subdirectory( fft ) add_subdirectory( fft )
add_subdirectory( field ) add_subdirectory( field )
......
tests/cuda/CMakeLists.txt 0 → 100644
View file @ 7da42ec8
###################################################################################################
#
# Tests for cuda
#
###################################################################################################
waLBerla_compile_test( FILES communication/GPUPackInfoTest.cpp DEPENDS blockforest )
waLBerla_execute_test( NAME GPUPackInfoTest )
waLBerla_compile_test( FILES communication/CommTest )
waLBerla_execute_test( NAME CommTest )
waLBerla_compile_test( FILES FieldTransferTest )
waLBerla_execute_test( NAME FieldTransferTest )
waLBerla_compile_test( FILES SimpleKernelTest.cpp Kernels.cu DEPENDS blockforest timeloop gui )
waLBerla_execute_test( NAME SimpleKernelTest )
waLBerla_compile_test( FILES CudaMPI DEPENDS blockforest timeloop gui )
waLBerla_execute_test( NAME CudaMPI )
waLBerla_compile_test( FILES FieldIndexing3DTest.cpp FieldIndexing3DTest.cu )
waLBerla_execute_test( NAME FieldIndexing3DTest )
tests/cuda/CudaMPI.cpp 0 → 100644
View file @ 7da42ec8
//======================================================================================================================
//
// 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 CudaMPI.h
//! \author Martin Bauer <martin.bauer@fau.de>
//
//======================================================================================================================
#include "blockforest/Initialization.h"
#include "core/debug/TestSubsystem.h"
#include "core/Environment.h"
#include "core/logging/Logging.h"
#include "core/mpi/Datatype.h"
#include "cuda/GPUField.h"
#include "field/communication/MPIDatatypes.h"
#include "field/AddToStorage.h"
#include "timeloop/SweepTimeloop.h"
#include "gui/Gui.h"
using namespace walberla;
void fullFieldTransfer()
{
Field<double,4> h_f1 ( 3, 4, 2, 42.0, field::fzyx );
Field<double,4> h_f2 ( 3, 4, 2, 27.0, field::fzyx );
cuda::GPUField<double> d_f ( 3, 4, 2, 4, 8.0, field::fzyx );
// Transfer h_f1 from CPU to GPU d_f
auto h_f1_datatype = mpi::Datatype ( field::communication::mpiDatatype( h_f1 ) );
auto h_f2_datatype = mpi::Datatype ( field::communication::mpiDatatype( h_f2 ) );
auto d_f_datatype = mpi::Datatype ( field::communication::mpiDatatype( d_f ) );
WALBERLA_LOG_DEVEL("ISend");
MPI_Request request1;
MPI_Isend( h_f1.data(), 1, h_f1_datatype, 0, 0, MPI_COMM_WORLD, &request1 );
WALBERLA_LOG_DEVEL("IRecv");
MPI_Request request2;
MPI_Irecv( d_f.data(), 1, d_f_datatype, 0, 0, MPI_COMM_WORLD, &request2 );
MPI_Wait( &request1, MPI_STATUS_IGNORE );
MPI_Wait( &request2, MPI_STATUS_IGNORE );
// Transfer GPU field d_f back to CPU into h_f2
MPI_Request request3;
WALBERLA_LOG_DEVEL("ISend");
MPI_Isend( d_f.data(), 1, d_f_datatype, 0, 0, MPI_COMM_WORLD , &request3 );
MPI_Request request4;
WALBERLA_LOG_DEVEL("IRecv");
MPI_Irecv( h_f2.data(), 1, h_f2_datatype, 0, 0, MPI_COMM_WORLD, &request4 );
MPI_Wait( &request3, MPI_STATUS_IGNORE );
MPI_Wait( &request4, MPI_STATUS_IGNORE );
WALBERLA_CHECK_EQUAL( h_f1, h_f2 );
}
void blockStorageAndGui( int argc, char ** argv )
{
shared_ptr< StructuredBlockForest > blocks = blockforest::createUniformBlockGrid(
uint_c(1) , uint_c(1), uint_c(1), // number of blocks in x,y,z direction
uint_c(5) , uint_c(7), uint_c(3), // number of blocks in x,y,z direction
real_c(1), // dx: length of one cell in physical coordinates
false, // one block per process? - "false" means all blocks to one process
true, true, true ); // no periodicity
typedef GhostLayerField<real_t,1> ScalarField;
BlockDataID cpuFieldID1 = field::addToStorage<ScalarField>( blocks, "CPUField 1", real_c(42), field::fzyx, uint_c(1) );
BlockDataID cpuFieldID2 = field::addToStorage<ScalarField>( blocks, "CPUField 2", real_c(0), field::fzyx, uint_c(1) );
typedef cuda::GPUField<real_t> GPUField;
BlockDataID gpuFieldID = blocks->addStructuredBlockData< GPUField >(
[&] ( IBlock * block, StructuredBlockStorage * const s ) {
return new GPUField( s->getNumberOfXCells(*block),
s->getNumberOfYCells(*block),
s->getNumberOfZCells(*block),
1 , 1.0);
},
"GPU Field" );
for( auto blockIt = blocks->begin(); blockIt != blocks->end(); ++blockIt )
{
// get the field stored on the current block
ScalarField * h_f1 = blockIt->getData<ScalarField>( cpuFieldID1 );
ScalarField * h_f2 = blockIt->getData<ScalarField>( cpuFieldID2 );
GPUField * d_f = blockIt->getData<GPUField> ( gpuFieldID );
auto h_f1_datatype = mpi::Datatype ( field::communication::mpiDatatypeSliceBeforeGhostlayer( *h_f1, stencil::W, 1, true ) );
auto h_f2_datatype = mpi::Datatype ( field::communication::mpiDatatypeSliceBeforeGhostlayer( *h_f2, stencil::W, 1, true ) );
auto d_f_datatype = mpi::Datatype ( field::communication::mpiDatatypeSliceBeforeGhostlayer( *d_f , stencil::W, 1, true ) );
MPI_Sendrecv( const_cast<double *>( h_f1->data() ), 1, h_f1_datatype, 0, 0,
d_f->data(), 1, d_f_datatype , 0, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE );
MPI_Sendrecv( d_f->data(), 1, d_f_datatype, 0, 0,
h_f2->data(), 1, h_f2_datatype, 0, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE );
}
SweepTimeloop timeloop( blocks, 4 );
GUI gui( timeloop, blocks, argc, argv );
gui.run();
}
int main( int argc, char ** argv )
{
debug::enterTestMode();
walberla::Environment walberlaEnv( argc, argv );
fullFieldTransfer();
//blockStorageAndGui(argc, argv);
return 0;
}
tests/cuda/CudaTestCommon.h 0 → 100644
View file @ 7da42ec8
//======================================================================================================================
//
// 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
//! \author Paulo Carvalho <prcjunior@inf.ufpr.br>
//
//======================================================================================================================
#pragma once
#include "core/DataTypes.h"
#include "core/debug/CheckFunctions.h"
#include "field/GhostLayerField.h"
#ifdef DBG_PRINT_ON
#define DBG_PRINT_FIELD( f ) printField( f )
#define DBG_PRINT( fmt, ... ) printf( fmt, ##__VA_ARGS__ )
#else
#define DBG_PRINT_FIELD( f )
#define DBG_PRINT(fmt, ...)
#endif
template<typename Field_T>
void printField( Field_T& field )
{
using namespace walberla;
cell_idx_t fs = 0;
cell_idx_t zs = -(cell_idx_t)field.nrOfGhostLayers();
cell_idx_t ys = -(cell_idx_t)field.nrOfGhostLayers();
cell_idx_t xs = -(cell_idx_t)field.nrOfGhostLayers();
cell_idx_t nf = (cell_idx_t)field.fSize();
cell_idx_t nz = (cell_idx_t)(field.zSize() + field.nrOfGhostLayers());
cell_idx_t ny = (cell_idx_t)(field.ySize() + field.nrOfGhostLayers());
cell_idx_t nx = (cell_idx_t)(field.xSize() + field.nrOfGhostLayers());
for ( cell_idx_t f = fs; f < nf; ++f ) {
std::cout << "{";
for ( cell_idx_t z = zs; z < nz; ++z ) {
std::cout << ( z == zs ? "[" : " [" );
for ( cell_idx_t y = ys; y < ny; ++y ) {
std::cout << "(";
for ( cell_idx_t x = xs; x < nx; ++x ) {
std::cout << field( x, y, z, f ) << ( x == nx-1 ? "" : " " );
}
std::cout << ( y == ny-1 ? ")" : ") " );
}
std::cout << "]\n";
}
std::cout << "}\n";
}
}
#define CHECK_FIELD_EQUAL( f1, f2 ) WALBERLA_CHECK( checkFieldEqual( f1, f2 ), "Field differ" )
template< typename Field_T >
bool checkFieldEqual( Field_T& field1, Field_T& field2 )
{
using namespace walberla;
WALBERLA_ASSERT( field1.xSize() == field2.xSize() &&
field1.ySize() == field2.ySize() &&
field1.zSize() == field2.zSize() &&
field1.fSize() == field2.fSize() );
WALBERLA_FOR_ALL_CELLS_XYZ( &field2,
for ( uint_t f = 0; f < field1.fSize(); ++f )
{
if ( field1.get( x, y, z, f ) != field2.get( x, y, z, f ) )
{
return false;
}
}
)
return true;
}
tests/cuda/FieldIndexing3DTest.cpp 0 → 100644
View file @ 7da42ec8
//======================================================================================================================
//
// 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
//! \author Paulo Carvalho <prcjunior@inf.ufpr.br>
//
//======================================================================================================================
#include "core/debug/TestSubsystem.h"
#include "core/Environment.h"
#include "core/mpi/Datatype.h"
#include "field/GhostLayerField.h"
#include "cuda/GPUField.h"
#include "cuda/FieldCopy.h"
#include "cuda/Kernel.h"
#include "cuda/FieldIndexing3D.h"
//#define DBG_PRINT_ON
#include "CudaTestCommon.h"
#include "FieldIndexing3DTest.h"
using namespace walberla;
typedef cuda::FieldIndexing3D<int> FieldIdx3D_T;
typedef GhostLayerField<int , F_SIZE> HostField_T;
typedef cuda::GPUField<int> GPUField_T;
void xyzTest()
{
const HostField_T emptyField( X_SIZE, Y_SIZE, Z_SIZE, GL_SIZE, -1, LAYOUT );
GPUField_T deviceField( X_SIZE, Y_SIZE, Z_SIZE, F_SIZE, 1, LAYOUT );
cuda::fieldCpy( deviceField, emptyField );
auto setValue = cuda::make_kernel( &setValueKernel );
setValue.addFieldIndexingParam( FieldIdx3D_T::xyz( deviceField ) );
setValue();
HostField_T resultField( X_SIZE, Y_SIZE, Z_SIZE, GL_SIZE, -1, LAYOUT );
cuda::fieldCpy( resultField, deviceField );
HostField_T expectedField( X_SIZE, Y_SIZE, Z_SIZE, GL_SIZE, -1, LAYOUT );
WALBERLA_FOR_ALL_CELLS_XYZ( &expectedField,
for ( uint_t f = 0; f < expectedField.fSize(); ++f )
{
expectedField.get( x, y, z, f ) = IDX4D( x, y, z, f );
}
)
DBG_PRINT_FIELD( resultField );
CHECK_FIELD_EQUAL( resultField, expectedField );
}
void sliceBeforeGhostLayerXYZTest()
{
const HostField_T emptyField( X_SIZE, Y_SIZE, Z_SIZE, GL_SIZE, -1, LAYOUT );
GPUField_T deviceField( X_SIZE, Y_SIZE, Z_SIZE, F_SIZE, 1, LAYOUT );
cuda::fieldCpy( deviceField, emptyField );
auto setValue = cuda::make_kernel( &setValueKernel );
setValue.addFieldIndexingParam( FieldIdx3D_T::sliceBeforeGhostLayerXYZ( deviceField, 1, stencil::B, true ) );
setValue();
HostField_T resultField( X_SIZE, Y_SIZE, Z_SIZE, GL_SIZE, -1, LAYOUT );
cuda::fieldCpy( resultField, deviceField );
HostField_T expectedField( X_SIZE, Y_SIZE, Z_SIZE, GL_SIZE, -1, LAYOUT );
CellInterval ci;
expectedField.getSliceBeforeGhostLayer( stencil::B, ci, 1, true );
WALBERLA_FOR_ALL_CELLS_IN_INTERVAL_XYZ( ci,
for ( uint_t f = 0; f < expectedField.fSize(); ++f )
{
expectedField.get( x, y, z, f ) = IDX4D( x - ci.xMin(), y - ci.yMin(), z - ci.zMin(), f );
}
)
DBG_PRINT_FIELD( resultField );
CHECK_FIELD_EQUAL( resultField, expectedField );
}
int main( int argc, char ** argv )
{
debug::enterTestMode();
walberla::Environment walberlaEnv( argc, argv );
xyzTest();
sliceBeforeGhostLayerXYZTest();
return 0;
}
tests/cuda/FieldIndexing3DTest.cu 0 → 100644
View file @ 7da42ec8
//======================================================================================================================
//
// 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
//! \author Paulo Carvalho <prcjunior@inf.ufpr.br>
//
//======================================================================================================================
#include "FieldIndexing3DTest.h"
namespace walberla {
__global__ void setValueKernel( FieldAccessor3D_T fa )
{
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
unsigned int z = blockIdx.z * blockDim.z + threadIdx.z;
fa.set( blockIdx, threadIdx );
if ( fa.isValidPosition() )
{
for ( int f = 0; f < F_SIZE; ++f )
{
fa.get(f) = IDX4D( x, y, z, f );
}
}
}
} // namespace walberla
tests/cuda/FieldIndexing3DTest.h 0 → 100644
View file @ 7da42ec8
//======================================================================================================================
//
// 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
//! \author Paulo Carvalho <prcjunior@inf.ufpr.br>
//
//======================================================================================================================
#pragma once
#include "cuda/FieldAccessor3D.h"
#define X_SIZE (128-2)
#define Y_SIZE (128-2)
#define Z_SIZE (128-2)
#define F_SIZE 19
#define LAYOUT field::fzyx
#define GL_SIZE 1
#define YOFFSET ( X_SIZE )
#define ZOFFSET ( ( Y_SIZE ) * ( YOFFSET ) )
#define FOFFSET ( ( Z_SIZE ) * ( ZOFFSET ) )
#define IDX4D( x, y, z, f ) ( (int)( (f) * (FOFFSET) + (z) * (Z_SIZE) + (y) * (YOFFSET) + (x) ) )
namespace walberla {
typedef cuda::FieldAccessor3D<int> FieldAccessor3D_T;
__global__ void setValueKernel( FieldAccessor3D_T fa );
} // namespace walberla
tests/cuda/FieldTransferTest.cpp 0 → 100644
View file @ 7da42ec8
//======================================================================================================================
//
// 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 FieldTransferTest.h
//! \author Martin Bauer <martin.bauer@fau.de>
//
//======================================================================================================================
#include "core/debug/TestSubsystem.h"
#include "core/Environment.h"
#include "field/Field.h"
#include "cuda/GPUField.h"
#include "cuda/FieldCopy.h"
using namespace walberla;
void simpleTransfer()
{
Field<double,4> h_f1 ( 16, 20, 30, 42.0, field::fzyx );
Field<double,4> h_f2 ( 16, 20, 30, 0.0, field::fzyx );
cuda::GPUField<double> d_f ( 16,20,30,4,0, field::fzyx );
WALBERLA_CHECK_EQUAL( h_f1.xSize() ,d_f.xSize() );
WALBERLA_CHECK_EQUAL( h_f1.ySize() ,d_f.ySize() );
WALBERLA_CHECK_EQUAL( h_f1.zSize() ,d_f.zSize() );
WALBERLA_CHECK_EQUAL( h_f1.fSize() ,d_f.fSize() );
WALBERLA_CHECK_EQUAL( h_f1.layout() ,d_f.layout() );
cuda::fieldCpy( d_f, h_f1 );
cuda::fieldCpy( h_f2, d_f );
WALBERLA_CHECK_EQUAL( h_f1, h_f2 );
}
int main( int argc, char ** argv )
{
debug::enterTestMode();
walberla::Environment walberlaEnv( argc, argv );
simpleTransfer();
return 0;
}
tests/cuda/Kernels.cu 0 → 100644
View file @ 7da42ec8
#include <iostream>
#include "cuda/FieldAccessor.h"
#include "cuda/FieldIndexing.h"
namespace walberla {
namespace cuda {
template<typename T>
class GPUField;
}
__global__ void kernel_double( cuda::FieldAccessor<double> f )
{
f.set( blockIdx, threadIdx );
f.get() *= 2.0;
}
void kernel_double_field( const cuda::GPUField<double> & field )
{
using namespace std;
cuda::FieldIndexing<double> iter = cuda::FieldIndexing<double>::sliceBeforeGhostLayerXYZ( field, 1, stencil::E, true );
std::cout << "Kernel call dims "
<< iter.blockDim().x << ","
<< iter.gridDim().x << ","
<< iter.gridDim().y << ","
<< iter.gridDim().z << endl;
kernel_double<<< iter.gridDim(), iter.blockDim() >>> ( iter.gpuAccess() );
}
} // namespace walberla
tests/cuda/SimpleKernelTest.cpp 0 → 100644
View file @ 7da42ec8
//======================================================================================================================
//
// 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 FieldTransferTest.h
//! \author Martin Bauer <martin.bauer@fau.de>
//
//======================================================================================================================
#include "cuda/FieldIndexing.h"
#include "blockforest/Initialization.h"
#include "core/debug/TestSubsystem.h"
#include "core/Environment.h"
#include "field/GhostLayerField.h"
#include "cuda/GPUField.h"
#include "cuda/FieldCopy.h"
#include "cuda/Kernel.h"
#include "gui/Gui.h"
#include "timeloop/SweepTimeloop.h"
using namespace walberla;
namespace walberla{
void kernel_double_field( const cuda::GPUField<double> & field );
void kernel_double( cuda::FieldAccessor<double> f );
}
GhostLayerField<real_t,1> * createCPUField( IBlock* const block, StructuredBlockStorage* const storage )
{
return new GhostLayerField<real_t,1> (
storage->getNumberOfXCells( *block ), // number of cells in x direction
storage->getNumberOfYCells( *block ), // number of cells in y direction
storage->getNumberOfZCells( *block ), // number of cells in z direction
1, // number of ghost layers
real_t(1), // initial value
field::fzyx);
}
cuda::GPUField<real_t> * createGPUField( IBlock* const block, StructuredBlockStorage* const storage )
{
return new cuda::GPUField<real_t> (
storage->getNumberOfXCells( *block ), // number of cells in x direction
storage->getNumberOfYCells( *block ), // number of cells in y direction
storage->getNumberOfZCells( *block ), // number of cells in z direction
1, // fSize
1, // number of ghost layers
field::fzyx );
}
int main( int argc, char ** argv )
{
walberla::Environment env( argc, argv );
debug::enterTestMode();
shared_ptr< StructuredBlockForest > blocks = blockforest::createUniformBlockGrid (
uint_t(1), uint_t(1), uint_t(1), // number of blocks in x,y,z direction
uint_t(14), uint_t(14), uint_t(14), // how many cells per block (x,y,z)
real_c(0.5), // dx: length of one cell in physical coordinates
false, // one block per process - "false" means all blocks to one process
false, false, false ); // no periodicity
BlockDataID cpuFieldID = blocks->addStructuredBlockData< GhostLayerField<real_t,1> > ( &createCPUField, "CPUField" );
BlockDataID gpuFieldID = blocks->addStructuredBlockData< cuda::GPUField<real_t> > ( &createGPUField, "GPUField" );
for ( auto blockIterator = blocks->begin(); blockIterator != blocks->end(); ++blockIterator )
{
IBlock & currentBlock = *blockIterator;
// get the field stored on the current block
auto cpuField = currentBlock.getData< GhostLayerField<real_t,1> > ( cpuFieldID );
auto gpuField = currentBlock.getData< cuda::GPUField<real_t> > ( gpuFieldID );
cuda::fieldCpy( *gpuField, *cpuField );
auto myKernel = cuda::make_kernel( &kernel_double );
auto indexing = cuda::FieldIndexing<double>::sliceBeforeGhostLayerXYZ( *gpuField, 1, stencil::W, true );
myKernel.addFieldIndexingParam(indexing);
myKernel();
cuda::fieldCpy( *cpuField, *gpuField );
WALBERLA_ASSERT_EQUAL( cpuField->get(0,0,0), 2 );
}
//SweepTimeloop timeloop ( blocks, uint_t(1) );
//timeloop.run();
//GUI gui ( timeloop, blocks, argc, argv );
//gui.run();
return 0;
}
tests/cuda/communication/CommTest.cpp 0 → 100644
View file @ 7da42ec8
This diff is collapsed.
tests/cuda/communication/GPUPackInfoTest.cpp 0 → 100644
View file @ 7da42ec8
This diff is collapsed.