Fix issues so all host/device combinations build and pass tests.

- Add more metadata to mock specializations for testing iterator in testing/copy.cu. - Add missing include to shuffle unit test. - Specialize wrapped_function for void return types. - MSVC is not a fan of the pattern `return static_cast<void>(expr);`. - Replace deprecated `tbb/tbb_thread.h` with `<thread>`. - Fix overcounting of initial value in tbb scans. - Apparently reverse_join may be called before operator() - Use `thrust::advance` instead of `+=` for generic iterators. - Wrap the OMP flags in -Xcompiler for NVCC - Extend ASSERT_STATIC_ASSERT skip for HOST=OMP, too - Add missing header caught by tbb.cuda configs. - Fix 'unsafe API' warnings in examples on MSVC: s/fopen/fstream/
NVIDIA · Jun 1, 2020 · 4d4fae1 · 4d4fae1
1 parent e478243
commit 4d4fae1
Show file tree

Hide file tree

Showing 12 changed files with 167 additions and 87 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -236,6 +236,9 @@ if ("MSVC" STREQUAL "${CMAKE_CXX_COMPILER_ID}")
   # object files:
   append_option_if_available("/bigobj" THRUST_CXX_WARNINGS)
 
+  # "Oh right, this is Visual Studio."
+  add_compile_definitions("NOMINMAX")
+
   set(THRUST_TREAT_FILE_AS_CXX "/TP")
 else ()
   append_option_if_available("-Werror" THRUST_CXX_WARNINGS)
@@ -679,6 +682,13 @@ foreach (THRUST_EXAMPLE_SOURCE IN LISTS THRUST_EXAMPLES)
     endif ()
   endif ()
 
