Commit d1678f7c authored by Tomáš Oberhuber's avatar Tomáš Oberhuber
Browse files

The CUDA version with the smart pointers seems to work.

parent 11645636

src/core/tnlSharedPointer.h

+3 −3
Original line number Diff line number Diff line
@@ -371,7 +371,7 @@ class tnlSharedPointer< Object, tnlCuda, lazy, false > : public tnlSmartPointer 
      template< typename... Args >

      bool recreate( Args... args )

      {

         std::cerr << "Creating new shared pointer..." << std::endl;

         //std::cerr << "Creating new shared pointer..." << std::endl;

         if( ! this->counter )

         {

            this->counter = new int;
@@ -500,7 +500,7 @@ class tnlSharedPointer< Object, tnlCuda, lazy, false > : public tnlSmartPointer 
#ifdef HAVE_CUDA

         if( this->modified )

         {

            std::cerr << "Synchronizing data ( " << sizeof( ObjectType ) << " bytes ) to adress " << this->cuda_pointer << "..." << 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 );
@@ -539,7 +539,7 @@ class tnlSharedPointer< Object, tnlCuda, lazy, false > : public tnlSmartPointer 
                  cudaFree( this->cuda_pointer );

               checkCudaDevice;

#endif         

               std::cerr << "Deleting data..." << std::endl;

               //std::cerr << "Deleting data..." << std::endl;

            }

         }

src/core/tnlSmartPointersRegister.cpp

+1 −1
Original line number Diff line number Diff line
@@ -29,7 +29,7 @@ void tnlSmartPointersRegister::insert( tnlSmartPointer* pointer, int deviceId ) 
{

   tnlAssert( deviceId >= 0 && deviceId < this->devicesCount,

              std::cerr << "deviceId = " << deviceId << " devicesCount = " << this->devicesCount );

   std::cerr << "Inserting pointer " << pointer << " to the register..." << std::endl;

   //std::cerr << "Inserting pointer " << pointer << " to the register..." << std::endl;

   pointersOnDevices[ deviceId ].push_back( pointer );

}

src/mesh/grids/tnlGridTraverser_impl.h

+1 −3
Original line number Diff line number Diff line
@@ -353,8 +353,6 @@ 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/solvers/ode/tnlMersonSolver_impl.h

+1 −1
Original line number Diff line number Diff line
@@ -416,7 +416,7 @@ typename Problem :: RealType tnlMersonSolver< Problem > :: computeError( const R 
                                                              &_k5[ gridOffset ],

                                                              &_kAux[ gridOffset ] );

         cudaThreadSynchronize();

         eps = Max( eps, kAux.max() );

         eps = Max( eps, kAux->max() );

      }

#endif

   }

src/solvers/pde/tnlExplicitUpdater.h

+7 −11
Original line number Diff line number Diff line
@@ -148,7 +148,6 @@ class tnlExplicitUpdater 
                                              TraverserUserData& userData,

                                              const GridEntity& entity )

            {

               printf( "BC size %d ptr %p \n", sizeof( BoundaryConditions ), &userData.boundaryConditions.template getData< DeviceType >() );

               ( userData.u.template modifyData< DeviceType >() )( entity ) = ( userData.boundaryConditions.template getData< DeviceType >() )

               ( userData.u.template getData< DeviceType >(),

                 entity,
@@ -169,21 +168,18 @@ class tnlExplicitUpdater 
                                              TraverserUserData& userData,

                                              const EntityType& entity )

            {

               printf( "DIF.OP. size %d ptr %p \n", sizeof( DifferentialOperator ), &userData.differentialOperator.template getData< DeviceType >() );

               printf( "RHS. size %d ptr %p \n", sizeof( RightHandSide ), &userData.rightHandSide.template getData< DeviceType >() );

               /*( *userData.fu )( entity ) =

                  ( userData.differentialOperatorPointer.template getData< DeviceType >() )(

                     *userData.u, 

               ( userData.fu.template modifyData< tnlCuda >() )( entity ) =

                  ( userData.differentialOperator.template getData< DeviceType >() )(

                     userData.u.template getData< tnlCuda >(), 

                     entity,

                     *userData.time );

                     userData.time );



               typedef tnlFunctionAdapter< MeshType, RightHandSide > FunctionAdapter;

               (  *userData.fu )( entity ) += 

               (  userData.fu.template modifyData< tnlCuda >() )( entity ) += 

                  FunctionAdapter::getValue(

                     userData.rightHandSidePointer.template getData< DeviceType >(),

                     userData.rightHandSide.template getData< DeviceType >(),

                     entity,

                     *userData.time );

                */

                     userData.time );

            }

      };