Workarounds for msvc 2015.

- MSVC isn't a fan of `decltype(...)::some_member` syntax. - WAR by aliasing the `decltype(...)` and doing `NewAlias::some_member` - Missing `template` keyword when rebinding pointer in `async/reduce.h` - Silence warning C4494 `declspec(allocator) used on non-pointer/ref type` - Bug in MSVC STL: microsoft/STL#696 - Disable async sort algo + test on MSVC - Triage. Looks like a bug in cudafe? See NVIDIA#1098. - Add pointer<T>::pointer_to(reference) - Required for C++11, hard compile error on MSVC. - Bring a definition of `atanh` into scope for complex number impl - Fix floating point literals be declared as floats instead of doubles - Fix `_Pragma` check for MSVC (MSVC uses `__pragma`) - Replace `std::remove_reference<T>::type&` with `std::add_lvalue_reference`. - Same behavior, and MSVC chokes on the other syntax when followed by `__host__`. - Remove constexpr markup from defaulted functions. - These are constexpr by default when possible, and the compilers were complaining about the markup in places. Other misc changes made for MSVC testing: - Bump C++ standard to C++14 (prepare for deprecation) - Add `-Xcudafe --display_error_number` to get useful diagnositics from cudafe.
alliepiper · Apr 8, 2020 · 1c9440e · 1c9440e
1 parent 5e64fb9
commit 1c9440e
Show file tree

Hide file tree

Showing 12 changed files with 82 additions and 50 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -64,7 +64,7 @@ endif ()
 add_definitions(-DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_${THRUST_DEVICE_SYSTEM})
 
 # Please note this also sets the default for the CUDA C++ version; see the comment below.
-set(CMAKE_CXX_STANDARD 11 CACHE STRING "The C++ version to be used.")
+set(CMAKE_CXX_STANDARD 14 CACHE STRING "The C++ version to be used.")
 set(CMAKE_CXX_EXTENSIONS OFF)
 
 message("-- C++ Standard version: ${CMAKE_CXX_STANDARD}")
@@ -179,6 +179,15 @@ if ("MSVC" STREQUAL "${CMAKE_CXX_COMPILER_ID}")
   # Disable warning about applying unary operator- to unsigned type.
   append_option_if_available("/wd4146" THRUST_CXX_WARNINGS)
 
+  # MSVC STL assumes that `allocator_traits`'s allocator will use raw pointers,
+  # and the `__DECLSPEC_ALLOCATOR` macro causes issues with thrust's universal
+  # allocators:
+  #   warning C4494: 'std::allocator_traits<_Alloc>::allocate' :
+  #      Ignoring __declspec(allocator) because the function return type is not
+  #      a pointer or reference
+  # See https://github.com/microsoft/STL/issues/696
+  append_option_if_available("/wd4494" THRUST_CXX_WARNINGS)
+
   set(THRUST_TREAT_FILE_AS_CXX "/TP")
 else ()
   append_option_if_available("-Werror" THRUST_CXX_WARNINGS)
@@ -245,6 +254,7 @@ if ("CUDA" STREQUAL "${THRUST_DEVICE_SYSTEM}")
   foreach (CXX_OPTION IN LISTS THRUST_CXX_WARNINGS)
     set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=${CXX_OPTION}")
   endforeach ()
+  set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcudafe --display_error_number")
 endif ()
 
 # For every public header, build a translation unit containing `#include <header>`

diff --git a/testing/async_sort.cu b/testing/async_sort.cu
@@ -1,6 +1,8 @@
 #include <thrust/detail/config.h>
 
-#if THRUST_CPP_DIALECT >= 2011 && !defined(THRUST_LEGACY_GCC)
+// Disabled on MSVC for GH issue #1098
+#if THRUST_CPP_DIALECT >= 2011 && !defined(THRUST_LEGACY_GCC) && \
+  THRUST_HOST_COMPILER != THRUST_HOST_COMPILER_MSVC
 
 #include <unittest/unittest.h>
 

diff --git a/thrust/async/sort.h b/thrust/async/sort.h
@@ -24,7 +24,9 @@
 #include <thrust/detail/cpp11_required.h>
 #include <thrust/detail/modern_gcc_required.h>
 
