Tunning experiments - included explicit updater entyties processor.

tests/benchmarks/heat-equation-benchmark/HeatEquationBenchmarkProblem_impl.h

@@ -525,20 +525,25 @@ getExplicitUpdate( const RealType& time,
       }
       if( this->cudaKernelType == "tunning" )
       {
+         using UserData = ExplicitUpdaterTraverserUserData< RealType,
+            MeshFunctionType,
+            DifferentialOperator,
+            BoundaryCondition,
+            RightHandSide >;
+         using ExplicitUpdaterType = TNL::ExplicitUpdater< Mesh, MeshFunctionType, DifferentialOperator, BoundaryCondition, RightHandSide >;
+         using IntertiorEntitiesProcessor = typename ExplicitUpdaterType::TraverserInteriorEntitiesProcessor;
+         using BoundaryEntitiesProcessor = typename ExplicitUpdaterType::TraverserBoundaryEntitiesProcessor;
+         
          const IndexType gridXSize = mesh->getDimensions().x();
          const IndexType gridYSize = mesh->getDimensions().y();
-         const RealType& hx_inv = mesh->template getSpaceStepsProducts< -2,  0 >();
-         const RealType& hy_inv = mesh->template getSpaceStepsProducts<  0, -2 >();
+         /*const RealType& hx_inv = mesh->template getSpaceStepsProducts< -2,  0 >();
+         const RealType& hy_inv = mesh->template getSpaceStepsProducts<  0, -2 >();*/
 
          dim3 cudaBlockSize( 16, 16 );
          dim3 cudaGridSize( gridXSize / 16 + ( gridXSize % 16 != 0 ),
                             gridYSize / 16 + ( gridYSize % 16 != 0 ) );
          
-         typedef ExplicitUpdaterTraverserUserData< RealType,
-                                           MeshFunctionType,
-                                           DifferentialOperator,
-                                           BoundaryCondition,
-                                           RightHandSide > UserData;
+         
          UserData userData;
          userData.time = time;
          userData.differentialOperator = &this->differentialOperatorPointer.template getData< Devices::Cuda >();
@@ -549,7 +554,7 @@ getExplicitUpdate( const RealType& time,
 
          TNL::Devices::Cuda::synchronizeDevice();
          int cudaErr;
-         _boundaryConditionsKernel< UserData, MeshType, RealType, IndexType >
+         _boundaryConditionsKernel< BoundaryEntitiesProcessor, UserData, MeshType, RealType, IndexType >
          <<< cudaGridSize, cudaBlockSize >>>
             ( &mesh.template getData< Devices::Cuda >(),
             userData );
@@ -563,7 +568,7 @@ getExplicitUpdate( const RealType& time,
           * Laplace operator
           */
          //cout << "Laplace operator ... " << endl;
-         _heatEquationKernel< UserData, MeshType, RealType, IndexType >
+         _heatEquationKernel< IntertiorEntitiesProcessor, UserData, MeshType, RealType, IndexType >
          <<< cudaGridSize, cudaBlockSize >>>
             ( &mesh.template getData< Devices::Cuda >(),
               userData );

tests/benchmarks/heat-equation-benchmark/Tuning/ExplicitUpdater.h

@@ -191,13 +191,10 @@ class ExplicitUpdater
                                               const IndexType& entityIndex,
                                               const typename MeshType::CoordinatesType& coordinates )
             {
+               typedef Functions::FunctionAdapter< MeshType, RightHandSide > FunctionAdapter;
                userData.fu[ entityIndex ] = 
                        ( *userData.differentialOperator )( mesh, userData.u, entityIndex, coordinates, userData.time );
-            
-               typedef Functions::FunctionAdapter< MeshType, RightHandSide > FunctionAdapter;
-               userData.fu[ entityIndex ] += 0.0;
-                  //FunctionAdapter::getValue( *userData.rightHandSide, entity, userData.time );               
-               
+                    //   + 0.0;
             }
             
       }; 

tests/benchmarks/heat-equation-benchmark/Tuning/tunning.h

@@ -34,7 +34,7 @@ struct Data
 
 #ifdef HAVE_CUDA
 
-template< typename UserData, typename Grid, typename Real, typename Index >
+template< typename BoundaryEntitiesProcessor, typename UserData, typename Grid, typename Real, typename Index >
 __global__ void _boundaryConditionsKernel( const Grid* grid,
                                            UserData userData )
 {
@@ -46,56 +46,55 @@ __global__ void _boundaryConditionsKernel( const Grid* grid,
    Coordinates coordinates( ( blockIdx.x ) * blockDim.x + threadIdx.x,
                              ( blockIdx.y ) * blockDim.y + threadIdx.y );
    
-   Index c = coordinates.y() * gridXSize + coordinates.x();
+   Index entityIndex = coordinates.y() * gridXSize + coordinates.x();
    
    if( coordinates.x() == 0 && coordinates.y() < gridYSize )
    {
-      //aux[ coordinates.y() * gridXSize ] = 0.0;
-      u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 ); //0.0;
+      //u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 );
+      BoundaryEntitiesProcessor::processEntity( *grid, userData, entityIndex, coordinates );
    }
    if( coordinates.x() == gridXSize - 1 && coordinates.y() < gridYSize )
    {
-      //aux[ coordinates.y() * gridXSize + gridYSize - 1 ] = 0.0;
-      //userData.u[ coordinates.y() * gridXSize + gridXSize - 1 ] = 0.0; //userData.u[ coordinates.y() * gridXSize + gridXSize - 1 ];
-      u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 ); //0.0;
+      //u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 );
+      BoundaryEntitiesProcessor::processEntity( *grid, userData, entityIndex, coordinates );
    }
    if( coordinates.y() == 0 && coordinates.x() < gridXSize )
    {
-      //aux[ coordinates.x() ] = 0.0; //userData.u[ coordinates.y() * gridXSize + 1 ];
-      //userData.u[ coordinates.x() ] = 0.0; //userData.u[ coordinates.y() * gridXSize + 1 ];
-      u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 ); //0.0;
+      //u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 );
+      BoundaryEntitiesProcessor::processEntity( *grid, userData, entityIndex, coordinates );
    }
    if( coordinates.y() == gridYSize -1  && coordinates.x() < gridXSize )
    {
-      //aux[ coordinates.y() * gridXSize + coordinates.x() ] = 0.0; //userData.u[ coordinates.y() * gridXSize + gridXSize - 1 ];      
-      //userData.u[ coordinates.y() * gridXSize + coordinates.x() ] = 0.0; //userData.u[ coordinates.y() * gridXSize + gridXSize - 1 ];      
-      u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 ); //0.0;
+      //u[ c ] = ( *bc )( *grid, u, c, coordinates, 0 );
+      BoundaryEntitiesProcessor::processEntity( *grid, userData, entityIndex, coordinates );
    }         
 }
 