+  if ("MSVC" STREQUAL "${CMAKE_CXX_COMPILER_ID}")
+    # Some examples use unsafe APIs (e.g. fopen) that MSVC will complain about
+    # unless this is set:
+    set_target_properties(${THRUST_EXAMPLE}
+      PROPERTIES COMPILE_DEFINITIONS "_CRT_SECURE_NO_WARNINGS")
+  endif()
+
   add_test(NAME ${THRUST_EXAMPLE}
     COMMAND ${CMAKE_COMMAND}
       -DTHRUST_EXAMPLE=${THRUST_EXAMPLE}

diff --git a/examples/discrete_voronoi.cu b/examples/discrete_voronoi.cu
@@ -4,10 +4,10 @@
 #include <thrust/extrema.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/random.h>
-#include <iostream>
 
+#include <iostream>
 #include <iomanip>
-#include <stdio.h>
+#include <fstream>
 #include <cmath>
 
 #include "include/timer.h"
@@ -135,21 +135,25 @@ void generate_random_sites(thrust::host_vector<int> &t, int Nb, int m, int n)
 //Export the tab to PGM image format
 void vector_to_pgm(thrust::host_vector<int> &t, int m, int n, const char *out)
 {
-    FILE *f;
+    assert(t.size() == m * n && "Vector size does not match image dims.");
 
-    f=fopen(out,"w+t");
-    fprintf(f,"P2\n");
-    fprintf(f,"%d %d\n 253\n",m,n);
+    std::fstream f(out, std::fstream::out);
+    f << "P2\n";
+    f << m << " " << n << "\n";
+    f << "253\n";
+
+    //Hash function to map values to [0,255]
+    auto to_grey_level = [&f](int in_value) -> int
+    {
+        return (71 * in_value) % 253;
+    };
 
-    for(int j = 0; j < n ; j++)
+    for (int value : t)
     {
-        for(int i = 0; i < m ; i++)
-        {
-            fprintf(f,"%d ",(int)(71*t[j*m+i])%253); //Hash function to map values to [0,255]
-        }
+      f << to_grey_level(value) << " ";
     }
-    fprintf(f,"\n");
-    fclose(f);
+    f << "\n";
+    f.close();
 }
 
 /************Main Jfa loop********************/

diff --git a/testing/copy.cu b/testing/copy.cu
@@ -724,6 +724,8 @@ struct only_set_when_expected_it
     __host__ __device__ only_set_when_expected_it operator*() const { return *this; }
     template<typename Difference>
     __host__ __device__ only_set_when_expected_it operator+(Difference) const { return *this; }
+    template<typename Difference>
+    __host__ __device__ only_set_when_expected_it operator+=(Difference) const { return *this; }
     template<typename Index>
     __host__ __device__ only_set_when_expected_it operator[](Index) const { return *this; }
 
@@ -739,11 +741,20 @@ struct only_set_when_expected_it
 
 namespace thrust
 {
+namespace detail
+{
+// We need this type to pass as a non-const ref for unary_transform_functor
+// to compile:
+template <>
+struct is_non_const_reference<only_set_when_expected_it> : thrust::true_type {};
+}
+
 template<>
 struct iterator_traits<only_set_when_expected_it>
 {
     typedef long long value_type;
     typedef only_set_when_expected_it reference;
+    typedef thrust::random_access_device_iterator_tag iterator_category;
 };
 }
 

diff --git a/testing/shuffle.cu b/testing/shuffle.cu
@@ -2,6 +2,7 @@
 
 #if THRUST_CPP_DIALECT >= 2011
 #include <thrust/random.h>
+#include <thrust/sequence.h>
 #include <thrust/shuffle.h>
 #include <thrust/sort.h>
 #include <unittest/unittest.h>

diff --git a/testing/unittest_static_assert.cu b/testing/unittest_static_assert.cu
@@ -22,7 +22,7 @@ struct static_assertion
 template<typename V>
 void TestStaticAssertAssert()
 {
-#if THRUST_DEVICE_SYSTEM != THRUST_DEVICE_SYSTEM_OMP
+#if THRUST_DEVICE_SYSTEM != THRUST_DEVICE_SYSTEM_OMP && THRUST_HOST_SYSTEM != THRUST_HOST_SYSTEM_OMP
     V test(10);
     ASSERT_STATIC_ASSERT(thrust::generate(test.begin(), test.end(), static_assertion<int>()));
 #endif

diff --git a/thrust/cmake/thrust-config.cmake b/thrust/cmake/thrust-config.cmake
@@ -551,6 +551,16 @@ macro(_thrust_find_TBB required)
   endif()
 endmacro()
 
+# Wrap the OpenMP flags for CUDA targets
+function(thrust_fixup_omp_target omp_target)
+  get_target_property(opts ${omp_target} INTERFACE_COMPILE_OPTIONS)
+  if (opts MATCHES "\\$<\\$<COMPILE_LANGUAGE:CXX>:([^>]*)>")
+    target_compile_options(${omp_target} INTERFACE
+      $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=${CMAKE_MATCH_1}>
+    )
+  endif()
+endfunction()
+
 # This must be a macro instead of a function to ensure that backends passed to
 # find_package(Thrust COMPONENTS [...]) have their full configuration loaded
 # into the current scope. This provides at least some remedy for CMake issue
@@ -568,6 +578,7 @@ macro(_thrust_find_OMP required)
     )
 
     if (TARGET OpenMP::OpenMP_CXX)
+      thrust_fixup_omp_target(OpenMP::OpenMP_CXX)
       thrust_set_OMP_target(OpenMP::OpenMP_CXX)
     else()
       thrust_debug("OpenMP::OpenMP_CXX not found!" internal)

diff --git a/thrust/detail/function.h b/thrust/detail/function.h
@@ -24,80 +24,137 @@ namespace thrust
 namespace detail
 {
 
-
-template<typename Function, typename Result>
-  struct wrapped_function
+template <typename Function, typename Result>
+struct wrapped_function
 {
   // mutable because Function::operator() might be const
   mutable Function m_f;
 
   inline __host__ __device__
   wrapped_function()
-    : m_f()
+      : m_f()
   {}
 
   inline __host__ __device__
-  wrapped_function(const Function &f)
-    : m_f(f)
+  wrapped_function(const Function& f)
+      : m_f(f)
   {}
 
   __thrust_exec_check_disable__
-  template<typename Argument>
+  template <typename Argument>
   inline __host__ __device__
-    Result operator()(Argument &x) const
+  Result operator()(Argument& x) const
   {
-    // we static cast to Result to handle void Result without error
-    // in case Function's result is non-void
     return static_cast<Result>(m_f(thrust::raw_reference_cast(x)));
   }
 
   __thrust_exec_check_disable__
-  template<typename Argument>
-    inline __host__ __device__ Result operator()(const Argument &x) const
+  template <typename Argument>
+  inline __host__ __device__
+  Result operator()(const Argument& x) const
   {
-    // we static cast to Result to handle void Result without error
-    // in case Function's result is non-void
     return static_cast<Result>(m_f(thrust::raw_reference_cast(x)));
   }
 
   __thrust_exec_check_disable__
-  template<typename Argument1, typename Argument2>
-    inline __host__ __device__ Result operator()(Argument1 &x, Argument2 &y) const
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  Result operator()(Argument1& x, Argument2& y) const
   {
-    // we static cast to Result to handle void Result without error
-    // in case Function's result is non-void
-    return static_cast<Result>(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)));
+    return static_cast<Result>(m_f(thrust::raw_reference_cast(x),
+                                   thrust::raw_reference_cast(y)));
   }
 
   __thrust_exec_check_disable__
-  template<typename Argument1, typename Argument2>
-    inline __host__ __device__ Result operator()(const Argument1 &x, Argument2 &y) const
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  Result operator()(const Argument1& x, Argument2& y) const
   {
-    // we static cast to Result to handle void Result without error
-    // in case Function's result is non-void
-    return static_cast<Result>(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)));
+    return static_cast<Result>(m_f(thrust::raw_reference_cast(x),
+                                   thrust::raw_reference_cast(y)));
   }
 
   __thrust_exec_check_disable__