-#if THRUST_CPP_DIALECT >= 2011 && !defined(THRUST_LEGACY_GCC)
+// Disabled on MSVC for GH issue #1098
+#if THRUST_CPP_DIALECT >= 2011 && !defined(THRUST_LEGACY_GCC) && \
+  THRUST_HOST_COMPILER != THRUST_HOST_COMPILER_MSVC
 
 #include <thrust/detail/static_assert.h>
 #include <thrust/detail/select_system.h>

diff --git a/thrust/detail/complex/catrig.h b/thrust/detail/complex/catrig.h
@@ -604,6 +604,7 @@ complex<double> catanh(complex<double> z)
   ay = fabs(y);
 
   /* This helps handle many cases. */
+  using std::atanh;
   if (y == 0 && ax <= 1)
     return (complex<double>(atanh(x), y));
 

diff --git a/thrust/detail/complex/catrigf.h b/thrust/detail/complex/catrigf.h
@@ -392,8 +392,8 @@ __host__ __device__ inline
 complex<float> catanhf(complex<float> z)
 {
   float x, y, ax, ay, rx, ry;
-  const volatile float pio2_lo = 6.1232339957367659e-17; /*  0x11a62633145c07.0p-106 */
-  const float pio2_hi = 1.5707963267948966e0;/*  0x1921fb54442d18.0p-52 */
+  const volatile float pio2_lo = 6.1232339957367659e-17f; /*  0x11a62633145c07.0p-106 */
+  const float pio2_hi = 1.5707963267948966e0f;/*  0x1921fb54442d18.0p-52 */
 
 
   x = z.real();
@@ -422,7 +422,7 @@ complex<float> catanhf(complex<float> z)
     return (complex<float>(real_part_reciprocal(x, y),
 			   copysignf(pio2_hi + pio2_lo, y)));
 
-  const float SQRT_3_EPSILON = 5.9801995673e-4; /*  0x9cc471.0p-34 */
+  const float SQRT_3_EPSILON = 5.9801995673e-4f; /*  0x9cc471.0p-34 */
   if (ax < SQRT_3_EPSILON / 2 && ay < SQRT_3_EPSILON / 2) {
     raise_inexact();
     return (z);

diff --git a/thrust/detail/config/exec_check_disable.h b/thrust/detail/config/exec_check_disable.h
@@ -28,7 +28,11 @@
 #if defined(__CUDACC__) && !defined(__NVCOMPILER_CUDA__) && \
     !(defined(__CUDA__) && defined(__clang__))
 
+#if THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC
+#define __thrust_exec_check_disable__ __pragma("nv_exec_check_disable")
+#else // MSVC
 #define __thrust_exec_check_disable__ _Pragma("nv_exec_check_disable")
+#endif // MSVC
 
 #else
 

diff --git a/thrust/detail/execute_with_dependencies.h b/thrust/detail/execute_with_dependencies.h
@@ -189,7 +189,7 @@ struct execute_with_allocator_and_dependencies
         return std::move(dependencies);
     }
 
-    typename std::remove_reference<Allocator>::type&
+    typename std::add_lvalue_reference<Allocator>::type
     __host__
     get_allocator()
     {

diff --git a/thrust/detail/pointer.h b/thrust/detail/pointer.h
@@ -211,6 +211,12 @@ template<typename Element, typename Tag, typename Reference, typename Derived>
     __host__ __device__
     explicit operator bool() const;
     #endif
+
+    __host__ __device__
+    static derived_type pointer_to(typename thrust::detail::pointer_traits_detail::pointer_to_param<Element>::type r)
+    {
+      return thrust::detail::pointer_traits<derived_type>::pointer_to(r);
+    }
 }; // end pointer
 
 // Output stream operator

diff --git a/thrust/detail/type_traits.h b/thrust/detail/type_traits.h
@@ -49,17 +49,14 @@ namespace detail
      // to the C++14 operator(), but we'd like standard traits to interoperate
      // with our version when tag dispatching.
      #if THRUST_CPP_DIALECT >= 2011
-     constexpr integral_constant() = default;
+     integral_constant() = default;
 
-     constexpr integral_constant(integral_constant const&) = default;
+     integral_constant(integral_constant const&) = default;
 
-     #if THRUST_CPP_DIALECT >= 2014
-     constexpr // In C++11, constexpr makes member functions const.
-     #endif
      integral_constant& operator=(integral_constant const&) = default;
 
      constexpr __host__ __device__
-     integral_constant(std::integral_constant<T, v>) {}
+     integral_constant(std::integral_constant<T, v>) noexcept {}
      #endif
 
      THRUST_CONSTEXPR __host__ __device__ operator value_type() const THRUST_NOEXCEPT { return value; }

diff --git a/thrust/system/cuda/detail/async/copy.h b/thrust/system/cuda/detail/async/copy.h
@@ -229,6 +229,10 @@ auto async_copy_n(
 template <
   typename FromPolicy, typename ToPolicy
 , typename ForwardIt, typename OutputIt
+  // MSVC2015 WAR: doesn't like decltype(...)::value in superclass definition
+, typename IsH2DCopy = decltype(is_host_to_device_copy(
+    std::declval<FromPolicy const&>()
+  , std::declval<ToPolicy const&>()))
 >
 struct is_buffered_trivially_relocatable_host_to_device_copy
   : thrust::integral_constant<
@@ -239,12 +243,7 @@ struct is_buffered_trivially_relocatable_host_to_device_copy
             typename iterator_traits<ForwardIt>::value_type
           , typename iterator_traits<OutputIt>::value_type
           >::value
-      && decltype(
-           is_host_to_device_copy(
-             std::declval<FromPolicy const&>()
-           , std::declval<ToPolicy const&>()
-           )
-         )::value
+      && IsH2DCopy::value
     >
 {};
 
@@ -334,6 +333,10 @@ auto async_copy_n(
 template <
   typename FromPolicy, typename ToPolicy
 , typename ForwardIt, typename OutputIt
+  // MSVC2015 WAR: doesn't like decltype(...)::value in superclass definition
+, typename IsD2HCopy = decltype(is_device_to_host_copy(
+    std::declval<FromPolicy const&>()
+  , std::declval<ToPolicy const&>()))
 >
 struct is_buffered_trivially_relocatable_device_to_host_copy
   : thrust::integral_constant<
@@ -344,12 +347,7 @@ struct is_buffered_trivially_relocatable_device_to_host_copy
             typename iterator_traits<ForwardIt>::value_type
           , typename iterator_traits<OutputIt>::value_type
           >::value
-      && decltype(
-           is_device_to_host_copy(
-             std::declval<FromPolicy const&>()
-           , std::declval<ToPolicy const&>()
-           )
-         )::value
+      && IsD2HCopy::value
     >
 {};
 

diff --git a/thrust/system/cuda/detail/async/reduce.h b/thrust/system/cuda/detail/async/reduce.h
@@ -75,7 +75,7 @@ auto async_reduce_n(
 
   using pointer
     = typename thrust::detail::allocator_traits<decltype(device_alloc)>::
-      rebind_traits<U>::pointer;
+      template rebind_traits<U>::pointer;
 
   unique_eager_future_promise_pair<U, pointer> fp;
 

diff --git a/thrust/system/cuda/detail/cross_system.h b/thrust/system/cuda/detail/cross_system.h
@@ -115,88 +115,100 @@ namespace cuda_cub {
     )
   )
 
-  template <typename ExecutionPolicy0, typename ExecutionPolicy1>
+  template <typename ExecutionPolicy0,
+            typename ExecutionPolicy1,
+            // MSVC2015 WAR: put decltype here instead of in trailing return type
+            typename Direction =
+              decltype(direction_of_copy(std::declval<ExecutionPolicy0>(),
+                                         std::declval<ExecutionPolicy1>()))>
   THRUST_CONSTEXPR __host__ __device__
   auto is_device_to_host_copy(
     ExecutionPolicy0 const& exec0
   , ExecutionPolicy1 const& exec1
   )
     noexcept -> 
       thrust::detail::integral_constant<
-        bool
-      ,    cudaMemcpyDeviceToHost
-        == decltype(direction_of_copy(exec0, exec1))::value
+        bool, cudaMemcpyDeviceToHost == Direction::value
       >
   {
     return {};
   }
 
-  template <typename ExecutionPolicy>
+  template <typename ExecutionPolicy,
+            // MSVC2015 WAR: put decltype here instead of in trailing return type
+            typename Direction =
+              decltype(direction_of_copy(std::declval<ExecutionPolicy>()))>
   THRUST_CONSTEXPR __host__ __device__
   auto is_device_to_host_copy(ExecutionPolicy const& exec)
     noexcept -> 
       thrust::detail::integral_constant<
-        bool
-      ,    cudaMemcpyDeviceToHost
-        == decltype(direction_of_copy(exec))::value
+        bool, cudaMemcpyDeviceToHost == Direction::value
       >
   {
     return {};
   }
 
-  template <typename ExecutionPolicy0, typename ExecutionPolicy1>
+  template <typename ExecutionPolicy0,
+            typename ExecutionPolicy1,
+            // MSVC2015 WAR: put decltype here instead of in trailing return type
+            typename Direction =
+              decltype(direction_of_copy(std::declval<ExecutionPolicy0>(),
+                                         std::declval<ExecutionPolicy1>()))>
   THRUST_CONSTEXPR __host__ __device__
   auto is_host_to_device_copy(
     ExecutionPolicy0 const& exec0
   , ExecutionPolicy1 const& exec1
   )
     noexcept -> 
       thrust::detail::integral_constant<
-        bool
-      ,    cudaMemcpyHostToDevice
-        == decltype(direction_of_copy(exec0, exec1))::value
+        bool, cudaMemcpyHostToDevice == Direction::value
       >
   {
     return {};
   }
 
-  template <typename ExecutionPolicy>
+  template <typename ExecutionPolicy,
+            // MSVC2015 WAR: put decltype here instead of in trailing return type
+            typename Direction =
+              decltype(direction_of_copy(std::declval<ExecutionPolicy>()))>
   THRUST_CONSTEXPR __host__ __device__
   auto is_host_to_device_copy(ExecutionPolicy const& exec)
     noexcept -> 
       thrust::detail::integral_constant<
-        bool
-      ,    cudaMemcpyHostToDevice
-        == decltype(direction_of_copy(exec))::value
+        bool, cudaMemcpyHostToDevice == Direction::value
       >
   {
     return {};
   }
 
-  template <typename ExecutionPolicy0, typename ExecutionPolicy1>
+  template <typename ExecutionPolicy0,
+            typename ExecutionPolicy1,
+            // MSVC2015 WAR: put decltype here instead of in trailing return type
+            typename Direction =
+              decltype(direction_of_copy(std::declval<ExecutionPolicy0>(),
+                                         std::declval<ExecutionPolicy1>()))>
   THRUST_CONSTEXPR __host__ __device__
   auto is_device_to_device_copy(
     ExecutionPolicy0 const& exec0
   , ExecutionPolicy1 const& exec1
   )
     noexcept -> 
       thrust::detail::integral_constant<
-        bool
-      ,    cudaMemcpyDeviceToDevice
-        == decltype(direction_of_copy(exec0, exec1))::value
+        bool, cudaMemcpyDeviceToDevice == Direction::value
       >
   {
     return {};
   }
 
-  template <typename ExecutionPolicy>
+  template <typename ExecutionPolicy,
+            // MSVC2015 WAR: put decltype here instead of in trailing return type
+            typename Direction =
+              decltype(direction_of_copy(std::declval<ExecutionPolicy>()))>
   THRUST_CONSTEXPR __host__ __device__
   auto is_device_to_device_copy(ExecutionPolicy const& exec)
     noexcept -> 
       thrust::detail::integral_constant<
-        bool
-      ,    cudaMemcpyDeviceToDevice
-        == decltype(direction_of_copy(exec))::value
+        bool, cudaMemcpyDeviceToDevice == Direction::value
       >
   {
     return {};