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

Finishing simple code explaining bad CUDA performance.

parent c0705f62
Loading
Loading
Loading
Loading
+4 −4
Original line number Diff line number Diff line
@@ -23,7 +23,7 @@
#include <mesh/grids/tnlGridEntityGetter.h>
#include <mesh/grids/tnlNeighbourGridEntityGetter.h>

#include <../tests/benchmarks/heat-equation-benchmark/TestGridEntity.h>
//#include <../tests/benchmarks/heat-equation-benchmark/TestGridEntity.h>
//#include <../tests/benchmarks/heat-equation-benchmark/TestNeighbourGridEntityGetter2D_impl.h>

template< typename Real,
@@ -53,11 +53,11 @@ class tnlGrid< 2, Real, Device, Index > : public tnlObject
   typedef MeshEntity< 0 > Vertex;
   
   
   template< int EntityDimensions, 
   /*template< int EntityDimensions, 
             typename Config = tnlGridEntityNoStencilStorage >//CrossStencilStorage< 1 > >
   using TestMeshEntity = TestGridEntity< ThisType, EntityDimensions, Config >;
   using TestMeshEntity = TestGridEntity< ThisType, EntityDimensions >;
   typedef TestMeshEntity< meshDimensions, tnlGridEntityCrossStencilStorage< 1 > > TestCell;
    
    */
   
   static constexpr int getMeshDimensions() { return meshDimensions; };

+1 −1
Original line number Diff line number Diff line
@@ -363,7 +363,7 @@ heatEquationTemplatedCompact( const GridType grid,
   //GridEntity entity( grid, coordinates, entityOrientation, entityBasis );
   //printf( "size = %d ", sizeof( GridEntity ) );
   //entity.refresh();
   typename GridType::TestCell entity( grid, coordinates, entityOrientation, entityBasis );
   //typename GridType::TestCell entity( grid, coordinates, entityOrientation, entityBasis );
   
   const IndexType tidX = begin.x() + ( gridXIdx * tnlCuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;
   const IndexType tidY = begin.y() + ( gridYIdx * tnlCuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y;
+7 −10
Original line number Diff line number Diff line
@@ -31,8 +31,7 @@ class TestNeighbourGridEntitiesStorage
};

template< typename Grid,          
          int EntityDimensions,
          typename Config >
          int EntityDimensions >
class TestGridEntity
{
};
@@ -42,9 +41,8 @@ template< int Dimensions,
          typename Real,
          typename Device,
          typename Index,          
          int EntityDimensions,
          typename Config >
class TestGridEntity< tnlGrid< Dimensions, Real, Device, Index >, EntityDimensions, Config >
          int EntityDimensions >
class TestGridEntity< tnlGrid< Dimensions, Real, Device, Index >, EntityDimensions >
{
   public:
      static const int entityDimensions = EntityDimensions;
@@ -56,9 +54,8 @@ class TestGridEntity< tnlGrid< Dimensions, Real, Device, Index >, EntityDimensio
template< int Dimensions,
          typename Real,
          typename Device,
          typename Index,
          typename Config >
class TestGridEntity< tnlGrid< Dimensions, Real, Device, Index >, Dimensions, Config >
          typename Index >
class TestGridEntity< tnlGrid< Dimensions, Real, Device, Index >, Dimensions >
{
   public:
      
@@ -68,7 +65,7 @@ class TestGridEntity< tnlGrid< Dimensions, Real, Device, Index >, Dimensions, Co
      typedef typename GridType::IndexType IndexType;
      typedef typename GridType::CoordinatesType CoordinatesType;
      typedef typename GridType::VertexType VertexType;
      typedef Config ConfigType;
      //typedef Config ConfigType;
      
      static const int meshDimensions = GridType::meshDimensions;
      
@@ -81,7 +78,7 @@ class TestGridEntity< tnlGrid< Dimensions, Real, Device, Index >, Dimensions, Co
      
      typedef tnlStaticVector< meshDimensions, IndexType > EntityOrientationType;
      typedef tnlStaticVector< meshDimensions, IndexType > EntityBasisType;
      typedef TestGridEntity< GridType, entityDimensions, Config > ThisType;
      typedef TestGridEntity< GridType, entityDimensions > ThisType;
      typedef TestNeighbourGridEntitiesStorage< ThisType > NeighbourGridEntitiesStorageType;
      
      __cuda_callable__ inline
+51 −586

File changed.

Preview size limit exceeded, changes collapsed.

+1 −1

File changed.

Contains only whitespace changes.