Debuging the shared pointers. (473f36d6) · Commits · TNL / tnl-dev

src/core/tnlSharedPointer.h

+94 −19

Original line number	Diff line number	Diff line
		@@ -22,14 +22,20 @@
		#include <core/tnlCuda.h>
		#include <core/tnlSmartPointer.h>

		template< typename Object, typename Device = typename Object::DeviceType >
		/***
		* Use the lazy mode if you do not want to call the object constructor in the
		* shared pointer constructor. You may call it later via the method recreate.
		*/
		template< typename Object,
		typename Device = typename Object::DeviceType,
		bool lazy = false >
		class tnlSharedPointer
		{
		static_assert( ! std::is_same< Device, void >::value, "The device cannot be void. You need to specify the device explicitly in your code." );
		};

		template< typename Object >
		class tnlSharedPointer< Object, tnlHost > : public tnlSmartPointer
		template< typename Object, bool lazy >
		class tnlSharedPointer< Object, tnlHost, lazy > : public tnlSmartPointer
		{
		public:

		@@ -37,9 +43,16 @@ class tnlSharedPointer< Object, tnlHost > : public tnlSmartPointer
		typedef tnlHost DeviceType;
		typedef tnlSharedPointer< Object, tnlHost > ThisType;

		explicit tnlSharedPointer()
		template< typename... Args >
		explicit tnlSharedPointer( Args... args )
		: counter( 0 ), pointer( 0 )
		{
		if( ! lazy )
		{
		this->counter = new int;
		this->pointer = new Object( args... );
		*( this->counter ) = 1;
		}
		}

		tnlSharedPointer( const ThisType& pointer )
		@@ -50,17 +63,32 @@ class tnlSharedPointer< Object, tnlHost > : public tnlSmartPointer
		}

		template< typename... Args >
		bool create( const Args... args )
		bool recreate( Args... args )
		{
		std::cerr << "Creating new shared pointer..." << std::endl;
		this->free();
		if( ! this->counter )
		{
		this->counter = new int;
		*this->counter = 1;
		this->pointer = new ObjectType( args... );
		return true;
		}
		if( *this->counter == 1 )
		{
		/****
		* The object is not shared
		*/
		this->pointer->~ObjectType();
		new ( this->pointer ) ObjectType( args... );
		return true;
		}
		( *this->counter )--;
		this->pointer = new Object( args... );
		this->counter = new int;
		if( ! this->pointer \|\| ! this->counter )
		return false;
		*( this->counter ) = 1;
		return true;

		}

		const Object* operator->() const
		@@ -89,6 +117,7 @@ class tnlSharedPointer< Object, tnlHost > : public tnlSmartPointer
		return *( this->pointer );
		}

		template< typename Device = tnlHost >
		Object& modifyData()
		{
		return *( this->pointer );
		@@ -147,8 +176,8 @@ class tnlSharedPointer< Object, tnlHost > : public tnlSmartPointer
		int* counter;
		};

		template< typename Object >
		class tnlSharedPointer< Object, tnlCuda > : public tnlSmartPointer
		template< typename Object, bool lazy >
		class tnlSharedPointer< Object, tnlCuda, lazy > : public tnlSmartPointer
		{
		public:

		@@ -156,10 +185,22 @@ class tnlSharedPointer< Object, tnlCuda > : public tnlSmartPointer
		typedef tnlHost DeviceType;
		typedef tnlSharedPointer< Object, tnlCuda > ThisType;

		explicit tnlSharedPointer()
		template< typename... Args >
		explicit tnlSharedPointer( Args... args )
		: counter( 0 ), cuda_pointer( 0 ),
		pointer( 0 ), modified( false )
		{
		if( ! lazy )
		{
		this->counter = new int;
		this->pointer = new Object( args... );
		#ifdef HAVE_CUDA
		this->cuda_pointer = tnlCuda::passToDevice( *this->pointer );
		if( ! checkCudaDevice )
		return;
		tnlCuda::insertSmartPointer( this );
		#endif
		}
		}

		tnlSharedPointer( const ThisType& pointer )
		@@ -171,12 +212,36 @@ class tnlSharedPointer< Object, tnlCuda > : public tnlSmartPointer
		*counter++;
		}


		template< typename... Args >
		bool create( const Args... args )
		bool recreate( Args... args )
		{
		std::cerr << "Creating new shared pointer..." << std::endl;
		this->free();
		if( ! this->counter )
		{
		this->counter = new int;
		*this->counter = 1;
		this->pointer = new ObjectType( args... );
		#ifdef HAVE_CUDA
		this->cuda_pointer = tnlCuda::passToDevice( *this->object );
		if( ! checkCudaDevice )
		return false;
		tnlCuda::insertSmartPointer( this );
		#endif
		return true;
		}
		if( *this->counter == 1 )
		{
		/****
		* The object is not shared
		*/
		this->pointer->~ObjectType();
		new ( this->pointer ) ObjectType( args... );
		#ifdef HAVE_CUDA
		cudaMemcpy( this->cuda_pointer, this->pointer, sizeof( Object ), cudaMemcpyHostToDevice );
		#endif
		return true;
		}

		this->modified = false;
		this->counter= new int;
		this->pointer = new Object( args... );
		@@ -228,11 +293,20 @@ class tnlSharedPointer< Object, tnlCuda > : public tnlSmartPointer
		return *( this->cuda_pointer );
		}

		template< typename Device = tnlHost >
		__cuda_callable__
		Object& modifyData()
		{
		if( std::is_same< Device, tnlHost >::value )
		{
		this->modified = true;
		return *( this->pointer );
		}
		if( std::is_same< Device, tnlCuda >::value )
		{
		return *( this->cuda_pointer );
		}
		}

		/*const ThisType& operator=( ThisType&& ptr )
		{
		@@ -270,12 +344,13 @@ class tnlSharedPointer< Object, tnlCuda > : public tnlSmartPointer
		#ifdef HAVE_CUDA
		if( this->modified )
		{
		//std::cerr << "Synchronizing data..." << std::endl;
		std::cerr << "Synchronizing data ( " << sizeof( ObjectType ) << " bytes ) to adress " << this->cuda_pointer << "..." << std::endl;
		tnlAssert( this->pointer, );
		tnlAssert( this->cuda_pointer, );
		cudaMemcpy( this->cuda_pointer, this->pointer, sizeof( ObjectType ), cudaMemcpyHostToDevice );
		if( ! checkCudaDevice )
		return false;
		this->modified = false;
		return true;
		}
		#else

src/mesh/grids/tnlGridTraverser_impl.h

+2 −0

Original line number	Diff line number	Diff line
		@@ -353,6 +353,8 @@ processEntities(
		UserData* kernelUserData = tnlCuda::passToDevice( userData );
		checkCudaDevice;

		std::cerr << "User data size is " << sizeof( UserData ) << std::endl;

		dim3 cudaBlockSize( 16, 16 );
		dim3 cudaBlocks;
		cudaBlocks.x = tnlCuda::getNumberOfBlocks( end.x() - begin.x() + 1, cudaBlockSize.x );

src/problems/tnlHeatEquationEocProblem_impl.h

+0 −3

Original line number	Diff line number	Diff line
		@@ -47,9 +47,6 @@ bool
		tnlHeatEquationEocProblem< Mesh, BoundaryCondition, RightHandSide, DifferentialOperator >::
		setup( const tnlParameterContainer& parameters )
		{
		this->boundaryConditionPointer.create();
		this->differentialOperatorPointer.create();
		this->rightHandSidePointer.create();
		if( ! this->boundaryConditionPointer->setup( parameters ) \|\|
		! this->rightHandSidePointer->setup( parameters ) )
		return false;

src/problems/tnlHeatEquationProblem.h

+2 −1

Original line number	Diff line number	Diff line
		@@ -47,6 +47,7 @@ class tnlHeatEquationProblem : public tnlPDEProblem< Mesh,
		typedef typename Mesh::DeviceType DeviceType;
		typedef typename DifferentialOperator::IndexType IndexType;
		typedef tnlMeshFunction< Mesh > MeshFunctionType;
		typedef tnlSharedPointer< MeshFunctionType, DeviceType > MeshFunctionPointer;
		typedef tnlPDEProblem< Mesh, RealType, DeviceType, IndexType > BaseType;
		typedef tnlCSRMatrix< RealType, DeviceType, IndexType > MatrixType;
		typedef tnlSharedPointer< DifferentialOperator > DifferentialOperatorPointer;
		@@ -109,7 +110,7 @@ class tnlHeatEquationProblem : public tnlPDEProblem< Mesh,

		protected:

		MeshFunctionType u;
		MeshFunctionPointer uPointer;

		DifferentialOperatorPointer differentialOperatorPointer;

src/problems/tnlHeatEquationProblem_impl.h

+11 −16

Original line number	Diff line number	Diff line
		@@ -89,10 +89,6 @@ bool
		tnlHeatEquationProblem< Mesh, BoundaryCondition, RightHandSide, DifferentialOperator >::
		setup( const tnlParameterContainer& parameters )
		{
		this->gpuTransferTimer.reset();
		this->boundaryConditionPointer.create();
		this->differentialOperatorPointer.create();
		this->rightHandSidePointer.create();
		if( ! this->boundaryConditionPointer->setup( parameters, "boundary-conditions-" ) \|\|
		! this->rightHandSidePointer->setup( parameters, "right-hand-side-" ) )
		return false;
		@@ -123,7 +119,7 @@ bindDofs( const MeshPointer& meshPointer,
		const DofVectorType& dofVector )
		{
		const IndexType dofs = meshPointer->template getEntitiesCount< typename MeshType::Cell >();
		this->u.bind( meshPointer, dofVector );
		this->uPointer->bind( meshPointer, dofVector );
		}

		template< typename Mesh,
		@@ -139,7 +135,7 @@ setInitialCondition( const tnlParameterContainer& parameters,
		{
		this->bindDofs( meshPointer, dofs );
		const tnlString& initialConditionFile = parameters.getParameter< tnlString >( "initial-condition" );
		if( ! this->u.boundLoad( initialConditionFile ) )
		if( ! this->uPointer->boundLoad( initialConditionFile ) )
		{
		cerr << "I am not able to load the initial condition from the file " << initialConditionFile << "." << endl;
		return false;
		@@ -160,7 +156,6 @@ setupLinearSystem( const MeshPointer& meshPointer,
		const IndexType dofs = this->getDofs( meshPointer );
		typedef typename MatrixPointer::ObjectType::CompressedRowsLengthsVector CompressedRowsLengthsVectorType;
		tnlSharedPointer< CompressedRowsLengthsVectorType > rowLengthsPointer;
		rowLengthsPointer.create();
		if( ! rowLengthsPointer->setSize( dofs ) )
		return false;
		tnlMatrixSetter< MeshType, DifferentialOperator, BoundaryCondition, CompressedRowsLengthsVectorType > matrixSetter;
		@@ -194,7 +189,7 @@ makeSnapshot( const RealType& time,
		//cout << "dofs = " << dofs << endl;
		tnlString fileName;
		FileNameBaseNumberEnding( "u-", step, 5, ".tnl", fileName );
		if( ! this->u.save( fileName ) )
		if( ! this->uPointer->save( fileName ) )
		return false;
		return true;
		}
		@@ -223,7 +218,7 @@ getExplicitRHS( const RealType& time,

		//cout << "u = " << u << endl;
		this->bindDofs( meshPointer, uDofs );
		MeshFunctionType fu( meshPointer, fuDofs );
		MeshFunctionPointer fuPointer( meshPointer, fuDofs );
		tnlExplicitUpdater< Mesh, MeshFunctionType, DifferentialOperator, BoundaryCondition, RightHandSide > explicitUpdater;
		explicitUpdater.setGPUTransferTimer( this->gpuTransferTimer );
		explicitUpdater.template update< typename Mesh::Cell >(
		@@ -232,8 +227,8 @@ getExplicitRHS( const RealType& time,
		this->differentialOperatorPointer,
		this->boundaryConditionPointer,
		this->rightHandSidePointer,
		this->u,
		fu );
		this->uPointer,
		fuPointer );
		/*tnlBoundaryConditionsSetter< MeshFunctionType, BoundaryCondition > boundaryConditionsSetter;
		boundaryConditionsSetter.template apply< typename Mesh::Cell >(
		this->boundaryCondition,
		@@ -254,14 +249,14 @@ template< typename Mesh,
		typename BoundaryCondition,
		typename RightHandSide,
		typename DifferentialOperator >
		template< typename Matrix >
		template< typename MatrixPointer >
		void
		tnlHeatEquationProblem< Mesh, BoundaryCondition, RightHandSide, DifferentialOperator >::
		assemblyLinearSystem( const RealType& time,
		const RealType& tau,
		const MeshPointer& meshPointer,
		const DofVectorType& dofs,
		Matrix& matrix,
		MatrixPointer& matrixPointer,
		DofVectorType& b,
		MeshDependentDataType& meshDependentData )
		{
		@@ -272,7 +267,7 @@ assemblyLinearSystem( const RealType& time,
		BoundaryCondition,
		RightHandSide,
		tnlBackwardTimeDiscretisation,
		Matrix,
		typename MatrixPointer::ObjectType,
		DofVectorType > systemAssembler;
		systemAssembler.template assembly< typename Mesh::Cell >(
		time,
		@@ -281,8 +276,8 @@ assemblyLinearSystem( const RealType& time,
		this->differentialOperatorPointer,
		this->boundaryConditionPointer,
		this->rightHandSidePointer,
		this->u,
		matrix,
		this->uPointer,
		matrixPointer,
		b );
		/*matrix.print( cout );
		cout << endl << b << endl;