-template< typename UserData, typename Grid, typename Real, typename Index >
+template< typename InteriorEntitiesProcessor, typename UserData, typename Grid, typename Real, typename Index >
 __global__ void _heatEquationKernel( const Grid* grid,
                                      UserData userData )
 {
-   Real* u = userData.u;
+   /*Real* u = userData.u;
    Real* fu = userData.fu;
-   const typename UserData::DifferentialOperatorType* op = userData.differentialOperator;
+   const typename UserData::DifferentialOperatorType* op = userData.differentialOperator;*/
 
    const Index& gridXSize = grid->getDimensions().x();
    const Index& gridYSize = grid->getDimensions().y();
    const Real& hx_inv = grid->template getSpaceStepsProducts< -2,  0 >();
    const Real& hy_inv = grid->template getSpaceStepsProducts<  0, -2 >();
    
-   typename Grid::CoordinatesType 
-      coordinates( blockIdx.x * blockDim.x + threadIdx.x, 
-                   blockIdx.y * blockDim.y + threadIdx.y );
+   using Coordinates = typename Grid::CoordinatesType;
+   Coordinates coordinates( blockIdx.x * blockDim.x + threadIdx.x, 
+                            blockIdx.y * blockDim.y + threadIdx.y );
+
    if( coordinates.x() > 0 && coordinates.x() < gridXSize - 1 &&
        coordinates.y() > 0 && coordinates.y() < gridYSize - 1 )
    {
-      const Index c = coordinates.y() * gridXSize + coordinates.x();
-      fu[ c ] = ( *op )( *grid, u, c, coordinates, 0 );
-         /*( ( userData.u[ c - 1 ]         - 2.0 * userData.u[ c ] + userData.u[ c + 1 ]         ) * hx_inv +
-                   ( userData.u[ c - gridXSize ] - 2.0 * userData.u[ c ] + userData.u[ c + gridXSize ] ) * hy_inv );*/
+      const Index entityIndex = coordinates.y() * gridXSize + coordinates.x();
+      InteriorEntitiesProcessor::processEntity( *grid, userData, entityIndex, coordinates );
+      
+         //fu[ entityIndex ] = ( *op )( *grid, userData.u, entityIndex, coordinates, userData.time ); // + 0.1;
+
    }  
 }
 

tests/benchmarks/heat-equation-benchmark/pure-c-rhs.h

@@ -82,6 +82,7 @@ __global__ void heatEquationKernel( const Real* u,
       const Index c = j * gridXSize + i;
       aux[ c ] = ( ( u[ c - 1 ]         - 2.0 * u[ c ] + u[ c + 1 ]         ) * hx_inv +
                    ( u[ c - gridXSize ] - 2.0 * u[ c ] + u[ c + gridXSize ] ) * hy_inv );
+      //aux[ c ] += 0.1;
       //aux[ c ] = ( ( __ldg( &u[ c - 1 ] ) - 2.0 * __ldg( &u[ c ] ) + __ldg( &u[ c + 1 ] ) ) * hx_inv +
       //                   ( __ldg( &u[ c - gridXSize ] ) - 2.0 * __ldg( &u[ c ] ) + __ldg( &u[ c + gridXSize ] ) ) * hy_inv );
    }