diff --git a/CMakeLists.txt b/CMakeLists.txt
index 3bf9ce1448180b21a792c4f3960829d5e800de71..cae38f0eca397d7f7daffac063d9ab224ff241dc 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 cd81cdbd298d34b5c8ea77c6061eb53b44c324ac..27f3b11b28ed46a9741c3593c573a243f1e0a81d 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 54a020a64e5b5130400c6e954ae92840cb0fe1df..3e37de372521cddb69db001ec05f6b238d644e15 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 c29d326ebe144ab3a13bd089bb9b0c71a41e8fde..08911855337817712c960047bf9688a7f134a752 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 f1c68071c13c573288d82277c67e1059f3450f67..a393b0cc48e800b56d63faff336434b71a516748 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 5e3a612b8e4709f5258c3eac772a8ec3b252f212..350bf384bbadae0eeb1045e1553010953f6a9390 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 a20faf0a589f4414c791a62585228a0dfc529bf9..f4bc910a2d04a99eb1356baa8545705c4a174cac 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 df36543f360564fce8905d97d198c2b40a2490a3..cfc9c44bb0e04982482787733076c71495a999b3 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 5620814a3df51713c431fadd9192254cce1064a8..3d7c6e5cdf72f292b76571b6eda2acf8dcda3da6 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 9c4b1724607c455a540f08c4f546272ee22a1318..53b748c4e6a7893d8fb645ffdcc271c80c8c5b58 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 ddb752e166041339943c566cb10d3d3261e43d5f..61f3c11c0e0684c892a02971bfc5a2d2df67979c 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 71b80cb6c0dc50e7326ea2e92a40d11c18c258a7..67c752cb13700628abcac76de2f066473baf26a7 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().
*/