Fixed mesh entity creation in CUDA grid traverser. (07c0214b) · Commits · TNL / tnl-dev

CMakeLists.txt

+1 −1

Original line number	Diff line number	Diff line
		@@ -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

src/TNL/Functions/MeshFunction_impl.h

+5 −0

Original line number	Diff line number	Diff line
		@@ -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;
		}

src/TNL/Meshes/GridDetails/GridTraverser_impl.h

+15 −11

Original line number	Diff line number	Diff line
		@@ -121,13 +121,13 @@ GridTraverser1D(
		typename GridType::CoordinatesType coordinates;

		coordinates.x() = kernelData->begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;

		if( coordinates.x() <= kernelData->end.x() )
		{
		GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );

		entity.refresh();
		if( coordinates.x() <= kernelData->end.x() )
		EntitiesProcessor::processEntity( entity.getMesh(), *userData, entity );
		}
		}

		template< typename Real,
		typename Index,
		@@ -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() )
		{

src/TNL/Problems/HeatEquationProblem_impl.h

+9 −2

Original line number	Diff line number	Diff line
		@@ -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;
		}