From 238c404b7ccc60c8a19e6a548f16f0c3926cf9bd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jakub=20Klinkovsk=C3=BD?= <klinkovsky@mmg.fjfi.cvut.cz> Date: Sun, 25 Aug 2019 15:08:15 +0200 Subject: [PATCH] Avoiding compiler warnings --- CMakeLists.txt | 5 +++ src/TNL/Assert.h | 16 +++++++++ src/TNL/Containers/NDArrayView.h | 5 +++ src/TNL/Containers/Subrange.h | 2 ++ src/TNL/Images/DicomImageInfo_impl.h | 4 ++- src/TNL/Images/DicomSeries_impl.h | 1 + .../DistributedMeshes/BufferEntitiesHelper.h | 16 ++++----- .../DistributedMeshes/CopyEntitiesHelper.h | 34 +++++++------------ .../DistributedGridIO_MeshFunction.h | 5 +-- src/TNL/Meshes/GridDetails/Grid1D.h | 3 ++ src/TNL/Meshes/GridDetails/Grid2D.h | 3 ++ src/TNL/Meshes/GridDetails/Grid3D.h | 3 ++ 12 files changed, 64 insertions(+), 33 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3bf9ce1448..cae38f0eca 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -102,6 +102,11 @@ if( CMAKE_CXX_COMPILER_ID STREQUAL "GNU" ) set( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-maybe-uninitialized" ) endif() +# disable false Clang warning: https://stackoverflow.com/q/57645872 +if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang") + set( CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-self-assign-overloaded" ) +endif() + # enable link time optimizations (but not in continuous integration) if( NOT DEFINED ENV{CI_JOB_NAME} ) if( CMAKE_CXX_COMPILER_ID STREQUAL "GNU" ) diff --git a/src/TNL/Assert.h b/src/TNL/Assert.h index cd81cdbd29..27f3b11b28 100644 --- a/src/TNL/Assert.h +++ b/src/TNL/Assert.h @@ -30,6 +30,14 @@ * Implemented by: Jakub Klinkovsky */ +// wrapper for nvcc pragma which disables warnings about __host__ __device__ +// functions: https://stackoverflow.com/q/55481202 +#ifdef __NVCC__ + #define TNL_NVCC_HD_WARNING_DISABLE #pragma hd_warning_disable +#else + #define TNL_NVCC_HD_WARNING_DISABLE +#endif + #if defined(NDEBUG) || defined(HAVE_MIC) // empty macros for optimized build @@ -280,6 +288,7 @@ cmpHelperOpFailure( const char* assertion, fatalFailure(); } +TNL_NVCC_HD_WARNING_DISABLE template< typename T1, typename T2 > __cuda_callable__ void cmpHelperTrue( const char* assertion, @@ -298,6 +307,7 @@ cmpHelperTrue( const char* assertion, expr1, "true", val1, true, "==" ); } +TNL_NVCC_HD_WARNING_DISABLE template< typename T1, typename T2 > __cuda_callable__ void cmpHelperFalse( const char* assertion, @@ -336,16 +346,22 @@ cmpHelper##op_name( const char* assertion, \ } // Implements the helper function for TNL_ASSERT_EQ +TNL_NVCC_HD_WARNING_DISABLE TNL_IMPL_CMP_HELPER_( EQ, == ); // Implements the helper function for TNL_ASSERT_NE +TNL_NVCC_HD_WARNING_DISABLE TNL_IMPL_CMP_HELPER_( NE, != ); // Implements the helper function for TNL_ASSERT_LE +TNL_NVCC_HD_WARNING_DISABLE TNL_IMPL_CMP_HELPER_( LE, <= ); // Implements the helper function for TNL_ASSERT_LT +TNL_NVCC_HD_WARNING_DISABLE TNL_IMPL_CMP_HELPER_( LT, < ); // Implements the helper function for TNL_ASSERT_GE +TNL_NVCC_HD_WARNING_DISABLE TNL_IMPL_CMP_HELPER_( GE, >= ); // Implements the helper function for TNL_ASSERT_GT +TNL_NVCC_HD_WARNING_DISABLE TNL_IMPL_CMP_HELPER_( GT, > ); #undef TNL_IMPL_CMP_HELPER_ diff --git a/src/TNL/Containers/NDArrayView.h b/src/TNL/Containers/NDArrayView.h index 54a020a64e..3e37de3725 100644 --- a/src/TNL/Containers/NDArrayView.h +++ b/src/TNL/Containers/NDArrayView.h @@ -69,6 +69,7 @@ public: // Copy-assignment does deep copy, just like regular array, but the sizes // must match (i.e. copy-assignment cannot resize). + TNL_NVCC_HD_WARNING_DISABLE __cuda_callable__ NDArrayView& operator=( const NDArrayView& other ) { @@ -79,7 +80,9 @@ public: } // Templated copy-assignment + TNL_NVCC_HD_WARNING_DISABLE template< typename OtherView > + __cuda_callable__ NDArrayView& operator=( const OtherView& other ) { static_assert( std::is_same< PermutationType, typename OtherView::PermutationType >::value, @@ -128,6 +131,7 @@ public: array = nullptr; } + TNL_NVCC_HD_WARNING_DISABLE __cuda_callable__ bool operator==( const NDArrayView& other ) const { @@ -137,6 +141,7 @@ public: return Algorithms::ArrayOperations< Device, Device >::compare( array, other.array, getStorageSize() ); } + TNL_NVCC_HD_WARNING_DISABLE __cuda_callable__ bool operator!=( const NDArrayView& other ) const { diff --git a/src/TNL/Containers/Subrange.h b/src/TNL/Containers/Subrange.h index c29d326ebe..0891185533 100644 --- a/src/TNL/Containers/Subrange.h +++ b/src/TNL/Containers/Subrange.h @@ -105,12 +105,14 @@ public: return i + begin; } + __cuda_callable__ bool operator==( const Subrange& other ) const { return begin == other.begin && end == other.end; } + __cuda_callable__ bool operator!=( const Subrange& other ) const { return ! (*this == other); diff --git a/src/TNL/Images/DicomImageInfo_impl.h b/src/TNL/Images/DicomImageInfo_impl.h index f1c68071c1..a393b0cc48 100644 --- a/src/TNL/Images/DicomImageInfo_impl.h +++ b/src/TNL/Images/DicomImageInfo_impl.h @@ -23,6 +23,8 @@ inline DicomImageInfo::DicomImageInfo( DicomHeader& dicomHeader ) : dicomHeader( dicomHeader ) { isObjectRetrieved = false; + width = 0; + height = 0; depth = 0; } @@ -105,4 +107,4 @@ inline int DicomImageInfo::getNumberOfSlices() } } // namespace Images -} // namespace TNL \ No newline at end of file +} // namespace TNL diff --git a/src/TNL/Images/DicomSeries_impl.h b/src/TNL/Images/DicomSeries_impl.h index 5e3a612b8e..350bf384bb 100644 --- a/src/TNL/Images/DicomSeries_impl.h +++ b/src/TNL/Images/DicomSeries_impl.h @@ -390,6 +390,7 @@ inline DicomHeader &DicomSeries::getHeader(int image) //check user argument if((image > 0) | (image <= dicomSeriesHeaders.getSize())) return *dicomSeriesHeaders.getElement(image); + throw std::out_of_range("image index out of range"); } inline bool DicomSeries::isDicomSeriesLoaded() diff --git a/src/TNL/Meshes/DistributedMeshes/BufferEntitiesHelper.h b/src/TNL/Meshes/DistributedMeshes/BufferEntitiesHelper.h index a20faf0a58..f4bc910a2d 100644 --- a/src/TNL/Meshes/DistributedMeshes/BufferEntitiesHelper.h +++ b/src/TNL/Meshes/DistributedMeshes/BufferEntitiesHelper.h @@ -27,9 +27,7 @@ template < typename MeshFunctionType, typename RealType=typename MeshFunctionType::MeshType::RealType, typename Device=typename MeshFunctionType::MeshType::DeviceType, typename Index=typename MeshFunctionType::MeshType::GlobalIndexType > -class BufferEntitiesHelper -{ -}; +class BufferEntitiesHelper; template < typename MeshFunctionType, @@ -53,14 +51,14 @@ class BufferEntitiesHelper< MeshFunctionType, MaskPointer, 1, RealType, Device, Index beginx=begin.x(); Index sizex=size.x(); - auto mesh = meshFunction.getMesh(); + auto* mesh = &meshFunction.getMeshPointer().template getData< Device >(); RealType* meshFunctionData = meshFunction.getData().getData(); const typename MaskPointer::ObjectType* mask( nullptr ); if( maskPointer ) mask = &maskPointer.template getData< Device >(); auto kernel = [tobuffer, mesh, buffer, isBoundary, meshFunctionData, mask, beginx ] __cuda_callable__ ( Index j ) { - typename MeshFunctionType::MeshType::Cell entity(mesh); + typename MeshFunctionType::MeshType::Cell entity(*mesh); entity.getCoordinates().x()=beginx+j; entity.refresh(); if( ! isBoundary || ! mask || ( *mask )[ entity.getIndex() ] ) @@ -99,7 +97,7 @@ class BufferEntitiesHelper< MeshFunctionType, MaskPointer, 2, RealType, Device, Index sizex=size.x(); Index sizey=size.y(); - auto mesh=meshFunction.getMesh(); + auto* mesh = &meshFunction.getMeshPointer().template getData< Device >(); RealType* meshFunctionData = meshFunction.getData().getData(); const typename MaskPointer::ObjectType* mask( nullptr ); if( maskPointer ) @@ -107,7 +105,7 @@ class BufferEntitiesHelper< MeshFunctionType, MaskPointer, 2, RealType, Device, auto kernel = [ tobuffer, mask, mesh, buffer, isBoundary, meshFunctionData, beginx, sizex, beginy] __cuda_callable__ ( Index i, Index j ) { - typename MeshFunctionType::MeshType::Cell entity(mesh); + typename MeshFunctionType::MeshType::Cell entity(*mesh); entity.getCoordinates().x() = beginx + i; entity.getCoordinates().y() = beginy + j; entity.refresh(); @@ -148,14 +146,14 @@ class BufferEntitiesHelper< MeshFunctionType, MaskPointer, 3, RealType, Device, Index sizey=size.y(); Index sizez=size.z(); - auto mesh=meshFunction.getMesh(); + auto* mesh = &meshFunction.getMeshPointer().template getData< Device >(); RealType * meshFunctionData=meshFunction.getData().getData(); const typename MaskPointer::ObjectType* mask( nullptr ); if( maskPointer ) mask = &maskPointer.template getData< Device >(); auto kernel = [ tobuffer, mesh, mask, buffer, isBoundary, meshFunctionData, beginx, sizex, beginy, sizey, beginz] __cuda_callable__ ( Index i, Index j, Index k ) { - typename MeshFunctionType::MeshType::Cell entity(mesh); + typename MeshFunctionType::MeshType::Cell entity(*mesh); entity.getCoordinates().x() = beginx + i; entity.getCoordinates().y() = beginy + j; entity.getCoordinates().z() = beginz + k; diff --git a/src/TNL/Meshes/DistributedMeshes/CopyEntitiesHelper.h b/src/TNL/Meshes/DistributedMeshes/CopyEntitiesHelper.h index df36543f36..cfc9c44bb0 100644 --- a/src/TNL/Meshes/DistributedMeshes/CopyEntitiesHelper.h +++ b/src/TNL/Meshes/DistributedMeshes/CopyEntitiesHelper.h @@ -20,15 +20,7 @@ namespace DistributedMeshes { template<typename MeshFunctionType, int dim=MeshFunctionType::getMeshDimension()> -class CopyEntitiesHelper -{ - public: - typedef typename MeshFunctionType::MeshType::CoordinatesType CoordinatesType; - static void Copy(MeshFunctionType &from, MeshFunctionType &to, CoordinatesType &fromBegin, CoordinatesType &toBegin, CoordinatesType &size) - { - } - -}; +class CopyEntitiesHelper; template<typename MeshFunctionType> @@ -43,12 +35,12 @@ class CopyEntitiesHelper<MeshFunctionType, 1> { auto toData=to.getData().getData(); auto fromData=from.getData().getData(); - auto fromMesh=from.getMesh(); - auto toMesh=to.getMesh(); + auto* fromMesh=&from.getMeshPointer().template getData< typename MeshFunctionType::MeshType::DeviceType >(); + auto* toMesh=&to.getMeshPointer().template getData< typename MeshFunctionType::MeshType::DeviceType >(); auto kernel = [fromData,toData, fromMesh, toMesh, fromBegin, toBegin] __cuda_callable__ ( Index i ) { - Cell fromEntity(fromMesh); - Cell toEntity(toMesh); + Cell fromEntity(*fromMesh); + Cell toEntity(*toMesh); toEntity.getCoordinates().x()=toBegin.x()+i; toEntity.refresh(); fromEntity.getCoordinates().x()=fromBegin.x()+i; @@ -75,12 +67,12 @@ class CopyEntitiesHelper<MeshFunctionType,2> { auto toData=to.getData().getData(); auto fromData=from.getData().getData(); - auto fromMesh=from.getMesh(); - auto toMesh=to.getMesh(); + auto* fromMesh=&from.getMeshPointer().template getData< typename MeshFunctionType::MeshType::DeviceType >(); + auto* toMesh=&to.getMeshPointer().template getData< typename MeshFunctionType::MeshType::DeviceType >(); auto kernel = [fromData,toData, fromMesh, toMesh, fromBegin, toBegin] __cuda_callable__ ( Index i, Index j ) { - Cell fromEntity(fromMesh); - Cell toEntity(toMesh); + Cell fromEntity(*fromMesh); + Cell toEntity(*toMesh); toEntity.getCoordinates().x()=toBegin.x()+i; toEntity.getCoordinates().y()=toBegin.y()+j; toEntity.refresh(); @@ -107,12 +99,12 @@ class CopyEntitiesHelper<MeshFunctionType,3> { auto toData=to.getData().getData(); auto fromData=from.getData().getData(); - auto fromMesh=from.getMesh(); - auto toMesh=to.getMesh(); + auto* fromMesh=&from.getMeshPointer().template getData< typename MeshFunctionType::MeshType::DeviceType >(); + auto* toMesh=&to.getMeshPointer().template getData< typename MeshFunctionType::MeshType::DeviceType >(); auto kernel = [fromData,toData, fromMesh, toMesh, fromBegin, toBegin] __cuda_callable__ ( Index i, Index j, Index k ) { - Cell fromEntity(fromMesh); - Cell toEntity(toMesh); + Cell fromEntity(*fromMesh); + Cell toEntity(*toMesh); toEntity.getCoordinates().x()=toBegin.x()+i; toEntity.getCoordinates().y()=toBegin.y()+j; toEntity.getCoordinates().z()=toBegin.z()+k; diff --git a/src/TNL/Meshes/DistributedMeshes/DistributedGridIO_MeshFunction.h b/src/TNL/Meshes/DistributedMeshes/DistributedGridIO_MeshFunction.h index 5620814a3d..3d7c6e5cdf 100644 --- a/src/TNL/Meshes/DistributedMeshes/DistributedGridIO_MeshFunction.h +++ b/src/TNL/Meshes/DistributedMeshes/DistributedGridIO_MeshFunction.h @@ -48,9 +48,10 @@ class DistributedGridIO< if(distrGrid==NULL) //not distributed { meshFunction.save(fileName); + return true; } - MeshType mesh=meshFunction.getMesh(); + const MeshType& mesh=meshFunction.getMesh(); PointType spaceSteps=mesh.getSpaceSteps(); PointType origin=mesh.getOrigin(); @@ -97,7 +98,7 @@ class DistributedGridIO< return true; } - MeshType mesh=meshFunction.getMesh(); + const MeshType& mesh=meshFunction.getMesh(); PointType spaceSteps=mesh.getSpaceSteps(); PointType origin=mesh.getOrigin(); diff --git a/src/TNL/Meshes/GridDetails/Grid1D.h b/src/TNL/Meshes/GridDetails/Grid1D.h index 9c4b172460..53b748c4e6 100644 --- a/src/TNL/Meshes/GridDetails/Grid1D.h +++ b/src/TNL/Meshes/GridDetails/Grid1D.h @@ -62,6 +62,9 @@ class Grid< 1, Real, Device, Index > : public Object Grid( const Index xSize ); + // empty destructor is needed only to avoid crappy nvcc warnings + ~Grid() {} + /** * \brief Returns type of grid Real (value), Device type and the type of Index. */ diff --git a/src/TNL/Meshes/GridDetails/Grid2D.h b/src/TNL/Meshes/GridDetails/Grid2D.h index ddb752e166..61f3c11c0e 100644 --- a/src/TNL/Meshes/GridDetails/Grid2D.h +++ b/src/TNL/Meshes/GridDetails/Grid2D.h @@ -62,6 +62,9 @@ class Grid< 2, Real, Device, Index > : public Object */ Grid( const Index xSize, const Index ySize ); + // empty destructor is needed only to avoid crappy nvcc warnings + ~Grid() {} + static String getType(); /** diff --git a/src/TNL/Meshes/GridDetails/Grid3D.h b/src/TNL/Meshes/GridDetails/Grid3D.h index 71b80cb6c0..67c752cb13 100644 --- a/src/TNL/Meshes/GridDetails/Grid3D.h +++ b/src/TNL/Meshes/GridDetails/Grid3D.h @@ -59,6 +59,9 @@ class Grid< 3, Real, Device, Index > : public Object Grid( const Index xSize, const Index ySize, const Index zSize ); + // empty destructor is needed only to avoid crappy nvcc warnings + ~Grid() {} + /** * \brief See Grid1D::getType(). */ -- GitLab