diff --git a/src/TNL/Functions/MeshFunction_impl.h b/src/TNL/Functions/MeshFunction_impl.h index 7164170f1ab51a20bcdeaeb9011cb97bb97b4590..8704fc6868b364c444d086916aa870556f41ce83 100644 --- a/src/TNL/Functions/MeshFunction_impl.h +++ b/src/TNL/Functions/MeshFunction_impl.h @@ -312,7 +312,7 @@ MeshFunction< Mesh, MeshEntityDimension, Real >:: getValue( const EntityType& meshEntity ) const { static_assert( EntityType::getEntityDimension() == MeshEntityDimension, "Calling with wrong EntityType -- entity dimensions do not match." ); - return this->data.getValue( meshEntity.getIndex() ); + return this->data.getElement( meshEntity.getIndex() ); } template< typename Mesh, @@ -325,7 +325,7 @@ setValue( const EntityType& meshEntity, const RealType& value ) { static_assert( EntityType::getEntityDimension() == MeshEntityDimension, "Calling with wrong EntityType -- entity dimensions do not match." ); - this->data.setValue( meshEntity.getIndex(), value ); + this->data.setElement( meshEntity.getIndex(), value ); } template< typename Mesh, diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_impl.h b/src/TNL/Meshes/GridDetails/GridTraverser_impl.h index 4443e161c3d3296f7a7924d927f00a3581248cec..b61837ae9e8673038c1ac263b0c94f37913dbdd5 100644 --- a/src/TNL/Meshes/GridDetails/GridTraverser_impl.h +++ b/src/TNL/Meshes/GridDetails/GridTraverser_impl.h @@ -381,23 +381,6 @@ GridTraverser2D( coordinates.x() = begin.x() + Devices::Cuda::getGlobalThreadIdx_x( gridIdx ); coordinates.y() = begin.y() + Devices::Cuda::getGlobalThreadIdx_y( gridIdx ); - /*if( processOnlyBoundaryEntities && - ( GridEntity::getMeshDimension() == 2 || GridEntity::getMeshDimension() == 0 ) ) - { - if( coordinates.x() == begin.x() || coordinates.x() == end.x() || - coordinates.y() == begin.y() || coordinates.y() == end.y() ) - { - GridEntity entity( *grid, coordinates, gridEntityParameters... ); - entity.refresh(); - EntitiesProcessor::processEntity - ( *grid, - *userData, - entity ); - } - return; - }*/ - - if( coordinates <= end ) { GridEntity entity( *grid, coordinates, gridEntityParameters... ); @@ -504,7 +487,7 @@ processEntities( { #ifdef HAVE_CUDA if( processOnlyBoundaryEntities && - ( GridEntity::getMeshDimension() == 2 || GridEntity::getMeshDimension() == 0 ) ) + ( GridEntity::getEntityDimension() == 2 || GridEntity::getEntityDimension() == 0 ) ) { dim3 cudaBlockSize( 256 ); dim3 cudaBlocksCountAlongX, cudaGridsCountAlongX, @@ -998,7 +981,7 @@ processEntities( { #ifdef HAVE_CUDA if( processOnlyBoundaryEntities && - ( GridEntity::getMeshDimension() == 3 || GridEntity::getMeshDimension() == 0 ) ) + ( GridEntity::getEntityDimension() == 3 || GridEntity::getEntityDimension() == 0 ) ) { dim3 cudaBlockSize( 16, 16 ); const IndexType entitiesAlongX = end.x() - begin.x() + 1; diff --git a/src/TNL/ParallelFor.h b/src/TNL/ParallelFor.h index b1e81ca97bf66576875aa1d0dcb34ecc0298d5df..7002ecafe0fd7cc97271bf2fd885f3bfb50ac331 100644 --- a/src/TNL/ParallelFor.h +++ b/src/TNL/ParallelFor.h @@ -76,6 +76,8 @@ struct ParallelFor< Devices::Cuda > Devices::Cuda::synchronizeDevice(); ParallelForKernel<<< gridSize, blockSize >>>( start, end, f, args... ); + cudaDeviceSynchronize(); + TNL_CHECK_CUDA_DEVICE; } #endif }