Fix/modernize thrust examples

thrust/examples/word_count.cmake

@@ -0,0 +1,21 @@
+target_compile_options(
+  ${example_target}
+  PRIVATE $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>: --extended-lambda>
+)
+
+# This check is actually not correct, because we must check the host compiler, not the CXX compiler.
+# We rely on these usually being the same ;)
+if (
+  "Clang" STREQUAL "${CMAKE_CXX_COMPILER_ID}"
+  AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL 13
+)
+  # When clang >= 13 is used as host compiler, we get the following warning:
+  #   nvcc_internal_extended_lambda_implementation:312:22: error: definition of implicit copy constructor for '__nv_hdl_wrapper_t<false, true, false, __nv_dl_tag<void (*)(), &TestAddressStabilityLambda, 2>, int (const int &)>' is deprecated because it has a user-declared copy assignment operator [-Werror,-Wdeprecated-copy]
+  #   312 | __nv_hdl_wrapper_t & operator=(const __nv_hdl_wrapper_t &in) = delete;
+  #       |                      ^
+  # Let's suppress it until NVBug 4980157 is resolved.
+  target_compile_options(
+    ${example_target}
+    PRIVATE $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>: -Wno-deprecated-copy>
+  )
+endif()

thrust/examples/word_count.cu

@@ -4,6 +4,8 @@
 #include <thrust/inner_product.h>
 #include <thrust/reduce.h>
 
+#include <cuda/std/iterator> // Required for std::begin/std::end
+
 #include <iostream>
 
 // This example computes the number of words in a text sample
@@ -18,15 +20,6 @@ __host__ __device__ bool is_alpha(const char c)
   return (c >= 'a' && c <= 'z') || (c >= 'A' && c <= 'Z');
 }
 
-// determines whether the right character begins a new word
-struct is_word_start
-{
-  __host__ __device__ bool operator()(const char& left, const char& right) const
-  {
-    return is_alpha(right) && !is_alpha(left);
-  }
-};
-
 int word_count(const thrust::device_vector<char>& input)
 {
   // check for empty string
@@ -35,14 +28,20 @@ int word_count(const thrust::device_vector<char>& input)
     return 0;
   }
 
+  // determines whether the right character begins a new word
+  auto is_word_start = [] __host__ __device__(const char& left, const char& right) {
+    return is_alpha(right) && !is_alpha(left);
+  };
+
   // compute the number characters that start a new word
   int wc = thrust::inner_product(
     input.begin(),
     input.end() - 1, // sequence of left characters
     input.begin() + 1, // sequence of right characters
     0, // initialize sum to 0
-    cuda::std::plus<int>(), // sum values together
-    is_word_start()); // how to compare the left and right characters
+    cuda::std::plus<int>{}, // sum values together
+    is_word_start // how to compare the left and right characters
+  );
 
   // if the first character is alphabetical, then it also begins a word
   if (is_alpha(input.front()))
@@ -65,16 +64,16 @@ int main()
     "  On the morrow he will leave me, as my hopes have flown before.'\n"
     "  Then the bird said, `Nevermore.'\n";
 
-  std::cout << "Text sample:" << std::endl;
-  std::cout << raw_input << std::endl;
+  std::cout << "Text sample:\n";
+  std::cout << raw_input << "\n";
 
   // transfer to device
-  thrust::device_vector<char> input(raw_input, raw_input + sizeof(raw_input));
+  thrust::device_vector<char> input(cuda::std::begin(raw_input), cuda::std::end(raw_input));
 
   // count words
   int wc = word_count(input);
 
-  std::cout << "Text sample contains " << wc << " words" << std::endl;
+  std::cout << "Text sample contains " << wc << " words\n";
 
   return 0;
 }