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

A keyword __cuda_callable__ was introduced. 

Min. interations feature was fixed in the GMRES method.
Timers for explicit and semi-implicit timestepping were added.
parent 490cf20e

examples/heat-equation/tnl-heat-equation-eoc.h

+1 −1
Original line number Diff line number Diff line
@@ -21,7 +21,7 @@ 
#include <solvers/tnlSolver.h>

#include <solvers/tnlFastBuildConfig.h>

#include <solvers/tnlConfigTags.h>

#include <functions/tnlTestFunction.h>

#include <functors/tnlTestFunction.h>

#include <operators/diffusion/tnlLinearDiffusion.h>

#include <operators/diffusion/tnlExactLinearDiffusion.h>

#include <operators/tnlAnalyticDirichletBoundaryConditions.h>

examples/heat-equation/tnl-heat-equation.h

+1 −1
Original line number Diff line number Diff line
@@ -26,7 +26,7 @@ 
#include <operators/tnlDirichletBoundaryConditions.h>

#include <operators/tnlAnalyticNeumannBoundaryConditions.h>

#include <operators/tnlNeumannBoundaryConditions.h>

#include <functions/tnlConstantFunction.h>

#include <functors/tnlConstantFunction.h>

#include <problems/tnlHeatEquationProblem.h>



//typedef tnlDefaultConfigTag BuildConfig;

src/CMakeLists.txt

+3 −3
Original line number Diff line number Diff line 
INCLUDE_DIRECTORIES( config )

ADD_SUBDIRECTORY( functions )

ADD_SUBDIRECTORY( functors )

ADD_SUBDIRECTORY( config )

ADD_SUBDIRECTORY( core )

ADD_SUBDIRECTORY( debug )
@@ -10,7 +10,7 @@ ADD_SUBDIRECTORY( problems ) 
ADD_SUBDIRECTORY( solvers )

ADD_SUBDIRECTORY( legacy )



set( tnl_SOURCES ${tnl_functions_SOURCES}

set( tnl_SOURCES ${tnl_functors_SOURCES}

                 ${tnl_config_SOURCES}

                 ${tnl_core_SOURCES}

                 ${tnl_legacy_SOURCES}
@@ -21,7 +21,7 @@ set( tnl_SOURCES ${tnl_functions_SOURCES} 
                 ${tnl_problems_SOURCES}

                  )



set( tnl_CUDA__SOURCES ${tnl_functions_CUDA__SOURCES}

set( tnl_CUDA__SOURCES ${tnl_functors_CUDA__SOURCES}

                       ${tnl_config_CUDA__SOURCES}

                       ${tnl_core_CUDA__SOURCES}

                       ${tnl_legacy_CUDA__SOURCES}

src/core/arrays/tnlArray.h

+5 −20
Original line number Diff line number Diff line
@@ -61,24 +61,15 @@ class tnlArray : public virtual tnlObject 


   void reset();



#ifdef HAVE_CUDA

   __device__ __host__

#endif

   Index getSize() const;

   __cuda_callable__ Index getSize() const;



   void setElement( const Index i, const Element& x );



   Element getElement( Index i ) const;



#ifdef HAVE_CUDA

   __device__ __host__

#endif

   Element& operator[] ( Index i );

   __cuda_callable__ Element& operator[] ( Index i );



#ifdef HAVE_CUDA

   __device__ __host__

#endif

   const Element& operator[] ( Index i ) const;

   __cuda_callable__ const Element& operator[] ( Index i ) const;



   tnlArray< Element, Device, Index >& operator = ( const tnlArray< Element, Device, Index >& array );
@@ -93,15 +84,9 @@ class tnlArray : public virtual tnlObject 


   void setValue( const Element& e );



#ifdef HAVE_CUDA

   __device__ __host__

#endif

   const Element* getData() const;

   __cuda_callable__ const Element* getData() const;



#ifdef HAVE_CUDA

   __device__ __host__

#endif

   Element* getData();

   __cuda_callable__ Element* getData();



   /*!

    * Returns true if non-zero size is set.

src/core/arrays/tnlArray_impl.h

+5 −15
Original line number Diff line number Diff line
@@ -145,9 +145,7 @@ void tnlArray< Element, Device, Index > :: reset() 
template< typename Element,

          typename Device,

          typename Index >

#ifdef HAVE_CUDA

   __device__ __host__

#endif

__cuda_callable__

Index tnlArray< Element, Device, Index > :: getSize() const

{

   return this -> size;
@@ -182,9 +180,7 @@ Element tnlArray< Element, Device, Index > :: getElement( Index i ) const 
template< typename Element,

          typename Device,

          typename Index >

#ifdef HAVE_CUDA

   __device__ __host__

#endif

__cuda_callable__

Element& tnlArray< Element, Device, Index > :: operator[] ( Index i )

{

   tnlAssert( 0 <= i && i < this -> getSize(),
@@ -198,9 +194,7 @@ Element& tnlArray< Element, Device, Index > :: operator[] ( Index i ) 
template< typename Element,

           typename Device,

           typename Index >

#ifdef HAVE_CUDA

   __device__ __host__

#endif

__cuda_callable__

const Element& tnlArray< Element, Device, Index > :: operator[] ( Index i ) const

{

   tnlAssert( 0 <= i && i < this -> getSize(),
@@ -295,9 +289,7 @@ void tnlArray< Element, Device, Index > :: setValue( const Element& e ) 
template< typename Element,

           typename Device,

           typename Index >

#ifdef HAVE_CUDA

   __device__ __host__

#endif

__cuda_callable__

const Element* tnlArray< Element, Device, Index > :: getData() const

{

   return this -> data;
@@ -306,9 +298,7 @@ const Element* tnlArray< Element, Device, Index > :: getData() const 
template< typename Element,

           typename Device,

           typename Index >

#ifdef HAVE_CUDA

   __device__ __host__

#endif

__cuda_callable__

Element* tnlArray< Element, Device, Index > :: getData()

{

   return this -> data;