-  template<typename Argument1, typename Argument2>
-    inline __host__ __device__ Result operator()(const Argument1 &x, const Argument2 &y) const
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  Result operator()(const Argument1& x, const Argument2& y) const
   {
-    // we static cast to Result to handle void Result without error
-    // in case Function's result is non-void
-    return static_cast<Result>(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)));
+    return static_cast<Result>(m_f(thrust::raw_reference_cast(x),
+                                   thrust::raw_reference_cast(y)));
   }
 
   __thrust_exec_check_disable__
-  template<typename Argument1, typename Argument2>
-    inline __host__ __device__ Result operator()(Argument1 &x, const Argument2 &y) const
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  Result operator()(Argument1& x, const Argument2& y) const
   {
-    // we static cast to Result to handle void Result without error
-    // in case Function's result is non-void
-    return static_cast<Result>(m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y)));
+    return static_cast<Result>(m_f(thrust::raw_reference_cast(x),
+                                   thrust::raw_reference_cast(y)));
   }
 }; // end wrapped_function
 
+// Specialize for void return types:
+template <typename Function>
+struct wrapped_function<Function, void>
+{
+  // mutable because Function::operator() might be const
+  mutable Function m_f;
+  inline __host__ __device__
+  wrapped_function()
+    : m_f()
+  {}
+
+  inline __host__ __device__
+  wrapped_function(const Function& f)
+    : m_f(f)
+  {}
+
+  __thrust_exec_check_disable__
+  template <typename Argument>
+  inline __host__ __device__
+  void operator()(Argument& x) const
+  {
+    m_f(thrust::raw_reference_cast(x));
+  }
 
-} // end detail
-} // end thrust
+  __thrust_exec_check_disable__
+  template <typename Argument>
+  inline __host__ __device__
+  void operator()(const Argument& x) const
+  {
+    m_f(thrust::raw_reference_cast(x));
+  }
+
+  __thrust_exec_check_disable__
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  void operator()(Argument1& x, Argument2& y) const
+  {
+    m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y));
+  }
+
+  __thrust_exec_check_disable__
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  void operator()(const Argument1& x, Argument2& y) const
+  {
+    m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y));
+  }
+  __thrust_exec_check_disable__
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  void operator()(const Argument1& x, const Argument2& y) const
+  {
+    m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y));
+  }
+  __thrust_exec_check_disable__
+  template <typename Argument1, typename Argument2>
+  inline __host__ __device__
+  void operator()(Argument1& x, const Argument2& y) const
+  {
+    m_f(thrust::raw_reference_cast(x), thrust::raw_reference_cast(y));
+  }
+}; // end wrapped_function
 
+} // namespace detail
+} // namespace thrust
diff --git a/thrust/detail/internal_functional.h b/thrust/detail/internal_functional.h
@@ -23,6 +23,7 @@
 
 #include <thrust/tuple.h>
 #include <thrust/iterator/iterator_traits.h>
+#include <thrust/detail/static_assert.h>
 #include <thrust/detail/type_traits.h>
 #include <thrust/iterator/detail/tuple_of_iterator_references.h>
 #include <thrust/detail/raw_reference_cast.h>

diff --git a/thrust/iterator/detail/zip_iterator_base.h b/thrust/iterator/detail/zip_iterator_base.h
@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <thrust/advance.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/iterator/iterator_facade.h>
 #include <thrust/iterator/iterator_categories.h>
@@ -45,12 +46,12 @@ class advance_iterator
 public:
   inline __host__ __device__
   advance_iterator(DiffType step) : m_step(step) {}
-  
+
   __thrust_exec_check_disable__
   template<typename Iterator>
   inline __host__ __device__
   void operator()(Iterator& it) const
-  { it += m_step; }
+  { thrust::advance(it, m_step); }
 
 private:
   DiffType m_step;

diff --git a/thrust/system/tbb/detail/reduce_by_key.inl b/thrust/system/tbb/detail/reduce_by_key.inl
@@ -27,8 +27,9 @@
 #include <thrust/detail/range/tail_flags.h>
 #include <tbb/blocked_range.h>
 #include <tbb/parallel_for.h>
-#include <tbb/tbb_thread.h>
+
 #include <cassert>
+#include <thread>
 
 
 namespace thrust
@@ -281,7 +282,7 @@ template<typename DerivedPolicy, typename Iterator1, typename Iterator2, typenam
   }
 
   // count the number of processors
-  const unsigned int p = thrust::max<unsigned int>(1u, ::tbb::tbb_thread::hardware_concurrency());
+  const unsigned int p = thrust::max<unsigned int>(1u, std::thread::hardware_concurrency());
 
   // generate O(P) intervals of sequential work
   // XXX oversubscribing is a tuning opportunity