From 07c0214b5ecab307742e0f9bb9181d49b257db5f Mon Sep 17 00:00:00 2001
From: Tomas Oberhuber <tomas.oberhuber@fjfi.cvut.cz>
Date: Sun, 2 Oct 2016 10:26:35 +0200
Subject: [PATCH] Fixed mesh entity creation in CUDA grid traverser.
---
CMakeLists.txt | 2 +-
src/TNL/Functions/MeshFunction_impl.h | 5 ++++
.../Meshes/GridDetails/GridTraverser_impl.h | 26 +++++++++++--------
src/TNL/Problems/HeatEquationProblem_impl.h | 11 ++++++--
4 files changed, 30 insertions(+), 14 deletions(-)
diff --git a/CMakeLists.txt b/CMakeLists.txt
index d2fd15a4f5..0cb47189db 100755
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -71,7 +71,7 @@ if( WITH_CUDA STREQUAL "yes" )
# disable false compiler warnings
# reference for the -Xcudafe flag: http://stackoverflow.com/questions/14831051/how-to-disable-compiler-warnings-with-nvcc/17095910#17095910
# list of possible tokens: http://www.ssl.berkeley.edu/~jimm/grizzly_docs/SSL/opt/intel/cc/9.0/lib/locale/en_US/mcpcom.msg
- set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} ; --expt-relaxed-constexpr -Xcudafe "\"--diag_suppress=code_is_unreachable --diag_suppress=implicit_return_from_non_void_function\"")
+ set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} ; -Wno-deprecated-gpu-targets --expt-relaxed-constexpr -Xcudafe "\"--diag_suppress=code_is_unreachable --diag_suppress=implicit_return_from_non_void_function\"")
#AddCompilerFlag( "-DHAVE_NOT_CXX11" ) # -U_GLIBCXX_ATOMIC_BUILTINS -U_GLIBCXX_USE_INT128 " )
set( ALL_CUDA_ARCHS -gencode arch=compute_20,code=sm_20
-gencode arch=compute_30,code=sm_30
diff --git a/src/TNL/Functions/MeshFunction_impl.h b/src/TNL/Functions/MeshFunction_impl.h
index 973f79ac55..97c2408478 100644
--- a/src/TNL/Functions/MeshFunction_impl.h
+++ b/src/TNL/Functions/MeshFunction_impl.h
@@ -134,6 +134,11 @@ setup( const MeshPointer& meshPointer,
if( ! this->load( fileName ) )
return false;
}
+ else
+ {
+ std::cerr << "Missing parameter " << prefix << "file." << std::endl;
+ return false;
+ }
return true;
}
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_impl.h b/src/TNL/Meshes/GridDetails/GridTraverser_impl.h
index 1f5bcb0608..663308e0c6 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser_impl.h
+++ b/src/TNL/Meshes/GridDetails/GridTraverser_impl.h
@@ -121,12 +121,12 @@ GridTraverser1D(
typename GridType::CoordinatesType coordinates;
coordinates.x() = kernelData->begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;
-
- GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
-
- entity.refresh();
if( coordinates.x() <= kernelData->end.x() )
+ {
+ GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
+ entity.refresh();
EntitiesProcessor::processEntity( entity.getMesh(), *userData, entity );
+ }
}
template< typename Real,
@@ -148,6 +148,7 @@ GridBoundaryTraverser1D(
if( threadIdx.x == 0 )
{
coordinates.x() = kernelData->begin.x();
+ printf( "thread %d coord %d \n", threadIdx.x, coordinates.x() );
GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
entity.refresh();
EntitiesProcessor::processEntity( entity.getMesh(), *userData, entity );
@@ -155,6 +156,7 @@ GridBoundaryTraverser1D(
if( threadIdx.x == 1 )
{
coordinates.x() = kernelData->end.x();
+ printf( "thread %d coord %d \n", threadIdx.x, coordinates.x() );
GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
entity.refresh();
EntitiesProcessor::processEntity( entity.getMesh(), *userData, entity );
@@ -325,11 +327,12 @@ GridTraverser2D(
coordinates.x() = kernelData->begin.x() + ( gridXIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;
coordinates.y() = kernelData->begin.y() + ( gridYIdx * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y;
- GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
+
- if( entity.getCoordinates().x() <= kernelData->end.x() &&
- entity.getCoordinates().y() <= kernelData->end.y() )
+ if( coordinates.x() <= kernelData->end.x() &&
+ coordinates.y() <= kernelData->end.y() )
{
+ GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
entity.refresh();
if( ! processOnlyBoundaryEntities || entity.isBoundaryEntity() )
{
@@ -526,12 +529,13 @@ GridTraverser3D(
coordinates.y() = kernelData->begin.y() + ( gridYIdx * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y;
coordinates.z() = kernelData->begin.z() + ( gridZIdx * Devices::Cuda::getMaxGridSize() + blockIdx.z ) * blockDim.z + threadIdx.z;
- GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
+
- if( entity.getCoordinates().x() <= kernelData->end.x() &&
- entity.getCoordinates().y() <= kernelData->end.y() &&
- entity.getCoordinates().z() <= kernelData->end.z() )
+ if( coordinates.x() <= kernelData->end.x() &&
+ coordinates.y() <= kernelData->end.y() &&
+ coordinates.z() <= kernelData->end.z() )
{
+ GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );
entity.refresh();
if( ! processOnlyBoundaryEntities || entity.isBoundaryEntity() )
{
diff --git a/src/TNL/Problems/HeatEquationProblem_impl.h b/src/TNL/Problems/HeatEquationProblem_impl.h
index c9e965fbc5..d7ab9b4a57 100644
--- a/src/TNL/Problems/HeatEquationProblem_impl.h
+++ b/src/TNL/Problems/HeatEquationProblem_impl.h
@@ -85,9 +85,16 @@ setup( const MeshPointer& meshPointer,
const Config::ParameterContainer& parameters,
const String& prefix )
{
- if( ! this->boundaryConditionPointer->setup( meshPointer, parameters, "boundary-conditions-" ) ||
- ! this->rightHandSidePointer->setup( parameters, "right-hand-side-" ) )
+ if( ! this->boundaryConditionPointer->setup( meshPointer, parameters, "boundary-conditions-" ) )
+ {
+ std::cerr << "I was not able to initialize the boundary conditions." << std::endl;
return false;
+ }
+ if( ! this->rightHandSidePointer->setup( parameters, "right-hand-side-" ) )
+ {
+ std::cerr << "I was not able to initialize the right-hand side function." << std::endl;
+ return false;
+ }
return true;
}
--
GitLab