From 9e075769ba3616bb65211505918ad407139eeeb5 Mon Sep 17 00:00:00 2001 From: Louis Mourgues-Auzemery Date: Tue, 6 Jan 2026 22:54:34 +0100 Subject: [PATCH 01/11] Modernize word_count.cu: Replace functor with __host__ __device__ lambda --- thrust/examples/word_count.cu | 37 ++++++++++++++++++----------------- 1 file changed, 19 insertions(+), 18 deletions(-) diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 260be025c4f..921ad25a402 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -5,6 +5,7 @@ #include #include +#include // Required for std::begin/std::end // This example computes the number of words in a text sample // with a single call to thrust::inner_product. The algorithm @@ -18,15 +19,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& input) { // check for empty string @@ -35,14 +27,22 @@ int word_count(const thrust::device_vector& input) return 0; } + // determines whether the right character begins a new word + // Define lambda locally for safety and clarity + 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(), // sum values together - is_word_start()); // how to compare the left and right characters + input.end() - 1, // sequence of left characters + input.begin() + 1, // sequence of right characters + 0, // initialize sum to 0 + cuda::std::plus{}, // sum values together (use {} for uniform init) + is_word_start // Pass the lambda object directly (NO parentheses) + ); // if the first character is alphabetical, then it also begins a word if (is_alpha(input.front())) @@ -65,16 +65,17 @@ 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 input(raw_input, raw_input + sizeof(raw_input)); + // Use std::begin/end for safer C++ iterator access + thrust::device_vector input(std::begin(raw_input), 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; } From e89e7c26b8f92e98a221ad7fa01ffb5585f6b110 Mon Sep 17 00:00:00 2001 From: Louis Mourgues-Auzemery Date: Tue, 6 Jan 2026 23:00:18 +0100 Subject: [PATCH 02/11] Modernize word_count.cu: Replace functor with __host__ __device__ lambda, and adapted to clang-format --- thrust/examples/word_count.cu | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 921ad25a402..451b7e04a68 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -29,19 +29,18 @@ int word_count(const thrust::device_vector& input) // determines whether the right character begins a new word // Define lambda locally for safety and clarity - auto is_word_start = [] __host__ __device__ (const char& left, const char& right) - { + 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 + input.end() - 1, // sequence of left characters + input.begin() + 1, // sequence of right characters + 0, // initialize sum to 0 cuda::std::plus{}, // sum values together (use {} for uniform init) - is_word_start // Pass the lambda object directly (NO parentheses) + is_word_start // Pass the lambda object directly (NO parentheses) ); // if the first character is alphabetical, then it also begins a word From c14a968013e3a186766fe2444f0ad3cd55ac31c2 Mon Sep 17 00:00:00 2001 From: Flawxd Date: Tue, 6 Jan 2026 23:55:23 +0100 Subject: [PATCH 03/11] Apply suggestions from code review Co-authored-by: Bernhard Manfred Gruber --- thrust/examples/word_count.cu | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 451b7e04a68..f5d899023e9 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -28,7 +28,6 @@ int word_count(const thrust::device_vector& input) } // determines whether the right character begins a new word - // Define lambda locally for safety and clarity auto is_word_start = [] __host__ __device__(const char& left, const char& right) { return is_alpha(right) && !is_alpha(left); }; @@ -39,7 +38,7 @@ int word_count(const thrust::device_vector& input) input.end() - 1, // sequence of left characters input.begin() + 1, // sequence of right characters 0, // initialize sum to 0 - cuda::std::plus{}, // sum values together (use {} for uniform init) + cuda::std::plus{}, // sum values together is_word_start // Pass the lambda object directly (NO parentheses) ); @@ -68,7 +67,6 @@ int main() std::cout << raw_input << "\n"; // transfer to device - // Use std::begin/end for safer C++ iterator access thrust::device_vector input(std::begin(raw_input), std::end(raw_input)); // count words From a80617f1813b33de4f44da9fc895da2465cf2e81 Mon Sep 17 00:00:00 2001 From: Louis Mourgues-Auzemery Date: Wed, 7 Jan 2026 00:05:16 +0100 Subject: [PATCH 04/11] Review suggestion for important comment --- thrust/examples/word_count.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index f5d899023e9..84ed7169a0e 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -39,7 +39,7 @@ int word_count(const thrust::device_vector& input) input.begin() + 1, // sequence of right characters 0, // initialize sum to 0 cuda::std::plus{}, // sum values together - is_word_start // Pass the lambda object directly (NO parentheses) + is_word_start // how to compare the left and right characters ); // if the first character is alphabetical, then it also begins a word From 98489e48f170687a35212401737ee006506d9be4 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 7 Jan 2026 12:12:53 +0100 Subject: [PATCH 05/11] Update thrust/examples/word_count.cu Co-authored-by: Michael Schellenberger Costa --- thrust/examples/word_count.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 84ed7169a0e..76fae7cf3d6 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -5,7 +5,7 @@ #include #include -#include // Required for std::begin/std::end +#include // Required for std::begin/std::end // This example computes the number of words in a text sample // with a single call to thrust::inner_product. The algorithm From a1042774e6d6c3b6ccb9d1a6d74a414a05c95d74 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 7 Jan 2026 12:13:04 +0100 Subject: [PATCH 06/11] Update thrust/examples/word_count.cu Co-authored-by: Michael Schellenberger Costa --- thrust/examples/word_count.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 76fae7cf3d6..14cc7c54fe8 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -67,7 +67,7 @@ int main() std::cout << raw_input << "\n"; // transfer to device - thrust::device_vector input(std::begin(raw_input), std::end(raw_input)); + thrust::device_vector input(cuda::std::begin(raw_input), cuda::std::end(raw_input)); // count words int wc = word_count(input); From 0185accc4f3adc8703713f07e7020e65e3497615 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Wed, 7 Jan 2026 11:52:33 +0000 Subject: [PATCH 07/11] [pre-commit.ci] auto code formatting --- thrust/examples/word_count.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 14cc7c54fe8..1d22cacaf9a 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -4,9 +4,10 @@ #include #include -#include #include // Required for std::begin/std::end +#include + // This example computes the number of words in a text sample // with a single call to thrust::inner_product. The algorithm // counts the number of characters which start a new word, i.e. From 06e50ff0c10b02593329fb7c2ec4f6af1dce96fd Mon Sep 17 00:00:00 2001 From: Louis Mourgues-Auzemery Date: Wed, 7 Jan 2026 13:09:45 +0100 Subject: [PATCH 08/11] Added the cmake file that was suggested for word_count --- thrust/examples/word_count.cmake | 4 ++++ 1 file changed, 4 insertions(+) create mode 100644 thrust/examples/word_count.cmake diff --git a/thrust/examples/word_count.cmake b/thrust/examples/word_count.cmake new file mode 100644 index 00000000000..a016b72c1b1 --- /dev/null +++ b/thrust/examples/word_count.cmake @@ -0,0 +1,4 @@ +target_compile_options( + ${example_target} + PRIVATE $<$: --extended-lambda> +) \ No newline at end of file From 3f87a1d8fb79cf613c8df3e15378155dc8e55094 Mon Sep 17 00:00:00 2001 From: "pre-commit-ci[bot]" <66853113+pre-commit-ci[bot]@users.noreply.github.com> Date: Wed, 7 Jan 2026 12:26:01 +0000 Subject: [PATCH 09/11] [pre-commit.ci] auto code formatting --- thrust/examples/word_count.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/word_count.cmake b/thrust/examples/word_count.cmake index a016b72c1b1..d26c340f299 100644 --- a/thrust/examples/word_count.cmake +++ b/thrust/examples/word_count.cmake @@ -1,4 +1,4 @@ target_compile_options( ${example_target} PRIVATE $<$: --extended-lambda> -) \ No newline at end of file +) From 3fd0ed9ec59342502e8ad547862d9ebef62ddcc5 Mon Sep 17 00:00:00 2001 From: Louis Mourgues-Auzemery Date: Wed, 7 Jan 2026 22:50:19 +0100 Subject: [PATCH 10/11] fixed cmake with suggestion --- thrust/examples/word_count.cmake | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/thrust/examples/word_count.cmake b/thrust/examples/word_count.cmake index d26c340f299..4ca997ca25e 100644 --- a/thrust/examples/word_count.cmake +++ b/thrust/examples/word_count.cmake @@ -2,3 +2,13 @@ target_compile_options( ${example_target} PRIVATE $<$: --extended-lambda> ) + +# 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, 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 $<$: -Wno-deprecated-copy> +) From 113bb158d9ef007e218f4eae0baaef600e00be91 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 9 Jan 2026 09:51:24 +0100 Subject: [PATCH 11/11] Update thrust/examples/word_count.cmake --- thrust/examples/word_count.cmake | 23 +++++++++++++++-------- 1 file changed, 15 insertions(+), 8 deletions(-) diff --git a/thrust/examples/word_count.cmake b/thrust/examples/word_count.cmake index 4ca997ca25e..9b8733dd322 100644 --- a/thrust/examples/word_count.cmake +++ b/thrust/examples/word_count.cmake @@ -3,12 +3,19 @@ target_compile_options( PRIVATE $<$: --extended-lambda> ) -# 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, 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 $<$: -Wno-deprecated-copy> +# 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, 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 $<$: -Wno-deprecated-copy> + ) +endif()