More corrections needed to suppress false positive host / device warnings.

vtkm/cont/ArrayHandlePermutation.h

@@ -36,6 +36,7 @@ class ArrayPortalPermutationExec
 public:
   typedef typename ValuePortalType::ValueType ValueType;
 
+  VTKM_SUPPRESS_EXEC_WARNINGS
   VTKM_EXEC_CONT_EXPORT
   ArrayPortalPermutationExec( )
     : IndexPortal(), ValuePortal() {  }
@@ -63,12 +64,14 @@ public:
     return this->IndexPortal.GetNumberOfValues();
   }
 
+  VTKM_SUPPRESS_EXEC_WARNINGS
   VTKM_EXEC_EXPORT
   ValueType Get(vtkm::Id index) const {
     vtkm::Id permutedIndex = this->IndexPortal.Get(index);
     return this->ValuePortal.Get(permutedIndex);
   }
 
+  VTKM_SUPPRESS_EXEC_WARNINGS
   VTKM_EXEC_EXPORT
   void Set(vtkm::Id index, const ValueType &value) const {
     vtkm::Id permutedIndex = this->IndexPortal.Get(index);

vtkm/cont/Field.h

@@ -41,12 +41,14 @@ namespace internal {
 struct RangeMin
 {
   template<typename T>
+  VTKM_EXEC_EXPORT
   T operator()(const T& a, const T& b)const { return vtkm::Min(a,b); }
 };
 
 struct RangeMax
 {
   template<typename T>
+  VTKM_EXEC_EXPORT
   T operator()(const T& a, const T& b)const { return vtkm::Max(a,b); }
 };
 

vtkm/exec/arg/ThreadIndicesTopologyMap.h

@@ -72,6 +72,7 @@ public:
   typedef typename ConnectivityType::IndicesType IndicesFromType;
   typedef typename ConnectivityType::CellShapeTag CellShapeTag;
 
+  VTKM_SUPPRESS_EXEC_WARNINGS
   template<typename OutToInArrayType, typename VisitArrayType>
   VTKM_EXEC_EXPORT
   ThreadIndicesTopologyMap(vtkm::Id threadIndex,

vtkm/exec/internal/VecFromPortal.h

@@ -67,6 +67,7 @@ public:
     }
   }
 
+  VTKM_SUPPRESS_EXEC_WARNINGS
   VTKM_EXEC_EXPORT
   ComponentType operator[](vtkm::IdComponent index) const
   {

vtkm/exec/internal/VecFromPortalPermute.h

@@ -65,6 +65,7 @@ public:
     }
   }
 
+  VTKM_SUPPRESS_EXEC_WARNINGS
   VTKM_EXEC_EXPORT
   ComponentType operator[](vtkm::IdComponent index) const
   {

vtkm/internal/Assume.h

@@ -44,16 +44,14 @@
   } while (false) /* do-while prevents extra semicolon warnings */
 
 // VTKM_ASSUME_IMPL is compiler-specific:
-#if defined(VTKM_MSVC)
-//Currently NVCC/VS can generate invalid PTX code when it encounters __assume.
-//So while this issue is being resolved we will disable VTKM_ASSUME when inside
-//CUDA code being built by Visual Studio
-#  if defined(__CUDA_ARCH__)
-#   define VTKM_ASSUME_IMPL(cond) do {} while (false) /* no-op */
-#  else
-#   define VTKM_ASSUME_IMPL(cond) __assume(cond)
-# endif
+#if defined(__CUDA_ARCH__)
+  //For all versions of CUDA this is a no-op while we look
+  //for a CUDA asm snippet that replicates this kind of behavior
+  #define VTKM_ASSUME_IMPL(cond) do {} while (false) /* no-op */
+#else
 
+#if defined(VTKM_MSVC)
+# define VTKM_ASSUME_IMPL(cond) __assume(cond)
 #elif defined(VTKM_ICC)
 # define VTKM_ASSUME_IMPL(cond) __assume(cond)
 #elif defined(VTKM_GCC) && ( __GNUC__ > 4 || (__GNUC__ == 4 && __GNUC_MINOR__ >= 5) )
@@ -65,4 +63,6 @@
 # define VTKM_ASSUME_IMPL(cond) do {} while (false) /* no-op */
 #endif
 
+#endif
+
 #endif // vtk_m_internal_Assume_h

vtkm/internal/ExportMacros.h

@@ -29,9 +29,11 @@
 #ifdef VTKM_CUDA
 #define VTKM_EXEC_EXPORT inline __device__ __host__
 #define VTKM_EXEC_CONT_EXPORT inline __device__ __host__
-#define VTKM_SUPPRESS_EXEC_WARNINGS \
-  #pragma hd_warning_disable \
-  #pragma nv_exec_check_disable
+#if __CUDAVER__ >= 75000
+#  define VTKM_SUPPRESS_EXEC_WARNINGS #pragma nv_exec_check_disable
+#else
+#  define VTKM_SUPPRESS_EXEC_WARNINGS #pragma hd_warning_disable
+#endif
 #define VTKM_EXEC_CONSTANT_EXPORT __device__ __constant__
 #else
 #define VTKM_EXEC_EXPORT inline

vtkm/worklet/WorkletMapTopology.h

@@ -123,6 +123,7 @@ public:
 
   /// Topology map worklets use topology map indices.
   ///
+  VTKM_SUPPRESS_EXEC_WARNINGS
   template<typename T, typename OutToInArrayType, typename VisitArrayType, typename InputDomainType>
   VTKM_EXEC_EXPORT
   vtkm::exec::arg::ThreadIndicesTopologyMap<InputDomainType>

vtkm/worklet/internal/WorkletBase.h

@@ -282,6 +282,7 @@ public:
   /// Worklet types can add additional indices by returning different object
   /// types.
   ///
+  VTKM_SUPPRESS_EXEC_WARNINGS
   template<typename T, typename OutToInArrayType, typename VisitArrayType, typename InputDomainType>
   VTKM_EXEC_EXPORT
   vtkm::exec::arg::ThreadIndicesBasic