From 6de8ae19160574ef29b8b7a2d56e6db1736f40b3 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Tue, 26 Aug 2025 21:54:42 -0500 Subject: [PATCH 01/19] Modernize Thrust examples following PR #753 patterns - Remove legacy include/host_device.h headers from 40 example files - Replace manual element assignment with std::initializer_list - Use range-based for loops where appropriate - Apply STL algorithms (std::generate) with lambdas - Use .size() instead of hardcoded array sizes - Improve semantic naming and inline usage - Maintain compatibility with current CUDA/Thrust version - Avoid thrust::enumerate (not available in current version) Files modernized: arbitrary_transformation.cu, basic_vector.cu, bounding_box.cu, bucket_sort2d.cu, constant_iterator.cu, counting_iterator.cu, device_ptr.cu, discrete_voronoi.cu, dot_products_with_zip.cu, expand.cu, histogram.cu, lambda.cu, lexicographical_sort.cu, max_abs_diff.cu, minmax.cu, mode.cu, monte_carlo.cu, monte_carlo_disjoint_sequences.cu, norm.cu, padded_grid_reduction.cu, permutation_iterator.cu, raw_reference_cast.cu, remove_points2d.cu, repeated_range.cu, saxpy.cu, scan_matrix_by_rows.cu, simple_moving_average.cu, sort.cu, sorting_aos_vs_soa.cu, stream_compaction.cu, sum_rows.cu, summary_statistics.cu, summed_area_table.cu, tiled_range.cu, transform_input_output_iterator.cu, transform_iterator.cu, transform_output_iterator.cu, uninitialized_vector.cu, weld_vertices.cu, word_count.cu --- thrust/examples/arbitrary_transformation.cu | 23 ++++------- thrust/examples/basic_vector.cu | 10 +---- thrust/examples/bounding_box.cu | 22 +++++------ thrust/examples/bucket_sort2d.cu | 2 +- thrust/examples/constant_iterator.cu | 6 +-- thrust/examples/counting_iterator.cu | 12 +----- thrust/examples/device_ptr.cu | 2 +- thrust/examples/discrete_voronoi.cu | 13 +++---- thrust/examples/dot_products_with_zip.cu | 6 +-- thrust/examples/expand.cu | 12 ++---- thrust/examples/histogram.cu | 7 ++-- thrust/examples/lambda.cu | 23 ++++++----- thrust/examples/lexicographical_sort.cu | 9 ++--- thrust/examples/max_abs_diff.cu | 13 ++----- thrust/examples/minmax.cu | 8 +--- thrust/examples/mode.cu | 7 +--- thrust/examples/monte_carlo.cu | 2 - .../monte_carlo_disjoint_sequences.cu | 2 - thrust/examples/norm.cu | 9 +---- thrust/examples/padded_grid_reduction.cu | 10 ++--- thrust/examples/permutation_iterator.cu | 14 +------ thrust/examples/raw_reference_cast.cu | 2 - thrust/examples/remove_points2d.cu | 13 ++----- thrust/examples/repeated_range.cu | 8 +--- thrust/examples/saxpy.cu | 14 +++---- thrust/examples/scan_matrix_by_rows.cu | 8 ++-- thrust/examples/simple_moving_average.cu | 11 ++---- thrust/examples/sort.cu | 39 +++++++------------ thrust/examples/sorting_aos_vs_soa.cu | 16 ++++---- thrust/examples/stream_compaction.cu | 2 - thrust/examples/sum_rows.cu | 12 ++---- thrust/examples/summary_statistics.cu | 6 +-- thrust/examples/summed_area_table.cu | 2 - thrust/examples/tiled_range.cu | 8 +--- .../transform_input_output_iterator.cu | 6 +-- thrust/examples/transform_iterator.cu | 2 - thrust/examples/transform_output_iterator.cu | 21 +++++----- thrust/examples/uninitialized_vector.cu | 2 - thrust/examples/weld_vertices.cu | 15 +++---- thrust/examples/word_count.cu | 5 +-- 40 files changed, 131 insertions(+), 273 deletions(-) diff --git a/thrust/examples/arbitrary_transformation.cu b/thrust/examples/arbitrary_transformation.cu index 2389955cb84c..976b6826dffa 100644 --- a/thrust/examples/arbitrary_transformation.cu +++ b/thrust/examples/arbitrary_transformation.cu @@ -7,7 +7,7 @@ #include -#include "include/host_device.h" +// #include "include/host_device.h" // Not needed for this example // This example shows how to implement an arbitrary transformation of // the form output[i] = F(first[i], second[i], third[i], ... ). @@ -62,21 +62,12 @@ struct arbitrary_functor2 int main() { - // allocate storage - thrust::device_vector A(5); - thrust::device_vector B(5); - thrust::device_vector C(5); + // allocate and initialize storage using initializer lists + thrust::device_vector A{3, 4, 0, 8, 2}; + thrust::device_vector B{6, 7, 2, 1, 8}; + thrust::device_vector C{2, 5, 7, 4, 3}; thrust::device_vector D1(5); - // clang-format off - // initialize input vectors - A[0] = 3; B[0] = 6; C[0] = 2; - A[1] = 4; B[1] = 7; C[1] = 5; - A[2] = 0; B[2] = 2; C[2] = 7; - A[3] = 8; B[3] = 1; C[3] = 4; - A[4] = 2; B[4] = 8; C[4] = 3; - // clang-format on - // apply the transformation thrust::for_each(thrust::make_zip_iterator(A.begin(), B.begin(), C.begin(), D1.begin()), thrust::make_zip_iterator(A.end(), B.end(), C.end(), D1.end()), @@ -84,7 +75,7 @@ int main() // print the output std::cout << "Tuple functor" << std::endl; - for (int i = 0; i < 5; i++) + for (size_t i = 0; i < A.size(); i++) { std::cout << A[i] << " + " << B[i] << " * " << C[i] << " = " << D1[i] << std::endl; } @@ -97,7 +88,7 @@ int main() // print the output std::cout << "N-ary functor" << std::endl; - for (int i = 0; i < 5; i++) + for (size_t i = 0; i < A.size(); i++) { std::cout << A[i] << " + " << B[i] << " * " << C[i] << " = " << D2[i] << std::endl; } diff --git a/thrust/examples/basic_vector.cu b/thrust/examples/basic_vector.cu index 7c1c39fac1e0..530e06680cb9 100644 --- a/thrust/examples/basic_vector.cu +++ b/thrust/examples/basic_vector.cu @@ -5,14 +5,8 @@ int main() { - // H has storage for 4 integers - thrust::host_vector H(4); - - // initialize individual elements - H[0] = 14; - H[1] = 20; - H[2] = 38; - H[3] = 46; + // H has storage for 4 integers with direct initialization + thrust::host_vector H{14, 20, 38, 46}; // H.size() returns the size of vector H std::cout << "H has size " << H.size() << std::endl; diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index cd73a9c3042a..c2a2a14f4ca9 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -4,7 +4,7 @@ #include #include -#include "include/host_device.h" +// #include "include/host_device.h" // Not needed for this example // This example shows how to compute a bounding box // for a set of points in two dimensions. @@ -54,7 +54,7 @@ struct bbox }; // reduce a pair of bounding boxes (a,b) to a bounding box containing a and b -struct bbox_reduction +struct bbox_union : public thrust::binary_function { __host__ __device__ bbox operator()(bbox a, bbox b) { @@ -78,21 +78,17 @@ int main() thrust::device_vector points(N); // generate some random points in the unit square - for (size_t i = 0; i < N; i++) - { - float x = u01(rng); - float y = u01(rng); - points[i] = point2d(x, y); - } + std::generate(points.begin(), points.end(), [&] { + auto x = u01(rng); + auto y = u01(rng); + return point2d(x, y); + }); // initial bounding box contains first point - bbox init = bbox(points[0], points[0]); - - // binary reduction operation - bbox_reduction binary_op; + bbox first_point = bbox(points[0], points[0]); // compute the bounding box for the point set - bbox result = thrust::reduce(points.begin(), points.end(), init, binary_op); + bbox result = thrust::reduce(points.begin(), points.end(), first_point, bbox_union()); // print output std::cout << "bounding box " << std::fixed; diff --git a/thrust/examples/bucket_sort2d.cu b/thrust/examples/bucket_sort2d.cu index bb1bf7c1b08e..eb00d0170953 100644 --- a/thrust/examples/bucket_sort2d.cu +++ b/thrust/examples/bucket_sort2d.cu @@ -9,7 +9,7 @@ #include #include -#include "include/host_device.h" +// #include "include/host_device.h" // Not needed for this example // define a 2d float vector using vec2 = thrust::tuple; diff --git a/thrust/examples/constant_iterator.cu b/thrust/examples/constant_iterator.cu index a2f01ee9f975..e3c21a939972 100644 --- a/thrust/examples/constant_iterator.cu +++ b/thrust/examples/constant_iterator.cu @@ -9,11 +9,7 @@ int main() { - thrust::device_vector data(4); - data[0] = 3; - data[1] = 7; - data[2] = 2; - data[3] = 5; + thrust::device_vector data{3, 7, 2, 5}; // add 10 to all values in data thrust::transform( diff --git a/thrust/examples/counting_iterator.cu b/thrust/examples/counting_iterator.cu index 96da36af783b..428db8b13cad 100644 --- a/thrust/examples/counting_iterator.cu +++ b/thrust/examples/counting_iterator.cu @@ -11,15 +11,7 @@ int main() // this example computes indices for all the nonzero values in a sequence // sequence of zero and nonzero values - thrust::device_vector stencil(8); - stencil[0] = 0; - stencil[1] = 1; - stencil[2] = 1; - stencil[3] = 0; - stencil[4] = 0; - stencil[5] = 1; - stencil[6] = 0; - stencil[7] = 1; + thrust::device_vector stencil{0, 1, 1, 0, 0, 1, 0, 1}; // storage for the nonzero indices thrust::device_vector indices(8); @@ -35,7 +27,7 @@ int main() // indices now contains [1,2,5,7] // print result - std::cout << "found " << (indices_end - indices.begin()) << " nonzero values at indices:\n"; + std::cout << "found " << std::distance(indices.begin(), indices_end) << " nonzero values at indices:\n"; thrust::copy(indices.begin(), indices_end, std::ostream_iterator(std::cout, "\n")); return 0; diff --git a/thrust/examples/device_ptr.cu b/thrust/examples/device_ptr.cu index be013cba810d..3fcb555be23d 100644 --- a/thrust/examples/device_ptr.cu +++ b/thrust/examples/device_ptr.cu @@ -15,7 +15,7 @@ int main() // device_ptr supports pointer arithmetic thrust::device_ptr first = d_ptr; thrust::device_ptr last = d_ptr + 10; - std::cout << "device array contains " << (last - first) << " values\n"; + std::cout << "device array contains " << std::distance(first, last) << " values\n"; // algorithms work as expected thrust::sequence(first, last); diff --git a/thrust/examples/discrete_voronoi.cu b/thrust/examples/discrete_voronoi.cu index e9887f6a9cfe..fa0c78e6f633 100644 --- a/thrust/examples/discrete_voronoi.cu +++ b/thrust/examples/discrete_voronoi.cu @@ -10,7 +10,6 @@ #include #include -#include "include/host_device.h" #include "include/timer.h" // Compute an approximate Voronoi Diagram with a Jump Flooding Algorithm (JFA) @@ -26,11 +25,11 @@ // Tuple = -struct minFunctor +struct voronoi_site_selector { int m, n, k; - __host__ __device__ minFunctor(int m, int n, int k) + __host__ __device__ voronoi_site_selector(int m, int n, int k) : m(m) , n(n) , k(k) @@ -125,9 +124,9 @@ void print(int m, int n, const thrust::device_vector& d_data) { thrust::host_vector h_data = d_data; - for (int i = 0; i < m; i++) + for (size_t i = 0; i < static_cast(m); i++) { - for (int j = 0; j < n; j++) + for (size_t j = 0; j < static_cast(n); j++) { std::cout << std::setw(4) << h_data[i * n + j] << " "; } @@ -140,7 +139,7 @@ void generate_random_sites(thrust::host_vector& t, int Nb, int m, int n) thrust::default_random_engine rng; thrust::uniform_int_distribution dist(0, m * n - 1); - for (int k = 0; k < Nb; k++) + for (size_t k = 0; k < static_cast(Nb); k++) { int index = dist(rng); t[index] = index + 1; @@ -199,7 +198,7 @@ void jfa(thrust::device_vector& in, thrust::device_vector& out, unsign thrust::counting_iterator(0)) + n * m, out.begin(), - minFunctor(m, n, k)); + voronoi_site_selector(m, n, k)); } /********************************************/ diff --git a/thrust/examples/dot_products_with_zip.cu b/thrust/examples/dot_products_with_zip.cu index 0c456a9d4b80..4b7a16e01c8b 100644 --- a/thrust/examples/dot_products_with_zip.cu +++ b/thrust/examples/dot_products_with_zip.cu @@ -5,7 +5,6 @@ #include #include -#include "include/host_device.h" // This example shows how thrust::zip_iterator can be used to create a // 'virtual' array of structures. In this case the structure is a 3d @@ -33,10 +32,7 @@ thrust::host_vector random_vector(const size_t N, unsigned int seed = thr thrust::default_random_engine rng(seed); thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector temp(N); - for (size_t i = 0; i < N; i++) - { - temp[i] = u01(rng); - } + std::generate(temp.begin(), temp.end(), [&]() { return u01(rng); }); return temp; } diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index 39cd750aa769..a979d0846eef 100644 --- a/thrust/examples/expand.cu +++ b/thrust/examples/expand.cu @@ -62,15 +62,11 @@ void print(const std::string& s, const Vector& v) int main() { - int counts[] = {3, 5, 2, 0, 1, 3, 4, 2, 4}; - int values[] = {1, 2, 3, 4, 5, 6, 7, 8, 9}; + thrust::device_vector d_counts = {3, 5, 2, 0, 1, 3, 4, 2, 4}; + thrust::device_vector d_values = {1, 2, 3, 4, 5, 6, 7, 8, 9}; - size_t input_size = sizeof(counts) / sizeof(int); - size_t output_size = thrust::reduce(counts, counts + input_size); - - // copy inputs to device - thrust::device_vector d_counts(counts, counts + input_size); - thrust::device_vector d_values(values, values + input_size); + size_t input_size = d_counts.size(); + size_t output_size = thrust::reduce(d_counts.begin(), d_counts.end()); thrust::device_vector d_output(output_size); // expand values according to counts diff --git a/thrust/examples/histogram.cu b/thrust/examples/histogram.cu index 1bc7455d0440..8141e958b773 100644 --- a/thrust/examples/histogram.cu +++ b/thrust/examples/histogram.cu @@ -149,15 +149,14 @@ int main() // generate random data on the host thrust::host_vector input(N); - for (int i = 0; i < N; i++) - { + std::generate(input.begin(), input.end(), [&]() { int sum = 0; for (int j = 0; j < S; j++) { sum += dist(rng); } - input[i] = sum / S; - } + return sum / S; + }); // demonstrate dense histogram method { diff --git a/thrust/examples/lambda.cu b/thrust/examples/lambda.cu index 783f7244ad95..4cf098435704 100644 --- a/thrust/examples/lambda.cu +++ b/thrust/examples/lambda.cu @@ -4,7 +4,6 @@ #include -#include "include/host_device.h" // This example demonstrates the use of placeholders to implement // the SAXPY operation (i.e. Y[i] = a * X[i] + Y[i]). @@ -41,14 +40,14 @@ struct saxpy_functor int main() { // input data - float a = 2.0f; - float x[4] = {1, 2, 3, 4}; - float y[4] = {1, 1, 1, 1}; + float a = 2.0f; + thrust::device_vector x_data = {1, 2, 3, 4}; + thrust::device_vector y_data = {1, 1, 1, 1}; // SAXPY implemented with a functor (function object) { - thrust::device_vector X(x, x + 4); - thrust::device_vector Y(y, y + 4); + thrust::device_vector X = x_data; + thrust::device_vector Y = y_data; thrust::transform( X.begin(), @@ -58,16 +57,16 @@ int main() saxpy_functor(a)); // functor std::cout << "SAXPY (functor method)" << std::endl; - for (size_t i = 0; i < 4; i++) + for (size_t i = 0; i < Y.size(); i++) { - std::cout << a << " * " << x[i] << " + " << y[i] << " = " << Y[i] << std::endl; + std::cout << a << " * " << x_data[i] << " + " << y_data[i] << " = " << Y[i] << std::endl; } } // SAXPY implemented with a placeholders { - thrust::device_vector X(x, x + 4); - thrust::device_vector Y(y, y + 4); + thrust::device_vector X = x_data; + thrust::device_vector Y = y_data; thrust::transform( X.begin(), @@ -77,9 +76,9 @@ int main() a * _1 + _2); // placeholder expression std::cout << "SAXPY (placeholder method)" << std::endl; - for (size_t i = 0; i < 4; i++) + for (size_t i = 0; i < Y.size(); i++) { - std::cout << a << " * " << x[i] << " + " << y[i] << " = " << Y[i] << std::endl; + std::cout << a << " * " << x_data[i] << " + " << y_data[i] << " = " << Y[i] << std::endl; } } diff --git a/thrust/examples/lexicographical_sort.cu b/thrust/examples/lexicographical_sort.cu index cb85a12da746..3aa608e0acfa 100644 --- a/thrust/examples/lexicographical_sort.cu +++ b/thrust/examples/lexicographical_sort.cu @@ -41,10 +41,7 @@ thrust::host_vector random_vector(size_t N) static thrust::default_random_engine rng; static thrust::uniform_int_distribution dist(0, 9); - for (size_t i = 0; i < N; i++) - { - vec[i] = dist(rng); - } + std::generate(vec.begin(), vec.end(), [&]() { return dist(rng); }); return vec; } @@ -59,7 +56,7 @@ int main() thrust::device_vector lower = random_vector(N); std::cout << "Unsorted Keys" << std::endl; - for (size_t i = 0; i < N; i++) + for (size_t i = 0; i < upper.size(); i++) { std::cout << "(" << upper[i] << "," << middle[i] << "," << lower[i] << ")" << std::endl; } @@ -82,7 +79,7 @@ int main() apply_permutation(upper, permutation); std::cout << "Sorted Keys" << std::endl; - for (size_t i = 0; i < N; i++) + for (size_t i = 0; i < upper.size(); i++) { std::cout << "(" << upper[i] << "," << middle[i] << "," << lower[i] << ")" << std::endl; } diff --git a/thrust/examples/max_abs_diff.cu b/thrust/examples/max_abs_diff.cu index 5614aaee48ba..6f2c04177b34 100644 --- a/thrust/examples/max_abs_diff.cu +++ b/thrust/examples/max_abs_diff.cu @@ -5,7 +5,7 @@ #include #include -#include "include/host_device.h" +// #include "include/host_device.h" // Legacy header - commented out // this example computes the maximum absolute difference // between the elements of two vectors @@ -21,15 +21,8 @@ struct abs_diff int main() { - thrust::device_vector d_a(4); - thrust::device_vector d_b(4); - - // clang-format off - d_a[0] = 1.0; d_b[0] = 2.0; - d_a[1] = 2.0; d_b[1] = 4.0; - d_a[2] = 3.0; d_b[2] = 3.0; - d_a[3] = 4.0; d_b[3] = 0.0; - // clang-format on + thrust::device_vector d_a = {1.0, 2.0, 3.0, 4.0}; + thrust::device_vector d_b = {2.0, 4.0, 3.0, 0.0}; // initial value of the reduction float init = 0; diff --git a/thrust/examples/minmax.cu b/thrust/examples/minmax.cu index 772257cf7af8..a412336ec2d4 100644 --- a/thrust/examples/minmax.cu +++ b/thrust/examples/minmax.cu @@ -5,7 +5,6 @@ #include #include -#include "include/host_device.h" // compute minimum and maximum values in a single reduction @@ -60,10 +59,7 @@ int main() // initialize data on host thrust::device_vector data(N); - for (size_t i = 0; i < data.size(); i++) - { - data[i] = dist(rng); - } + std::generate(data.begin(), data.end(), [&]() { return dist(rng); }); // setup arguments minmax_unary_op unary_op; @@ -77,7 +73,7 @@ int main() // print results std::cout << "[ "; - for (size_t i = 0; i < N; i++) + for (size_t i = 0; i < data.size(); i++) { std::cout << data[i] << " "; } diff --git a/thrust/examples/mode.cu b/thrust/examples/mode.cu index 6849c2c5fc70..a45f3c29857a 100644 --- a/thrust/examples/mode.cu +++ b/thrust/examples/mode.cu @@ -26,10 +26,7 @@ int main() // generate random data on the host thrust::host_vector h_data(N); - for (size_t i = 0; i < N; i++) - { - h_data[i] = dist(rng); - } + std::generate(h_data.begin(), h_data.end(), [&]() { return dist(rng); }); // transfer data to device thrust::device_vector d_data(h_data); @@ -70,7 +67,7 @@ int main() thrust::device_vector::iterator mode_iter; mode_iter = thrust::max_element(d_output_counts.begin(), d_output_counts.end()); - int mode = d_output_keys[mode_iter - d_output_counts.begin()]; + int mode = d_output_keys[std::distance(d_output_counts.begin(), mode_iter)]; int occurrences = *mode_iter; std::cout << "Modal value " << mode << " occurs " << occurrences << " times " << std::endl; diff --git a/thrust/examples/monte_carlo.cu b/thrust/examples/monte_carlo.cu index 041666844483..659bd9a892c5 100644 --- a/thrust/examples/monte_carlo.cu +++ b/thrust/examples/monte_carlo.cu @@ -7,8 +7,6 @@ #include #include -#include "include/host_device.h" - // we could vary M & N to find the perf sweet spot __host__ __device__ unsigned int hash(unsigned int a) diff --git a/thrust/examples/monte_carlo_disjoint_sequences.cu b/thrust/examples/monte_carlo_disjoint_sequences.cu index fa56abb9c0ac..8e429bca032b 100644 --- a/thrust/examples/monte_carlo_disjoint_sequences.cu +++ b/thrust/examples/monte_carlo_disjoint_sequences.cu @@ -6,8 +6,6 @@ #include #include -#include "include/host_device.h" - // The technique demonstrated in the example monte_carlo.cu // assigns an independently seeded random number generator to each // of 30K threads, and uses a hashing scheme based on thread index to diff --git a/thrust/examples/norm.cu b/thrust/examples/norm.cu index a706ee2970b1..3f7f62947a89 100644 --- a/thrust/examples/norm.cu +++ b/thrust/examples/norm.cu @@ -6,8 +6,6 @@ #include #include -#include "include/host_device.h" - // This example computes the norm [1] of a vector. The norm is // computed by squaring all numbers in the vector, summing the // squares, and taking the square root of the sum of squares. In @@ -31,11 +29,8 @@ struct square int main() { - // initialize host array - float x[4] = {1.0, 2.0, 3.0, 4.0}; - - // transfer to device - thrust::device_vector d_x(x, x + 4); + // initialize device vector directly + thrust::device_vector d_x = {1.0, 2.0, 3.0, 4.0}; // setup arguments square unary_op; diff --git a/thrust/examples/padded_grid_reduction.cu b/thrust/examples/padded_grid_reduction.cu index 943b03373729..aa7a91455452 100644 --- a/thrust/examples/padded_grid_reduction.cu +++ b/thrust/examples/padded_grid_reduction.cu @@ -12,8 +12,6 @@ #include -#include "include/host_device.h" - // This example computes the minimum and maximum values // over a padded grid. The padded values are not considered // during the reduction operation. @@ -82,9 +80,9 @@ int main() thrust::device_vector data(M * N, -1); // initialize valid values in grid - for (int i = 0; i < M; i++) + for (size_t i = 0; i < static_cast(M); i++) { - for (int j = 0; j < n; j++) + for (size_t j = 0; j < static_cast(n); j++) { data[i * N + j] = dist(rng); } @@ -93,10 +91,10 @@ int main() // print full grid std::cout << "padded grid" << std::endl; std::cout << std::fixed << std::setprecision(4); - for (int i = 0; i < M; i++) + for (size_t i = 0; i < static_cast(M); i++) { std::cout << " "; - for (int j = 0; j < N; j++) + for (size_t j = 0; j < static_cast(N); j++) { std::cout << data[i * N + j] << " "; } diff --git a/thrust/examples/permutation_iterator.cu b/thrust/examples/permutation_iterator.cu index aeac96f054b2..9889fc3f47e1 100644 --- a/thrust/examples/permutation_iterator.cu +++ b/thrust/examples/permutation_iterator.cu @@ -10,20 +10,10 @@ int main() { // gather locations - thrust::device_vector map(4); - map[0] = 3; - map[1] = 1; - map[2] = 0; - map[3] = 5; + thrust::device_vector map = {3, 1, 0, 5}; // array to gather from - thrust::device_vector source(6); - source[0] = 10; - source[1] = 20; - source[2] = 30; - source[3] = 40; - source[4] = 50; - source[5] = 60; + thrust::device_vector source = {10, 20, 30, 40, 50, 60}; // fuse gather with reduction: // sum = source[map[0]] + source[map[1]] + ... diff --git a/thrust/examples/raw_reference_cast.cu b/thrust/examples/raw_reference_cast.cu index 2b3229928149..32592f5ccc1f 100644 --- a/thrust/examples/raw_reference_cast.cu +++ b/thrust/examples/raw_reference_cast.cu @@ -5,8 +5,6 @@ #include -#include "include/host_device.h" - // This example illustrates how to use the raw_reference_cast to convert // system-specific reference wrappers into native references. // diff --git a/thrust/examples/remove_points2d.cu b/thrust/examples/remove_points2d.cu index 7b906b2a90d5..9af8c2cb2a51 100644 --- a/thrust/examples/remove_points2d.cu +++ b/thrust/examples/remove_points2d.cu @@ -2,8 +2,6 @@ #include #include -#include "include/host_device.h" - // This example generates random points in the // unit square [0,1)x[0,1) and then removes all // points where x^2 + y^2 > 1 @@ -41,16 +39,13 @@ int main() thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector x(N); thrust::host_vector y(N); - for (size_t i = 0; i < N; i++) - { - x[i] = u01(rng); - y[i] = u01(rng); - } + std::generate(x.begin(), x.end(), [&]() { return u01(rng); }); + std::generate(y.begin(), y.end(), [&]() { return u01(rng); }); // print the initial points std::cout << std::fixed; std::cout << "Generated " << N << " points" << std::endl; - for (size_t i = 0; i < N; i++) + for (size_t i = 0; i < x.size(); i++) { std::cout << "(" << x[i] << "," << y[i] << ")" << std::endl; } @@ -69,7 +64,7 @@ int main() // print the filtered points std::cout << "After stream compaction, " << new_size << " points remain" << std::endl; - for (size_t i = 0; i < new_size; i++) + for (size_t i = 0; i < x.size(); i++) { std::cout << "(" << x[i] << "," << y[i] << ")" << std::endl; } diff --git a/thrust/examples/repeated_range.cu b/thrust/examples/repeated_range.cu index e454d9873143..218ab98352a2 100644 --- a/thrust/examples/repeated_range.cu +++ b/thrust/examples/repeated_range.cu @@ -8,8 +8,6 @@ #include -#include "include/host_device.h" - // this example illustrates how to make repeated access to a range of values // examples: // repeated_range([0, 1, 2, 3], 1) -> [0, 1, 2, 3] @@ -69,11 +67,7 @@ protected: int main() { - thrust::device_vector data(4); - data[0] = 10; - data[1] = 20; - data[2] = 30; - data[3] = 40; + thrust::device_vector data{10, 20, 30, 40}; // print the initial data std::cout << "range "; diff --git a/thrust/examples/saxpy.cu b/thrust/examples/saxpy.cu index d336312920c1..48979b24068d 100644 --- a/thrust/examples/saxpy.cu +++ b/thrust/examples/saxpy.cu @@ -7,8 +7,6 @@ #include #include -#include "include/host_device.h" - // This example illustrates how to implement the SAXPY // operation (Y[i] = a * X[i] + Y[i]) using Thrust. // The saxpy_slow function demonstrates the most @@ -55,13 +53,13 @@ void saxpy_slow(float A, thrust::device_vector& X, thrust::device_vector< int main() { // initialize host arrays - float x[4] = {1.0, 1.0, 1.0, 1.0}; - float y[4] = {1.0, 2.0, 3.0, 4.0}; + thrust::host_vector x{1.0, 1.0, 1.0, 1.0}; + thrust::host_vector y{1.0, 2.0, 3.0, 4.0}; { // transfer to device - thrust::device_vector X(x, x + 4); - thrust::device_vector Y(y, y + 4); + thrust::device_vector X(x); + thrust::device_vector Y(y); // slow method saxpy_slow(2.0, X, Y); @@ -69,8 +67,8 @@ int main() { // transfer to device - thrust::device_vector X(x, x + 4); - thrust::device_vector Y(y, y + 4); + thrust::device_vector X(x); + thrust::device_vector Y(y); // fast method saxpy_fast(2.0, X, Y); diff --git a/thrust/examples/scan_matrix_by_rows.cu b/thrust/examples/scan_matrix_by_rows.cu index d19991553fdc..eaef3714604e 100644 --- a/thrust/examples/scan_matrix_by_rows.cu +++ b/thrust/examples/scan_matrix_by_rows.cu @@ -6,8 +6,6 @@ #include -#include "include/host_device.h" - // We have a matrix stored in a `thrust::device_vector`. We want to perform a // scan on each row of a matrix. @@ -16,7 +14,7 @@ __host__ void scan_matrix_by_rows0(thrust::device_vector& u, int n, int m) // Here, we launch a separate scan for each row in the matrix. This works, // but each kernel only does a small amount of work. It would be better if we // could launch one big kernel for the entire matrix. - for (int i = 0; i < n; ++i) + for (size_t i = 0; i < static_cast(n); ++i) { thrust::inclusive_scan(u.begin() + m * i, u.begin() + m * (i + 1), u.begin() + m * i); } @@ -71,9 +69,9 @@ int main() thrust::sequence(u1.begin(), u1.end()); scan_matrix_by_rows1(u1, n, m); - for (int i = 0; i < n; ++i) + for (size_t i = 0; i < static_cast(n); ++i) { - for (int j = 0; j < m; ++j) + for (size_t j = 0; j < static_cast(m); ++j) { assert(u0[j + m * i] == u1[j + m * i]); } diff --git a/thrust/examples/simple_moving_average.cu b/thrust/examples/simple_moving_average.cu index c1ed4fc17171..f61cd6c1185d 100644 --- a/thrust/examples/simple_moving_average.cu +++ b/thrust/examples/simple_moving_average.cu @@ -8,8 +8,6 @@ #include #include -#include "include/host_device.h" - // Efficiently computes the simple moving average (SMA) [1] of a data series // using a parallel prefix-sum or "scan" operation. // @@ -71,10 +69,7 @@ int main() thrust::device_vector data(n); thrust::default_random_engine rng; thrust::uniform_int_distribution dist(0, 10); - for (size_t i = 0; i < n; i++) - { - data[i] = static_cast(dist(rng)); - } + std::generate(data.begin(), data.end(), [&]() { return static_cast(dist(rng)); }); // allocate storage for averages thrust::device_vector averages(data.size() - (w - 1)); @@ -84,9 +79,9 @@ int main() // print data series std::cout << "data series: [ "; - for (size_t i = 0; i < data.size(); i++) + for (const auto& value : data) { - std::cout << data[i] << " "; + std::cout << value << " "; } std::cout << "]" << std::endl; diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index 75ade90ef254..0d24e1f83b89 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -2,10 +2,10 @@ #include #include +#include #include #include - -#include "include/host_device.h" +#include // Helper routines @@ -13,68 +13,57 @@ void initialize(thrust::device_vector& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); - for (size_t i = 0; i < v.size(); i++) - { - v[i] = dist(rng); - } + std::generate(v.begin(), v.end(), [&]() { return dist(rng); }); } void initialize(thrust::device_vector& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(2, 19); - for (size_t i = 0; i < v.size(); i++) - { - v[i] = dist(rng) / 2.0f; - } + std::generate(v.begin(), v.end(), [&]() { return dist(rng) / 2.0f; }); } void initialize(thrust::device_vector>& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(0, 9); - for (size_t i = 0; i < v.size(); i++) - { + std::generate(v.begin(), v.end(), [&]() { int a = dist(rng); int b = dist(rng); - v[i] = thrust::make_pair(a, b); - } + return thrust::make_pair(a, b); + }); } void initialize(thrust::device_vector& v1, thrust::device_vector& v2) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); - for (size_t i = 0; i < v1.size(); i++) - { - v1[i] = dist(rng); - v2[i] = static_cast(i); - } + std::generate(v1.begin(), v1.end(), [&]() { return dist(rng); }); + std::iota(v2.begin(), v2.end(), 0); } void print(const thrust::device_vector& v) { - for (size_t i = 0; i < v.size(); i++) + for (const auto& value : v) { - std::cout << " " << v[i]; + std::cout << " " << value; } std::cout << "\n"; } void print(const thrust::device_vector& v) { - for (size_t i = 0; i < v.size(); i++) + for (const auto& value : v) { - std::cout << " " << std::fixed << std::setprecision(1) << v[i]; + std::cout << " " << std::fixed << std::setprecision(1) << value; } std::cout << "\n"; } void print(const thrust::device_vector>& v) { - for (size_t i = 0; i < v.size(); i++) + for (const auto& p : v) { - thrust::pair p = v[i]; std::cout << " (" << p.first << "," << p.second << ")"; } std::cout << "\n"; diff --git a/thrust/examples/sorting_aos_vs_soa.cu b/thrust/examples/sorting_aos_vs_soa.cu index 4532d0c657e9..857bd0dab4e2 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -3,9 +3,9 @@ #include #include +#include #include -#include "include/host_device.h" #include "include/timer.h" // This examples compares sorting performance using Array of Structures (AoS) @@ -35,10 +35,7 @@ void initialize_keys(thrust::device_vector& keys) thrust::host_vector h_keys(keys.size()); - for (size_t i = 0; i < h_keys.size(); i++) - { - h_keys[i] = dist(rng); - } + std::generate(h_keys.begin(), h_keys.end(), [&]() { return dist(rng); }); keys = h_keys; } @@ -50,10 +47,11 @@ void initialize_keys(thrust::device_vector& structures) thrust::host_vector h_structures(structures.size()); - for (size_t i = 0; i < h_structures.size(); i++) - { - h_structures[i].key = dist(rng); - } + std::generate(h_structures.begin(), h_structures.end(), [&]() { + MyStruct s; + s.key = dist(rng); + return s; + }); structures = h_structures; } diff --git a/thrust/examples/stream_compaction.cu b/thrust/examples/stream_compaction.cu index ace71894c681..4cdaa9fb6d3f 100644 --- a/thrust/examples/stream_compaction.cu +++ b/thrust/examples/stream_compaction.cu @@ -8,8 +8,6 @@ #include #include -#include "include/host_device.h" - // this functor returns true if the argument is odd, and false otherwise template struct is_odd diff --git a/thrust/examples/sum_rows.cu b/thrust/examples/sum_rows.cu index 718355a22dfe..937a393dbb56 100644 --- a/thrust/examples/sum_rows.cu +++ b/thrust/examples/sum_rows.cu @@ -5,10 +5,9 @@ #include #include +#include #include -#include "include/host_device.h" - // convert a linear index to a row index template struct linear_index_to_row_index @@ -34,10 +33,7 @@ int main() // initialize data thrust::device_vector array(R * C); - for (size_t i = 0; i < array.size(); i++) - { - array[i] = dist(rng); - } + std::generate(array.begin(), array.end(), [&]() { return dist(rng); }); // allocate storage for row sums and indices thrust::device_vector row_sums(R); @@ -54,10 +50,10 @@ int main() ::cuda::std::plus()); // print data - for (int i = 0; i < R; i++) + for (size_t i = 0; i < static_cast(R); i++) { std::cout << "[ "; - for (int j = 0; j < C; j++) + for (size_t j = 0; j < static_cast(C); j++) { std::cout << array[i * C + j] << " "; } diff --git a/thrust/examples/summary_statistics.cu b/thrust/examples/summary_statistics.cu index 14674ef3f446..de7ef8b07772 100644 --- a/thrust/examples/summary_statistics.cu +++ b/thrust/examples/summary_statistics.cu @@ -8,8 +8,6 @@ #include #include -#include "include/host_device.h" - // This example computes several statistical properties of a data // series in a single reduction. The algorithm is described in detail here: // http://en.wikipedia.org/wiki/Algorithms_for_calculating_variance#Parallel_algorithm @@ -135,10 +133,10 @@ int main() using T = float; // initialize host array - T h_x[] = {4, 7, 13, 16}; + thrust::host_vector h_x{4, 7, 13, 16}; // transfer to device - thrust::device_vector d_x(h_x, h_x + sizeof(h_x) / sizeof(T)); + thrust::device_vector d_x(h_x); // setup arguments summary_stats_unary_op unary_op; diff --git a/thrust/examples/summed_area_table.cu b/thrust/examples/summed_area_table.cu index 80e958afc007..8946d080dd67 100644 --- a/thrust/examples/summed_area_table.cu +++ b/thrust/examples/summed_area_table.cu @@ -9,8 +9,6 @@ #include #include -#include "include/host_device.h" - // This example computes a summed area table using segmented scan // http://en.wikipedia.org/wiki/Summed_area_table diff --git a/thrust/examples/tiled_range.cu b/thrust/examples/tiled_range.cu index 7f67a409cb47..ac27e85d537f 100644 --- a/thrust/examples/tiled_range.cu +++ b/thrust/examples/tiled_range.cu @@ -8,8 +8,6 @@ #include -#include "include/host_device.h" - // this example illustrates how to tile a range multiple times // examples: // tiled_range([0, 1, 2, 3], 1) -> [0, 1, 2, 3] @@ -69,11 +67,7 @@ protected: int main() { - thrust::device_vector data(4); - data[0] = 10; - data[1] = 20; - data[2] = 30; - data[3] = 40; + thrust::device_vector data{10, 20, 30, 40}; // print the initial data std::cout << "range "; diff --git a/thrust/examples/transform_input_output_iterator.cu b/thrust/examples/transform_input_output_iterator.cu index a9630014d3f5..f02bf01872b3 100644 --- a/thrust/examples/transform_input_output_iterator.cu +++ b/thrust/examples/transform_input_output_iterator.cu @@ -7,8 +7,6 @@ #include -#include "include/host_device.h" - // Base 2 fixed point class ScaledInteger { @@ -109,9 +107,9 @@ int main() std::cout << "] \n"; std::cout << "Result [ "; - for (size_t i = 0; i < size; i++) + for (const auto& value : C_h) { - std::cout << C_h[i] << " "; + std::cout << value << " "; } std::cout << "] \n"; diff --git a/thrust/examples/transform_iterator.cu b/thrust/examples/transform_iterator.cu index 3288058fa843..e710b898427f 100644 --- a/thrust/examples/transform_iterator.cu +++ b/thrust/examples/transform_iterator.cu @@ -8,8 +8,6 @@ #include #include -#include "include/host_device.h" - // this functor clamps a value to the range [lo, hi] template struct clamp diff --git a/thrust/examples/transform_output_iterator.cu b/thrust/examples/transform_output_iterator.cu index 840b35939c44..c62da242882f 100644 --- a/thrust/examples/transform_output_iterator.cu +++ b/thrust/examples/transform_output_iterator.cu @@ -5,7 +5,6 @@ #include -#include "include/host_device.h" struct Functor { template @@ -19,15 +18,15 @@ struct Functor int main() { - float u[4] = {4, 3, 2, 1}; - float v[4] = {-1, 1, 1, -1}; - int idx[3] = {3, 0, 1}; - float w[3] = {0, 0, 0}; + thrust::host_vector u{4, 3, 2, 1}; + thrust::host_vector v{-1, 1, 1, -1}; + thrust::host_vector idx{3, 0, 1}; + thrust::host_vector w{0, 0, 0}; - thrust::device_vector U(u, u + 4); - thrust::device_vector V(v, v + 4); - thrust::device_vector IDX(idx, idx + 3); - thrust::device_vector W(w, w + 3); + thrust::device_vector U(u); + thrust::device_vector V(v); + thrust::device_vector IDX(idx); + thrust::device_vector W(w); // gather multiple elements and apply a function before writing result in memory thrust::gather(IDX.begin(), @@ -36,9 +35,9 @@ int main() thrust::make_transform_output_iterator(W.begin(), Functor())); std::cout << "result= [ "; - for (size_t i = 0; i < 3; i++) + for (const auto& value : W) { - std::cout << W[i] << " "; + std::cout << value << " "; } std::cout << "] \n"; diff --git a/thrust/examples/uninitialized_vector.cu b/thrust/examples/uninitialized_vector.cu index 3f5a255657c3..a76cbb3228f6 100644 --- a/thrust/examples/uninitialized_vector.cu +++ b/thrust/examples/uninitialized_vector.cu @@ -11,8 +11,6 @@ #include -#include "include/host_device.h" - // uninitialized_allocator is an allocator which // derives from device_allocator and which has a // no-op construct member function diff --git a/thrust/examples/weld_vertices.cu b/thrust/examples/weld_vertices.cu index ff1bff568c2b..87c982d0014d 100644 --- a/thrust/examples/weld_vertices.cu +++ b/thrust/examples/weld_vertices.cu @@ -41,15 +41,12 @@ int main() // allocate memory for input mesh representation thrust::device_vector input(9); - input[0] = vec2(0, 0); // First Triangle - input[1] = vec2(1, 0); - input[2] = vec2(0, 1); - input[3] = vec2(1, 0); // Second Triangle - input[4] = vec2(1, 1); - input[5] = vec2(0, 1); - input[6] = vec2(1, 0); // Third Triangle - input[7] = vec2(2, 0); - input[8] = vec2(1, 1); + thrust::host_vector h_input{ + vec2(0, 0), vec2(1, 0), vec2(0, 1), // First Triangle + vec2(1, 0), vec2(1, 1), vec2(0, 1), // Second Triangle + vec2(1, 0), vec2(2, 0), vec2(1, 1) // Third Triangle + }; + input = h_input; // allocate space for output mesh representation thrust::device_vector vertices = input; diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 6f9f7fcdc6f2..346cce8c8670 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -5,8 +5,6 @@ #include -#include "include/host_device.h" - // 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. @@ -70,7 +68,8 @@ int main() std::cout << raw_input << std::endl; // transfer to device - thrust::device_vector input(raw_input, raw_input + sizeof(raw_input)); + thrust::host_vector h_input(raw_input, raw_input + sizeof(raw_input)); + thrust::device_vector input(h_input); // count words int wc = word_count(input); From e6147285387e6285256dd51bb0fbbad70e4b5ee2 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Tue, 26 Aug 2025 22:03:34 -0500 Subject: [PATCH 02/19] removed debug comments --- thrust/examples/arbitrary_transformation.cu | 2 -- thrust/examples/bounding_box.cu | 2 -- thrust/examples/bucket_sort2d.cu | 2 -- 3 files changed, 6 deletions(-) diff --git a/thrust/examples/arbitrary_transformation.cu b/thrust/examples/arbitrary_transformation.cu index 976b6826dffa..395bbfeb04c1 100644 --- a/thrust/examples/arbitrary_transformation.cu +++ b/thrust/examples/arbitrary_transformation.cu @@ -7,8 +7,6 @@ #include -// #include "include/host_device.h" // Not needed for this example - // This example shows how to implement an arbitrary transformation of // the form output[i] = F(first[i], second[i], third[i], ... ). // In this example, we use a function with 3 inputs and 1 output. diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index c2a2a14f4ca9..4b8c208c7d19 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -4,8 +4,6 @@ #include #include -// #include "include/host_device.h" // Not needed for this example - // This example shows how to compute a bounding box // for a set of points in two dimensions. diff --git a/thrust/examples/bucket_sort2d.cu b/thrust/examples/bucket_sort2d.cu index eb00d0170953..926e1d3a292d 100644 --- a/thrust/examples/bucket_sort2d.cu +++ b/thrust/examples/bucket_sort2d.cu @@ -9,8 +9,6 @@ #include #include -// #include "include/host_device.h" // Not needed for this example - // define a 2d float vector using vec2 = thrust::tuple; From b6efb9bc3db1a9f7d056d3a1fc046cc98cae8df1 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Tue, 26 Aug 2025 22:29:04 -0500 Subject: [PATCH 03/19] clang format --- thrust/examples/bounding_box.cu | 2 +- thrust/examples/dot_products_with_zip.cu | 5 +++-- thrust/examples/lambda.cu | 3 +-- thrust/examples/lexicographical_sort.cu | 4 +++- thrust/examples/minmax.cu | 5 +++-- thrust/examples/mode.cu | 4 +++- thrust/examples/remove_points2d.cu | 8 ++++++-- thrust/examples/simple_moving_average.cu | 4 +++- thrust/examples/sort.cu | 12 +++++++++--- thrust/examples/sorting_aos_vs_soa.cu | 5 ++++- thrust/examples/sum_rows.cu | 4 +++- thrust/examples/weld_vertices.cu | 12 +++++++++--- 12 files changed, 48 insertions(+), 20 deletions(-) diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index 4b8c208c7d19..4de8f6188404 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -52,7 +52,7 @@ struct bbox }; // reduce a pair of bounding boxes (a,b) to a bounding box containing a and b -struct bbox_union : public thrust::binary_function +struct bbox_union : public thrust::binary_function { __host__ __device__ bbox operator()(bbox a, bbox b) { diff --git a/thrust/examples/dot_products_with_zip.cu b/thrust/examples/dot_products_with_zip.cu index 4b7a16e01c8b..3764e1eef88e 100644 --- a/thrust/examples/dot_products_with_zip.cu +++ b/thrust/examples/dot_products_with_zip.cu @@ -5,7 +5,6 @@ #include #include - // This example shows how thrust::zip_iterator can be used to create a // 'virtual' array of structures. In this case the structure is a 3d // vector type (Float3) whose (x,y,z) components will be stored in @@ -32,7 +31,9 @@ thrust::host_vector random_vector(const size_t N, unsigned int seed = thr thrust::default_random_engine rng(seed); thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector temp(N); - std::generate(temp.begin(), temp.end(), [&]() { return u01(rng); }); + std::generate(temp.begin(), temp.end(), [&]() { + return u01(rng); + }); return temp; } diff --git a/thrust/examples/lambda.cu b/thrust/examples/lambda.cu index 4cf098435704..2b8fde4cf5de 100644 --- a/thrust/examples/lambda.cu +++ b/thrust/examples/lambda.cu @@ -4,7 +4,6 @@ #include - // This example demonstrates the use of placeholders to implement // the SAXPY operation (i.e. Y[i] = a * X[i] + Y[i]). // @@ -40,7 +39,7 @@ struct saxpy_functor int main() { // input data - float a = 2.0f; + float a = 2.0f; thrust::device_vector x_data = {1, 2, 3, 4}; thrust::device_vector y_data = {1, 1, 1, 1}; diff --git a/thrust/examples/lexicographical_sort.cu b/thrust/examples/lexicographical_sort.cu index 3aa608e0acfa..38cb8f26ef90 100644 --- a/thrust/examples/lexicographical_sort.cu +++ b/thrust/examples/lexicographical_sort.cu @@ -41,7 +41,9 @@ thrust::host_vector random_vector(size_t N) static thrust::default_random_engine rng; static thrust::uniform_int_distribution dist(0, 9); - std::generate(vec.begin(), vec.end(), [&]() { return dist(rng); }); + std::generate(vec.begin(), vec.end(), [&]() { + return dist(rng); + }); return vec; } diff --git a/thrust/examples/minmax.cu b/thrust/examples/minmax.cu index a412336ec2d4..be4d3972da4e 100644 --- a/thrust/examples/minmax.cu +++ b/thrust/examples/minmax.cu @@ -5,7 +5,6 @@ #include #include - // compute minimum and maximum values in a single reduction // minmax_pair stores the minimum and maximum @@ -59,7 +58,9 @@ int main() // initialize data on host thrust::device_vector data(N); - std::generate(data.begin(), data.end(), [&]() { return dist(rng); }); + std::generate(data.begin(), data.end(), [&]() { + return dist(rng); + }); // setup arguments minmax_unary_op unary_op; diff --git a/thrust/examples/mode.cu b/thrust/examples/mode.cu index a45f3c29857a..41b4949bf4dc 100644 --- a/thrust/examples/mode.cu +++ b/thrust/examples/mode.cu @@ -26,7 +26,9 @@ int main() // generate random data on the host thrust::host_vector h_data(N); - std::generate(h_data.begin(), h_data.end(), [&]() { return dist(rng); }); + std::generate(h_data.begin(), h_data.end(), [&]() { + return dist(rng); + }); // transfer data to device thrust::device_vector d_data(h_data); diff --git a/thrust/examples/remove_points2d.cu b/thrust/examples/remove_points2d.cu index 9af8c2cb2a51..cba2eb92d792 100644 --- a/thrust/examples/remove_points2d.cu +++ b/thrust/examples/remove_points2d.cu @@ -39,8 +39,12 @@ int main() thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector x(N); thrust::host_vector y(N); - std::generate(x.begin(), x.end(), [&]() { return u01(rng); }); - std::generate(y.begin(), y.end(), [&]() { return u01(rng); }); + std::generate(x.begin(), x.end(), [&]() { + return u01(rng); + }); + std::generate(y.begin(), y.end(), [&]() { + return u01(rng); + }); // print the initial points std::cout << std::fixed; diff --git a/thrust/examples/simple_moving_average.cu b/thrust/examples/simple_moving_average.cu index f61cd6c1185d..9be723045992 100644 --- a/thrust/examples/simple_moving_average.cu +++ b/thrust/examples/simple_moving_average.cu @@ -69,7 +69,9 @@ int main() thrust::device_vector data(n); thrust::default_random_engine rng; thrust::uniform_int_distribution dist(0, 10); - std::generate(data.begin(), data.end(), [&]() { return static_cast(dist(rng)); }); + std::generate(data.begin(), data.end(), [&]() { + return static_cast(dist(rng)); + }); // allocate storage for averages thrust::device_vector averages(data.size() - (w - 1)); diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index 0d24e1f83b89..e9695649fbf3 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -13,14 +13,18 @@ void initialize(thrust::device_vector& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); - std::generate(v.begin(), v.end(), [&]() { return dist(rng); }); + std::generate(v.begin(), v.end(), [&]() { + return dist(rng); + }); } void initialize(thrust::device_vector& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(2, 19); - std::generate(v.begin(), v.end(), [&]() { return dist(rng) / 2.0f; }); + std::generate(v.begin(), v.end(), [&]() { + return dist(rng) / 2.0f; + }); } void initialize(thrust::device_vector>& v) @@ -38,7 +42,9 @@ void initialize(thrust::device_vector& v1, thrust::device_vector& v2) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); - std::generate(v1.begin(), v1.end(), [&]() { return dist(rng); }); + std::generate(v1.begin(), v1.end(), [&]() { + return dist(rng); + }); std::iota(v2.begin(), v2.end(), 0); } diff --git a/thrust/examples/sorting_aos_vs_soa.cu b/thrust/examples/sorting_aos_vs_soa.cu index 857bd0dab4e2..2a7c8a8e62ff 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -4,6 +4,7 @@ #include #include + #include #include "include/timer.h" @@ -35,7 +36,9 @@ void initialize_keys(thrust::device_vector& keys) thrust::host_vector h_keys(keys.size()); - std::generate(h_keys.begin(), h_keys.end(), [&]() { return dist(rng); }); + std::generate(h_keys.begin(), h_keys.end(), [&]() { + return dist(rng); + }); keys = h_keys; } diff --git a/thrust/examples/sum_rows.cu b/thrust/examples/sum_rows.cu index 937a393dbb56..c5382cab3691 100644 --- a/thrust/examples/sum_rows.cu +++ b/thrust/examples/sum_rows.cu @@ -33,7 +33,9 @@ int main() // initialize data thrust::device_vector array(R * C); - std::generate(array.begin(), array.end(), [&]() { return dist(rng); }); + std::generate(array.begin(), array.end(), [&]() { + return dist(rng); + }); // allocate storage for row sums and indices thrust::device_vector row_sums(R); diff --git a/thrust/examples/weld_vertices.cu b/thrust/examples/weld_vertices.cu index 87c982d0014d..1e13afbb1f22 100644 --- a/thrust/examples/weld_vertices.cu +++ b/thrust/examples/weld_vertices.cu @@ -42,9 +42,15 @@ int main() thrust::device_vector input(9); thrust::host_vector h_input{ - vec2(0, 0), vec2(1, 0), vec2(0, 1), // First Triangle - vec2(1, 0), vec2(1, 1), vec2(0, 1), // Second Triangle - vec2(1, 0), vec2(2, 0), vec2(1, 1) // Third Triangle + vec2(0, 0), + vec2(1, 0), + vec2(0, 1), // First Triangle + vec2(1, 0), + vec2(1, 1), + vec2(0, 1), // Second Triangle + vec2(1, 0), + vec2(2, 0), + vec2(1, 1) // Third Triangle }; input = h_input; From 4731c69f2a536c5bca9aa30ba51556ccde9ea543 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Thu, 28 Aug 2025 22:12:07 -0500 Subject: [PATCH 04/19] Modernized to thrust::generate, cuda::std::distance, thrust::sequence where necessary --- thrust/examples/arbitrary_transformation.cu | 2 +- thrust/examples/basic_vector.cu | 2 +- thrust/examples/bounding_box.cu | 30 ++++++++++++++------ thrust/examples/counting_iterator.cu | 2 +- thrust/examples/device_ptr.cu | 2 +- thrust/examples/dot_products_with_zip.cu | 2 +- thrust/examples/expand.cu | 2 +- thrust/examples/histogram.cu | 2 +- thrust/examples/lexicographical_sort.cu | 2 +- thrust/examples/max_abs_diff.cu | 4 +-- thrust/examples/minmax.cu | 5 ++-- thrust/examples/mode.cu | 4 +-- thrust/examples/remove_points2d.cu | 4 +-- thrust/examples/scan_by_key.cu | 2 -- thrust/examples/simple_moving_average.cu | 6 ++-- thrust/examples/sort.cu | 22 ++++++++++---- thrust/examples/sorting_aos_vs_soa.cu | 4 +-- thrust/examples/strided_range.cu | 2 -- thrust/examples/sum_rows.cu | 5 ++-- thrust/examples/transform_output_iterator.cu | 1 + thrust/examples/weld_vertices.cu | 1 + thrust/examples/word_count.cu | 1 + 22 files changed, 65 insertions(+), 42 deletions(-) diff --git a/thrust/examples/arbitrary_transformation.cu b/thrust/examples/arbitrary_transformation.cu index 395bbfeb04c1..e16be0f731d6 100644 --- a/thrust/examples/arbitrary_transformation.cu +++ b/thrust/examples/arbitrary_transformation.cu @@ -60,7 +60,7 @@ struct arbitrary_functor2 int main() { - // allocate and initialize storage using initializer lists + // allocate and initialize thrust::device_vector A{3, 4, 0, 8, 2}; thrust::device_vector B{6, 7, 2, 1, 8}; thrust::device_vector C{2, 5, 7, 4, 3}; diff --git a/thrust/examples/basic_vector.cu b/thrust/examples/basic_vector.cu index 530e06680cb9..f2cc84925f5b 100644 --- a/thrust/examples/basic_vector.cu +++ b/thrust/examples/basic_vector.cu @@ -5,7 +5,7 @@ int main() { - // H has storage for 4 integers with direct initialization + // H holds 4 integers thrust::host_vector H{14, 20, 38, 46}; // H.size() returns the size of vector H diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index 4de8f6188404..ddadc08a789c 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -1,5 +1,6 @@ #include #include +#include #include #include #include @@ -66,27 +67,38 @@ struct bbox_union : public thrust::binary_function } }; +// functor to generate random points +struct random_point_generator +{ + mutable unsigned int seed; + + random_point_generator() + : seed(0) + {} + + __host__ __device__ point2d operator()() const + { + thrust::default_random_engine rng(seed++); + thrust::uniform_real_distribution u01(0.0f, 1.0f); + return point2d(u01(rng), u01(rng)); + } +}; + int main() { const size_t N = 40; - thrust::default_random_engine rng; - thrust::uniform_real_distribution u01(0.0f, 1.0f); // allocate storage for points thrust::device_vector points(N); // generate some random points in the unit square - std::generate(points.begin(), points.end(), [&] { - auto x = u01(rng); - auto y = u01(rng); - return point2d(x, y); - }); + thrust::generate(points.begin(), points.end(), random_point_generator()); // initial bounding box contains first point - bbox first_point = bbox(points[0], points[0]); + bbox first_point{points[0], points[0]}; // compute the bounding box for the point set - bbox result = thrust::reduce(points.begin(), points.end(), first_point, bbox_union()); + bbox result = thrust::reduce(points.begin(), points.end(), first_point, bbox_union{}); // print output std::cout << "bounding box " << std::fixed; diff --git a/thrust/examples/counting_iterator.cu b/thrust/examples/counting_iterator.cu index 428db8b13cad..48c4ac90810c 100644 --- a/thrust/examples/counting_iterator.cu +++ b/thrust/examples/counting_iterator.cu @@ -27,7 +27,7 @@ int main() // indices now contains [1,2,5,7] // print result - std::cout << "found " << std::distance(indices.begin(), indices_end) << " nonzero values at indices:\n"; + std::cout << "found " << cuda::std::distance(indices.begin(), indices_end) << " nonzero values at indices:\n"; thrust::copy(indices.begin(), indices_end, std::ostream_iterator(std::cout, "\n")); return 0; diff --git a/thrust/examples/device_ptr.cu b/thrust/examples/device_ptr.cu index 3fcb555be23d..c165f4261365 100644 --- a/thrust/examples/device_ptr.cu +++ b/thrust/examples/device_ptr.cu @@ -15,7 +15,7 @@ int main() // device_ptr supports pointer arithmetic thrust::device_ptr first = d_ptr; thrust::device_ptr last = d_ptr + 10; - std::cout << "device array contains " << std::distance(first, last) << " values\n"; + std::cout << "device array contains " << cuda::std::distance(first, last) << " values\n"; // algorithms work as expected thrust::sequence(first, last); diff --git a/thrust/examples/dot_products_with_zip.cu b/thrust/examples/dot_products_with_zip.cu index 3764e1eef88e..b5fbed13f5d8 100644 --- a/thrust/examples/dot_products_with_zip.cu +++ b/thrust/examples/dot_products_with_zip.cu @@ -31,7 +31,7 @@ thrust::host_vector random_vector(const size_t N, unsigned int seed = thr thrust::default_random_engine rng(seed); thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector temp(N); - std::generate(temp.begin(), temp.end(), [&]() { + thrust::generate(temp.begin(), temp.end(), [&]() { return u01(rng); }); return temp; diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index a979d0846eef..053ae1cc73aa 100644 --- a/thrust/examples/expand.cu +++ b/thrust/examples/expand.cu @@ -40,7 +40,7 @@ OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator // compute max-scan over the output indices, filling in the holes thrust::inclusive_scan( - output_indices.begin(), output_indices.end(), output_indices.begin(), ::cuda::maximum()); + output_indices.begin(), output_indices.end(), output_indices.begin(), thrust::maximum()); // gather input values according to index array (output = first2[output_indices]) thrust::gather(output_indices.begin(), output_indices.end(), first2, output); diff --git a/thrust/examples/histogram.cu b/thrust/examples/histogram.cu index 8141e958b773..2e76b2791fb9 100644 --- a/thrust/examples/histogram.cu +++ b/thrust/examples/histogram.cu @@ -149,7 +149,7 @@ int main() // generate random data on the host thrust::host_vector input(N); - std::generate(input.begin(), input.end(), [&]() { + thrust::generate(input.begin(), input.end(), [&]() { int sum = 0; for (int j = 0; j < S; j++) { diff --git a/thrust/examples/lexicographical_sort.cu b/thrust/examples/lexicographical_sort.cu index 38cb8f26ef90..9432fd05a615 100644 --- a/thrust/examples/lexicographical_sort.cu +++ b/thrust/examples/lexicographical_sort.cu @@ -41,7 +41,7 @@ thrust::host_vector random_vector(size_t N) static thrust::default_random_engine rng; static thrust::uniform_int_distribution dist(0, 9); - std::generate(vec.begin(), vec.end(), [&]() { + thrust::generate(vec.begin(), vec.end(), [&]() { return dist(rng); }); diff --git a/thrust/examples/max_abs_diff.cu b/thrust/examples/max_abs_diff.cu index 6f2c04177b34..7d829b23db5d 100644 --- a/thrust/examples/max_abs_diff.cu +++ b/thrust/examples/max_abs_diff.cu @@ -5,8 +5,6 @@ #include #include -// #include "include/host_device.h" // Legacy header - commented out - // this example computes the maximum absolute difference // between the elements of two vectors @@ -28,7 +26,7 @@ int main() float init = 0; // binary operations - ::cuda::maximum binary_op1; + thrust::maximum binary_op1; abs_diff binary_op2; float max_abs_diff = thrust::inner_product(d_a.begin(), d_a.end(), d_b.begin(), init, binary_op1, binary_op2); diff --git a/thrust/examples/minmax.cu b/thrust/examples/minmax.cu index be4d3972da4e..911cfe9b679c 100644 --- a/thrust/examples/minmax.cu +++ b/thrust/examples/minmax.cu @@ -57,10 +57,11 @@ int main() thrust::uniform_int_distribution dist(10, 99); // initialize data on host - thrust::device_vector data(N); - std::generate(data.begin(), data.end(), [&]() { + thrust::host_vector host_data(N); + thrust::generate(host_data.begin(), host_data.end(), [&]() { return dist(rng); }); + thrust::device_vector data = host_data; // setup arguments minmax_unary_op unary_op; diff --git a/thrust/examples/mode.cu b/thrust/examples/mode.cu index 41b4949bf4dc..bc390dff0af5 100644 --- a/thrust/examples/mode.cu +++ b/thrust/examples/mode.cu @@ -26,7 +26,7 @@ int main() // generate random data on the host thrust::host_vector h_data(N); - std::generate(h_data.begin(), h_data.end(), [&]() { + thrust::generate(h_data.begin(), h_data.end(), [&]() { return dist(rng); }); @@ -69,7 +69,7 @@ int main() thrust::device_vector::iterator mode_iter; mode_iter = thrust::max_element(d_output_counts.begin(), d_output_counts.end()); - int mode = d_output_keys[std::distance(d_output_counts.begin(), mode_iter)]; + int mode = d_output_keys[cuda::std::distance(d_output_counts.begin(), mode_iter)]; int occurrences = *mode_iter; std::cout << "Modal value " << mode << " occurs " << occurrences << " times " << std::endl; diff --git a/thrust/examples/remove_points2d.cu b/thrust/examples/remove_points2d.cu index cba2eb92d792..7be88c51ec33 100644 --- a/thrust/examples/remove_points2d.cu +++ b/thrust/examples/remove_points2d.cu @@ -39,10 +39,10 @@ int main() thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector x(N); thrust::host_vector y(N); - std::generate(x.begin(), x.end(), [&]() { + thrust::generate(x.begin(), x.end(), [&]() { return u01(rng); }); - std::generate(y.begin(), y.end(), [&]() { + thrust::generate(y.begin(), y.end(), [&]() { return u01(rng); }); diff --git a/thrust/examples/scan_by_key.cu b/thrust/examples/scan_by_key.cu index d275362da9eb..9a8128db2e7e 100644 --- a/thrust/examples/scan_by_key.cu +++ b/thrust/examples/scan_by_key.cu @@ -4,8 +4,6 @@ #include -#include "include/host_device.h" - // BinaryPredicate for the head flag segment representation // equivalent to ::cuda::std::not_fn(thrust::project2nd())); template diff --git a/thrust/examples/simple_moving_average.cu b/thrust/examples/simple_moving_average.cu index 9be723045992..0e47d15651f2 100644 --- a/thrust/examples/simple_moving_average.cu +++ b/thrust/examples/simple_moving_average.cu @@ -1,5 +1,6 @@ #include #include +#include #include #include #include @@ -66,12 +67,13 @@ int main() size_t w = 4; // generate random data series - thrust::device_vector data(n); + thrust::host_vector host_data(n); thrust::default_random_engine rng; thrust::uniform_int_distribution dist(0, 10); - std::generate(data.begin(), data.end(), [&]() { + thrust::generate(host_data.begin(), host_data.end(), [&]() { return static_cast(dist(rng)); }); + thrust::device_vector data = host_data; // allocate storage for averages thrust::device_vector averages(data.size() - (w - 1)); diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index e9695649fbf3..58a3eae9b7e0 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -1,4 +1,5 @@ #include +#include #include #include @@ -13,39 +14,47 @@ void initialize(thrust::device_vector& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); - std::generate(v.begin(), v.end(), [&]() { + thrust::host_vector host_data(v.size()); + thrust::generate(host_data.begin(), host_data.end(), [&]() { return dist(rng); }); + v = host_data; } void initialize(thrust::device_vector& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(2, 19); - std::generate(v.begin(), v.end(), [&]() { + thrust::host_vector host_data(v.size()); + thrust::generate(host_data.begin(), host_data.end(), [&]() { return dist(rng) / 2.0f; }); + v = host_data; } void initialize(thrust::device_vector>& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(0, 9); - std::generate(v.begin(), v.end(), [&]() { + thrust::host_vector> host_data(v.size()); + thrust::generate(host_data.begin(), host_data.end(), [&]() { int a = dist(rng); int b = dist(rng); return thrust::make_pair(a, b); }); + v = host_data; } void initialize(thrust::device_vector& v1, thrust::device_vector& v2) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); - std::generate(v1.begin(), v1.end(), [&]() { + thrust::host_vector host_data(v1.size()); + thrust::generate(host_data.begin(), host_data.end(), [&]() { return dist(rng); }); - std::iota(v2.begin(), v2.end(), 0); + v1 = host_data; + thrust::sequence(v2.begin(), v2.end(), 0); } void print(const thrust::device_vector& v) @@ -70,7 +79,8 @@ void print(const thrust::device_vector>& v) { for (const auto& p : v) { - std::cout << " (" << p.first << "," << p.second << ")"; + thrust::pair local_p = p; + std::cout << " (" << local_p.first << "," << local_p.second << ")"; } std::cout << "\n"; } diff --git a/thrust/examples/sorting_aos_vs_soa.cu b/thrust/examples/sorting_aos_vs_soa.cu index 2a7c8a8e62ff..201c37419146 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -36,7 +36,7 @@ void initialize_keys(thrust::device_vector& keys) thrust::host_vector h_keys(keys.size()); - std::generate(h_keys.begin(), h_keys.end(), [&]() { + thrust::generate(h_keys.begin(), h_keys.end(), [&]() { return dist(rng); }); @@ -50,7 +50,7 @@ void initialize_keys(thrust::device_vector& structures) thrust::host_vector h_structures(structures.size()); - std::generate(h_structures.begin(), h_structures.end(), [&]() { + thrust::generate(h_structures.begin(), h_structures.end(), [&]() { MyStruct s; s.key = dist(rng); return s; diff --git a/thrust/examples/strided_range.cu b/thrust/examples/strided_range.cu index 7389e1f0ae3e..7ad9fd390aa7 100644 --- a/thrust/examples/strided_range.cu +++ b/thrust/examples/strided_range.cu @@ -8,8 +8,6 @@ #include -#include "include/host_device.h" - // this example illustrates how to make strided access to a range of values // examples: // strided_range([0, 1, 2, 3, 4, 5, 6], 1) -> [0, 1, 2, 3, 4, 5, 6] diff --git a/thrust/examples/sum_rows.cu b/thrust/examples/sum_rows.cu index c5382cab3691..16c127974eae 100644 --- a/thrust/examples/sum_rows.cu +++ b/thrust/examples/sum_rows.cu @@ -32,10 +32,11 @@ int main() thrust::uniform_int_distribution dist(10, 99); // initialize data - thrust::device_vector array(R * C); - std::generate(array.begin(), array.end(), [&]() { + thrust::host_vector host_array(R * C); + thrust::generate(host_array.begin(), host_array.end(), [&]() { return dist(rng); }); + thrust::device_vector array = host_array; // allocate storage for row sums and indices thrust::device_vector row_sums(R); diff --git a/thrust/examples/transform_output_iterator.cu b/thrust/examples/transform_output_iterator.cu index c62da242882f..3f6af6aa5558 100644 --- a/thrust/examples/transform_output_iterator.cu +++ b/thrust/examples/transform_output_iterator.cu @@ -1,6 +1,7 @@ #include #include #include +#include #include #include diff --git a/thrust/examples/weld_vertices.cu b/thrust/examples/weld_vertices.cu index 1e13afbb1f22..08fc2fd32076 100644 --- a/thrust/examples/weld_vertices.cu +++ b/thrust/examples/weld_vertices.cu @@ -1,5 +1,6 @@ #include #include +#include #include #include #include diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index 346cce8c8670..c4cc7e15f952 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -1,5 +1,6 @@ #include #include +#include #include #include From 60238f3d1b57a84b19604cba27b3bc49a4fcd6d2 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Fri, 29 Aug 2025 02:24:39 -0500 Subject: [PATCH 05/19] moved to member variables --- thrust/examples/bounding_box.cu | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index ddadc08a789c..549488f2c410 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -1,6 +1,5 @@ #include #include -#include #include #include #include @@ -70,16 +69,16 @@ struct bbox_union : public thrust::binary_function // functor to generate random points struct random_point_generator { - mutable unsigned int seed; + mutable thrust::default_random_engine rng; + mutable thrust::uniform_real_distribution u01; random_point_generator() - : seed(0) + : rng(0) + , u01(0.0f, 1.0f) {} __host__ __device__ point2d operator()() const { - thrust::default_random_engine rng(seed++); - thrust::uniform_real_distribution u01(0.0f, 1.0f); return point2d(u01(rng), u01(rng)); } }; @@ -95,10 +94,10 @@ int main() thrust::generate(points.begin(), points.end(), random_point_generator()); // initial bounding box contains first point - bbox first_point{points[0], points[0]}; + bbox init = bbox(points[0], points[0]); // compute the bounding box for the point set - bbox result = thrust::reduce(points.begin(), points.end(), first_point, bbox_union{}); + bbox result = thrust::reduce(points.begin(), points.end(), init, bbox_union{}); // print output std::cout << "bounding box " << std::fixed; From 0c289758a23c202a36abc9ad69fc7a7a0c998992 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Fri, 29 Aug 2025 02:47:23 -0500 Subject: [PATCH 06/19] direct constructor syntax --- thrust/examples/bounding_box.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index 549488f2c410..81591f47f424 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -94,7 +94,7 @@ int main() thrust::generate(points.begin(), points.end(), random_point_generator()); // initial bounding box contains first point - bbox init = bbox(points[0], points[0]); + bbox init(points[0], points[0]); // compute the bounding box for the point set bbox result = thrust::reduce(points.begin(), points.end(), init, bbox_union{}); From ba240a932190a7062833999d11936d762d43dccf Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Fri, 29 Aug 2025 13:10:09 -0500 Subject: [PATCH 07/19] remove unused includes --- thrust/examples/sort.cu | 1 - thrust/examples/sorting_aos_vs_soa.cu | 2 -- thrust/examples/sum_rows.cu | 1 - 3 files changed, 4 deletions(-) diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index 58a3eae9b7e0..b8c8f740d48b 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -3,7 +3,6 @@ #include #include -#include #include #include #include diff --git a/thrust/examples/sorting_aos_vs_soa.cu b/thrust/examples/sorting_aos_vs_soa.cu index 201c37419146..91d189a9a004 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -3,8 +3,6 @@ #include #include -#include - #include #include "include/timer.h" diff --git a/thrust/examples/sum_rows.cu b/thrust/examples/sum_rows.cu index 16c127974eae..fb9fd1860051 100644 --- a/thrust/examples/sum_rows.cu +++ b/thrust/examples/sum_rows.cu @@ -5,7 +5,6 @@ #include #include -#include #include // convert a linear index to a row index From 38d677128745af792610dd7c2ac28ad69492c8a6 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Fri, 29 Aug 2025 14:44:25 -0500 Subject: [PATCH 08/19] additional fixes --- thrust/examples/bounding_box.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index 81591f47f424..f3743e19070e 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -28,7 +28,7 @@ struct bbox // construct an empty box __host__ __device__ bbox() {} - // construct a box from a single point + // construct a box from a single pointgit branch -a __host__ __device__ bbox(const point2d& point) : lower_left(point) , upper_right(point) From 027fc15326c52c8847b099d6ef5553abce87efb9 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Fri, 29 Aug 2025 14:54:30 -0500 Subject: [PATCH 09/19] fix warnings --- thrust/examples/expand.cu | 2 +- thrust/examples/max_abs_diff.cu | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index 053ae1cc73aa..c171cb6244a7 100644 --- a/thrust/examples/expand.cu +++ b/thrust/examples/expand.cu @@ -40,7 +40,7 @@ OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator // compute max-scan over the output indices, filling in the holes thrust::inclusive_scan( - output_indices.begin(), output_indices.end(), output_indices.begin(), thrust::maximum()); + output_indices.begin(), output_indices.end(), output_indices.begin(), cuda::maximum()); // gather input values according to index array (output = first2[output_indices]) thrust::gather(output_indices.begin(), output_indices.end(), first2, output); diff --git a/thrust/examples/max_abs_diff.cu b/thrust/examples/max_abs_diff.cu index 7d829b23db5d..ee98c2881430 100644 --- a/thrust/examples/max_abs_diff.cu +++ b/thrust/examples/max_abs_diff.cu @@ -26,7 +26,7 @@ int main() float init = 0; // binary operations - thrust::maximum binary_op1; + cuda::maximum binary_op1; abs_diff binary_op2; float max_abs_diff = thrust::inner_product(d_a.begin(), d_a.end(), d_b.begin(), init, binary_op1, binary_op2); From 88376b81c41060bf0bf4cbd361cdac9e54e5a273 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Fri, 29 Aug 2025 15:26:14 -0500 Subject: [PATCH 10/19] fix --- thrust/examples/bounding_box.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index f3743e19070e..5832a6a75ae7 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -52,7 +52,7 @@ struct bbox }; // reduce a pair of bounding boxes (a,b) to a bounding box containing a and b -struct bbox_union : public thrust::binary_function +struct bbox_union { __host__ __device__ bbox operator()(bbox a, bbox b) { From 48f456af9f018e2619dd722b792372e1fc39a75e Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 4 Sep 2025 08:54:05 +0200 Subject: [PATCH 11/19] Foxe include `host_device` for non-cuda builds --- thrust/examples/CMakeLists.txt | 5 +++++ thrust/examples/include/host_device.h | 2 ++ 2 files changed, 7 insertions(+) diff --git a/thrust/examples/CMakeLists.txt b/thrust/examples/CMakeLists.txt index aa374a3891f1..6ac23b2cb7f8 100644 --- a/thrust/examples/CMakeLists.txt +++ b/thrust/examples/CMakeLists.txt @@ -92,6 +92,11 @@ function(thrust_add_example target_name_var example_name example_src thrust_targ thrust_configure_cuda_target(${example_target} RDC ${THRUST_FORCE_RDC}) endif() + # We do not want to explicitly include `host_device.h` if not needed, so force include the file for non CUDA targets + target_compile_options(${example_target} PRIVATE + $<$:-include include/host_device.h> + ) + # Add to the active configuration's meta target add_dependencies(${config_meta_target} ${example_target}) diff --git a/thrust/examples/include/host_device.h b/thrust/examples/include/host_device.h index 40b65ace0eaf..aebb753aeba8 100644 --- a/thrust/examples/include/host_device.h +++ b/thrust/examples/include/host_device.h @@ -16,6 +16,8 @@ #pragma once +#include + #if !_CCCL_HAS_CUDA_COMPILER() # ifndef __host__ From a24acc2e51c126e5f313f5bb861443675f9d44b3 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 4 Sep 2025 08:54:55 +0200 Subject: [PATCH 12/19] Minor fixes --- thrust/examples/expand.cu | 1 - thrust/examples/sort.cu | 1 + 2 files changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index c171cb6244a7..326e820fce86 100644 --- a/thrust/examples/expand.cu +++ b/thrust/examples/expand.cu @@ -65,7 +65,6 @@ int main() thrust::device_vector d_counts = {3, 5, 2, 0, 1, 3, 4, 2, 4}; thrust::device_vector d_values = {1, 2, 3, 4, 5, 6, 7, 8, 9}; - size_t input_size = d_counts.size(); size_t output_size = thrust::reduce(d_counts.begin(), d_counts.end()); thrust::device_vector d_output(output_size); diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index b8c8f740d48b..e6265d17d0ed 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -1,6 +1,7 @@ #include #include #include +#include #include #include From bd59431cee0a98753f84be5c3220a4661c73a328 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Thu, 4 Sep 2025 15:11:27 +0200 Subject: [PATCH 13/19] Use separate flag for MSVC --- thrust/examples/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/thrust/examples/CMakeLists.txt b/thrust/examples/CMakeLists.txt index 6ac23b2cb7f8..2b8f92679233 100644 --- a/thrust/examples/CMakeLists.txt +++ b/thrust/examples/CMakeLists.txt @@ -94,6 +94,7 @@ function(thrust_add_example target_name_var example_name example_src thrust_targ # We do not want to explicitly include `host_device.h` if not needed, so force include the file for non CUDA targets target_compile_options(${example_target} PRIVATE + $<$:/FI include/host_device.h> $<$:-include include/host_device.h> ) From 5e3292967cf06104be7ef669d3378c26b1bb3be6 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Thu, 11 Sep 2025 23:01:34 -0500 Subject: [PATCH 14/19] fix race conditions --- thrust/examples/bounding_box.cu | 8 +++++-- thrust/examples/discrete_voronoi.cu | 4 ++-- thrust/examples/dot_products_with_zip.cu | 7 +++--- thrust/examples/expand.cu | 14 ++++++++++-- thrust/examples/histogram.cu | 7 +++--- thrust/examples/lexicographical_sort.cu | 7 +++--- thrust/examples/max_abs_diff.cu | 12 ++++++++-- thrust/examples/minmax.cu | 7 +++--- thrust/examples/mode.cu | 7 +++--- thrust/examples/padded_grid_reduction.cu | 8 +++---- thrust/examples/remove_points2d.cu | 11 +++++----- thrust/examples/simple_moving_average.cu | 7 +++--- thrust/examples/sort.cu | 28 ++++++++++++++---------- thrust/examples/sorting_aos_vs_soa.cu | 14 +++++++----- thrust/examples/sum_rows.cu | 7 +++--- thrust/examples/word_count.cu | 3 +-- 16 files changed, 92 insertions(+), 59 deletions(-) diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index 5832a6a75ae7..f26d0e9c6fa1 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -28,7 +28,7 @@ struct bbox // construct an empty box __host__ __device__ bbox() {} - // construct a box from a single pointgit branch -a + // construct a box from a single point __host__ __device__ bbox(const point2d& point) : lower_left(point) , upper_right(point) @@ -91,7 +91,11 @@ int main() thrust::device_vector points(N); // generate some random points in the unit square - thrust::generate(points.begin(), points.end(), random_point_generator()); + random_point_generator gen; + for (size_t i = 0; i < points.size(); ++i) + { + points[i] = gen(); + } // initial bounding box contains first point bbox init(points[0], points[0]); diff --git a/thrust/examples/discrete_voronoi.cu b/thrust/examples/discrete_voronoi.cu index fa0c78e6f633..23756da90858 100644 --- a/thrust/examples/discrete_voronoi.cu +++ b/thrust/examples/discrete_voronoi.cu @@ -124,9 +124,9 @@ void print(int m, int n, const thrust::device_vector& d_data) { thrust::host_vector h_data = d_data; - for (size_t i = 0; i < static_cast(m); i++) + for (int i = 0; i < m; i++) { - for (size_t j = 0; j < static_cast(n); j++) + for (int j = 0; j < n; j++) { std::cout << std::setw(4) << h_data[i * n + j] << " "; } diff --git a/thrust/examples/dot_products_with_zip.cu b/thrust/examples/dot_products_with_zip.cu index b5fbed13f5d8..710160e8e1a6 100644 --- a/thrust/examples/dot_products_with_zip.cu +++ b/thrust/examples/dot_products_with_zip.cu @@ -31,9 +31,10 @@ thrust::host_vector random_vector(const size_t N, unsigned int seed = thr thrust::default_random_engine rng(seed); thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector temp(N); - thrust::generate(temp.begin(), temp.end(), [&]() { - return u01(rng); - }); + for (size_t i = 0; i < N; ++i) + { + temp[i] = u01(rng); + } return temp; } diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index 326e820fce86..7eead73ce87d 100644 --- a/thrust/examples/expand.cu +++ b/thrust/examples/expand.cu @@ -8,6 +8,16 @@ #include #include +// Custom maximum functor for compatibility +template +struct max_functor +{ + __host__ __device__ T operator()(const T& a, const T& b) + { + return (a > b) ? a : b; + } +}; + // This example demonstrates how to expand an input sequence by // replicating each element a variable number of times. For example, // @@ -40,7 +50,7 @@ OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator // compute max-scan over the output indices, filling in the holes thrust::inclusive_scan( - output_indices.begin(), output_indices.end(), output_indices.begin(), cuda::maximum()); + output_indices.begin(), output_indices.end(), output_indices.begin(), max_functor()); // gather input values according to index array (output = first2[output_indices]) thrust::gather(output_indices.begin(), output_indices.end(), first2, output); @@ -65,7 +75,7 @@ int main() thrust::device_vector d_counts = {3, 5, 2, 0, 1, 3, 4, 2, 4}; thrust::device_vector d_values = {1, 2, 3, 4, 5, 6, 7, 8, 9}; - size_t output_size = thrust::reduce(d_counts.begin(), d_counts.end()); + const size_t output_size = thrust::reduce(d_counts.begin(), d_counts.end()); thrust::device_vector d_output(output_size); // expand values according to counts diff --git a/thrust/examples/histogram.cu b/thrust/examples/histogram.cu index 2e76b2791fb9..1bc7455d0440 100644 --- a/thrust/examples/histogram.cu +++ b/thrust/examples/histogram.cu @@ -149,14 +149,15 @@ int main() // generate random data on the host thrust::host_vector input(N); - thrust::generate(input.begin(), input.end(), [&]() { + for (int i = 0; i < N; i++) + { int sum = 0; for (int j = 0; j < S; j++) { sum += dist(rng); } - return sum / S; - }); + input[i] = sum / S; + } // demonstrate dense histogram method { diff --git a/thrust/examples/lexicographical_sort.cu b/thrust/examples/lexicographical_sort.cu index 9432fd05a615..6ff55cb3708b 100644 --- a/thrust/examples/lexicographical_sort.cu +++ b/thrust/examples/lexicographical_sort.cu @@ -41,9 +41,10 @@ thrust::host_vector random_vector(size_t N) static thrust::default_random_engine rng; static thrust::uniform_int_distribution dist(0, 9); - thrust::generate(vec.begin(), vec.end(), [&]() { - return dist(rng); - }); + for (size_t i = 0; i < N; i++) + { + vec[i] = dist(rng); + } return vec; } diff --git a/thrust/examples/max_abs_diff.cu b/thrust/examples/max_abs_diff.cu index ee98c2881430..40762fc01ee7 100644 --- a/thrust/examples/max_abs_diff.cu +++ b/thrust/examples/max_abs_diff.cu @@ -2,7 +2,6 @@ #include #include -#include #include // this example computes the maximum absolute difference @@ -17,6 +16,15 @@ struct abs_diff } }; +template +struct max_functor +{ + __host__ __device__ T operator()(const T& a, const T& b) + { + return (a > b) ? a : b; + } +}; + int main() { thrust::device_vector d_a = {1.0, 2.0, 3.0, 4.0}; @@ -26,7 +34,7 @@ int main() float init = 0; // binary operations - cuda::maximum binary_op1; + max_functor binary_op1; abs_diff binary_op2; float max_abs_diff = thrust::inner_product(d_a.begin(), d_a.end(), d_b.begin(), init, binary_op1, binary_op2); diff --git a/thrust/examples/minmax.cu b/thrust/examples/minmax.cu index 911cfe9b679c..9022805a44e3 100644 --- a/thrust/examples/minmax.cu +++ b/thrust/examples/minmax.cu @@ -58,9 +58,10 @@ int main() // initialize data on host thrust::host_vector host_data(N); - thrust::generate(host_data.begin(), host_data.end(), [&]() { - return dist(rng); - }); + for (size_t i = 0; i < host_data.size(); i++) + { + host_data[i] = dist(rng); + } thrust::device_vector data = host_data; // setup arguments diff --git a/thrust/examples/mode.cu b/thrust/examples/mode.cu index bc390dff0af5..0d57106a7dab 100644 --- a/thrust/examples/mode.cu +++ b/thrust/examples/mode.cu @@ -26,9 +26,10 @@ int main() // generate random data on the host thrust::host_vector h_data(N); - thrust::generate(h_data.begin(), h_data.end(), [&]() { - return dist(rng); - }); + for (size_t i = 0; i < h_data.size(); i++) + { + h_data[i] = dist(rng); + } // transfer data to device thrust::device_vector d_data(h_data); diff --git a/thrust/examples/padded_grid_reduction.cu b/thrust/examples/padded_grid_reduction.cu index aa7a91455452..8fe446643627 100644 --- a/thrust/examples/padded_grid_reduction.cu +++ b/thrust/examples/padded_grid_reduction.cu @@ -80,9 +80,9 @@ int main() thrust::device_vector data(M * N, -1); // initialize valid values in grid - for (size_t i = 0; i < static_cast(M); i++) + for (int i = 0; i < M; i++) { - for (size_t j = 0; j < static_cast(n); j++) + for (int j = 0; j < n; j++) { data[i * N + j] = dist(rng); } @@ -91,10 +91,10 @@ int main() // print full grid std::cout << "padded grid" << std::endl; std::cout << std::fixed << std::setprecision(4); - for (size_t i = 0; i < static_cast(M); i++) + for (int i = 0; i < M; i++) { std::cout << " "; - for (size_t j = 0; j < static_cast(N); j++) + for (int j = 0; j < N; j++) { std::cout << data[i * N + j] << " "; } diff --git a/thrust/examples/remove_points2d.cu b/thrust/examples/remove_points2d.cu index 7be88c51ec33..7121e2e54f59 100644 --- a/thrust/examples/remove_points2d.cu +++ b/thrust/examples/remove_points2d.cu @@ -39,12 +39,11 @@ int main() thrust::uniform_real_distribution u01(0.0f, 1.0f); thrust::host_vector x(N); thrust::host_vector y(N); - thrust::generate(x.begin(), x.end(), [&]() { - return u01(rng); - }); - thrust::generate(y.begin(), y.end(), [&]() { - return u01(rng); - }); + for (size_t i = 0; i < N; i++) + { + x[i] = u01(rng); + y[i] = u01(rng); + } // print the initial points std::cout << std::fixed; diff --git a/thrust/examples/simple_moving_average.cu b/thrust/examples/simple_moving_average.cu index 0e47d15651f2..e52d8a0ae985 100644 --- a/thrust/examples/simple_moving_average.cu +++ b/thrust/examples/simple_moving_average.cu @@ -70,9 +70,10 @@ int main() thrust::host_vector host_data(n); thrust::default_random_engine rng; thrust::uniform_int_distribution dist(0, 10); - thrust::generate(host_data.begin(), host_data.end(), [&]() { - return static_cast(dist(rng)); - }); + for (size_t i = 0; i < host_data.size(); i++) + { + host_data[i] = static_cast(dist(rng)); + } thrust::device_vector data = host_data; // allocate storage for averages diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index e6265d17d0ed..15f98a70c3b5 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -15,9 +15,10 @@ void initialize(thrust::device_vector& v) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); thrust::host_vector host_data(v.size()); - thrust::generate(host_data.begin(), host_data.end(), [&]() { - return dist(rng); - }); + for (size_t i = 0; i < host_data.size(); i++) + { + host_data[i] = dist(rng); + } v = host_data; } @@ -26,9 +27,10 @@ void initialize(thrust::device_vector& v) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(2, 19); thrust::host_vector host_data(v.size()); - thrust::generate(host_data.begin(), host_data.end(), [&]() { - return dist(rng) / 2.0f; - }); + for (size_t i = 0; i < host_data.size(); i++) + { + host_data[i] = dist(rng) / 2.0f; + } v = host_data; } @@ -37,11 +39,12 @@ void initialize(thrust::device_vector>& v) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(0, 9); thrust::host_vector> host_data(v.size()); - thrust::generate(host_data.begin(), host_data.end(), [&]() { + for (size_t i = 0; i < host_data.size(); i++) + { int a = dist(rng); int b = dist(rng); - return thrust::make_pair(a, b); - }); + host_data[i] = thrust::make_pair(a, b); + } v = host_data; } @@ -50,9 +53,10 @@ void initialize(thrust::device_vector& v1, thrust::device_vector& v2) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); thrust::host_vector host_data(v1.size()); - thrust::generate(host_data.begin(), host_data.end(), [&]() { - return dist(rng); - }); + for (size_t i = 0; i < host_data.size(); i++) + { + host_data[i] = dist(rng); + } v1 = host_data; thrust::sequence(v2.begin(), v2.end(), 0); } diff --git a/thrust/examples/sorting_aos_vs_soa.cu b/thrust/examples/sorting_aos_vs_soa.cu index 91d189a9a004..7e28d71a6d55 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -34,9 +34,10 @@ void initialize_keys(thrust::device_vector& keys) thrust::host_vector h_keys(keys.size()); - thrust::generate(h_keys.begin(), h_keys.end(), [&]() { - return dist(rng); - }); + for (size_t i = 0; i < h_keys.size(); i++) + { + h_keys[i] = dist(rng); + } keys = h_keys; } @@ -48,11 +49,12 @@ void initialize_keys(thrust::device_vector& structures) thrust::host_vector h_structures(structures.size()); - thrust::generate(h_structures.begin(), h_structures.end(), [&]() { + for (size_t i = 0; i < h_structures.size(); i++) + { MyStruct s; s.key = dist(rng); - return s; - }); + h_structures[i] = s; + } structures = h_structures; } diff --git a/thrust/examples/sum_rows.cu b/thrust/examples/sum_rows.cu index fb9fd1860051..1fa7fd2189eb 100644 --- a/thrust/examples/sum_rows.cu +++ b/thrust/examples/sum_rows.cu @@ -32,9 +32,10 @@ int main() // initialize data thrust::host_vector host_array(R * C); - thrust::generate(host_array.begin(), host_array.end(), [&]() { - return dist(rng); - }); + for (size_t i = 0; i < host_array.size(); i++) + { + host_array[i] = dist(rng); + } thrust::device_vector array = host_array; // allocate storage for row sums and indices diff --git a/thrust/examples/word_count.cu b/thrust/examples/word_count.cu index c4cc7e15f952..a41842bcab5f 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -69,8 +69,7 @@ int main() std::cout << raw_input << std::endl; // transfer to device - thrust::host_vector h_input(raw_input, raw_input + sizeof(raw_input)); - thrust::device_vector input(h_input); + thrust::device_vector input(raw_input, raw_input + sizeof(raw_input)); // count words int wc = word_count(input); From 0841ade22922242b59a89442fe76470a84dbe96b Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Thu, 11 Sep 2025 23:06:24 -0500 Subject: [PATCH 15/19] fix clang format --- thrust/examples/sort.cu | 4 ++-- thrust/examples/sorting_aos_vs_soa.cu | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index 15f98a70c3b5..9beb0eb91112 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -41,8 +41,8 @@ void initialize(thrust::device_vector>& v) thrust::host_vector> host_data(v.size()); for (size_t i = 0; i < host_data.size(); i++) { - int a = dist(rng); - int b = dist(rng); + int a = dist(rng); + int b = dist(rng); host_data[i] = thrust::make_pair(a, b); } v = host_data; diff --git a/thrust/examples/sorting_aos_vs_soa.cu b/thrust/examples/sorting_aos_vs_soa.cu index 7e28d71a6d55..7fdac79044c0 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -52,7 +52,7 @@ void initialize_keys(thrust::device_vector& structures) for (size_t i = 0; i < h_structures.size(); i++) { MyStruct s; - s.key = dist(rng); + s.key = dist(rng); h_structures[i] = s; } From 02b49b4b903f8c1b203d0e902cc59c65c619b00d Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Mon, 15 Sep 2025 16:01:34 -0500 Subject: [PATCH 16/19] use cuda::maximum --- thrust/examples/expand.cu | 12 +----------- thrust/examples/max_abs_diff.cu | 13 +++---------- 2 files changed, 4 insertions(+), 21 deletions(-) diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index 7eead73ce87d..133f922595f5 100644 --- a/thrust/examples/expand.cu +++ b/thrust/examples/expand.cu @@ -8,16 +8,6 @@ #include #include -// Custom maximum functor for compatibility -template -struct max_functor -{ - __host__ __device__ T operator()(const T& a, const T& b) - { - return (a > b) ? a : b; - } -}; - // This example demonstrates how to expand an input sequence by // replicating each element a variable number of times. For example, // @@ -50,7 +40,7 @@ OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator // compute max-scan over the output indices, filling in the holes thrust::inclusive_scan( - output_indices.begin(), output_indices.end(), output_indices.begin(), max_functor()); + output_indices.begin(), output_indices.end(), output_indices.begin(), cuda::maximum()); // gather input values according to index array (output = first2[output_indices]) thrust::gather(output_indices.begin(), output_indices.end(), first2, output); diff --git a/thrust/examples/max_abs_diff.cu b/thrust/examples/max_abs_diff.cu index 40762fc01ee7..24cd034f8097 100644 --- a/thrust/examples/max_abs_diff.cu +++ b/thrust/examples/max_abs_diff.cu @@ -2,6 +2,8 @@ #include #include +#include + #include // this example computes the maximum absolute difference @@ -16,15 +18,6 @@ struct abs_diff } }; -template -struct max_functor -{ - __host__ __device__ T operator()(const T& a, const T& b) - { - return (a > b) ? a : b; - } -}; - int main() { thrust::device_vector d_a = {1.0, 2.0, 3.0, 4.0}; @@ -34,7 +27,7 @@ int main() float init = 0; // binary operations - max_functor binary_op1; + cuda::maximum binary_op1{}; abs_diff binary_op2; float max_abs_diff = thrust::inner_product(d_a.begin(), d_a.end(), d_b.begin(), init, binary_op1, binary_op2); From 30a6d4f573204e0531bf0d1a23a2399dc9730c17 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 15 Sep 2025 23:20:53 +0200 Subject: [PATCH 17/19] Update thrust/examples/expand.cu --- thrust/examples/expand.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index 133f922595f5..acee30a51ddd 100644 --- a/thrust/examples/expand.cu +++ b/thrust/examples/expand.cu @@ -40,7 +40,7 @@ OutputIterator expand(InputIterator1 first1, InputIterator1 last1, InputIterator // compute max-scan over the output indices, filling in the holes thrust::inclusive_scan( - output_indices.begin(), output_indices.end(), output_indices.begin(), cuda::maximum()); + output_indices.begin(), output_indices.end(), output_indices.begin(), cuda::maximum{}); // gather input values according to index array (output = first2[output_indices]) thrust::gather(output_indices.begin(), output_indices.end(), first2, output); From 3557c0d49d199e5ebd9eb5f27c0b4a35390fba7c Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Mon, 15 Sep 2025 17:18:17 -0500 Subject: [PATCH 18/19] minor fixes and using ranges --- thrust/examples/minmax.cu | 8 +++--- thrust/examples/mode.cu | 4 +-- thrust/examples/scan_by_key.cu | 6 ++--- thrust/examples/scan_matrix_by_rows.cu | 6 ++--- thrust/examples/set_operations.cu | 6 ++--- thrust/examples/simple_moving_average.cu | 4 +-- thrust/examples/sort.cu | 32 +++++++++++++----------- thrust/examples/sorting_aos_vs_soa.cu | 8 +++--- thrust/examples/sum_rows.cu | 8 +++--- 9 files changed, 41 insertions(+), 41 deletions(-) diff --git a/thrust/examples/minmax.cu b/thrust/examples/minmax.cu index 9022805a44e3..006fe866d3f5 100644 --- a/thrust/examples/minmax.cu +++ b/thrust/examples/minmax.cu @@ -58,9 +58,9 @@ int main() // initialize data on host thrust::host_vector host_data(N); - for (size_t i = 0; i < host_data.size(); i++) + for (auto& e : host_data) { - host_data[i] = dist(rng); + e = dist(rng); } thrust::device_vector data = host_data; @@ -76,9 +76,9 @@ int main() // print results std::cout << "[ "; - for (size_t i = 0; i < data.size(); i++) + for (auto& e : host_data) { - std::cout << data[i] << " "; + std::cout << e << " "; } std::cout << "]" << std::endl; diff --git a/thrust/examples/mode.cu b/thrust/examples/mode.cu index 0d57106a7dab..05e023bb8ef7 100644 --- a/thrust/examples/mode.cu +++ b/thrust/examples/mode.cu @@ -26,9 +26,9 @@ int main() // generate random data on the host thrust::host_vector h_data(N); - for (size_t i = 0; i < h_data.size(); i++) + for (auto& e : h_data) { - h_data[i] = dist(rng); + e = dist(rng); } // transfer data to device diff --git a/thrust/examples/scan_by_key.cu b/thrust/examples/scan_by_key.cu index 9a8128db2e7e..daba51044007 100644 --- a/thrust/examples/scan_by_key.cu +++ b/thrust/examples/scan_by_key.cu @@ -18,11 +18,11 @@ struct head_flag_predicate template void print(const Vector& v) { - for (size_t i = 0; i < v.size(); i++) + for (const auto& e : v) { - std::cout << v[i] << " "; + std::cout << e << " "; } - std::cout << "\n"; + std::cout << std::endl; } int main() diff --git a/thrust/examples/scan_matrix_by_rows.cu b/thrust/examples/scan_matrix_by_rows.cu index eaef3714604e..01047814f104 100644 --- a/thrust/examples/scan_matrix_by_rows.cu +++ b/thrust/examples/scan_matrix_by_rows.cu @@ -14,7 +14,7 @@ __host__ void scan_matrix_by_rows0(thrust::device_vector& u, int n, int m) // Here, we launch a separate scan for each row in the matrix. This works, // but each kernel only does a small amount of work. It would be better if we // could launch one big kernel for the entire matrix. - for (size_t i = 0; i < static_cast(n); ++i) + for (int i = 0; i < n; ++i) { thrust::inclusive_scan(u.begin() + m * i, u.begin() + m * (i + 1), u.begin() + m * i); } @@ -69,9 +69,9 @@ int main() thrust::sequence(u1.begin(), u1.end()); scan_matrix_by_rows1(u1, n, m); - for (size_t i = 0; i < static_cast(n); ++i) + for (int i = 0; i < n; ++i) { - for (size_t j = 0; j < static_cast(m); ++j) + for (int j = 0; j < m; ++j) { assert(u0[j + m * i] == u1[j + m * i]); } diff --git a/thrust/examples/set_operations.cu b/thrust/examples/set_operations.cu index 8a99b7ca14cc..94ab201ffe3d 100644 --- a/thrust/examples/set_operations.cu +++ b/thrust/examples/set_operations.cu @@ -37,11 +37,11 @@ template void print(const String& s, const Vector& v) { std::cout << s << " ["; - for (size_t i = 0; i < v.size(); i++) + for (const auto& e : v) { - std::cout << " " << v[i]; + std::cout << " " << e; } - std::cout << " ]\n"; + std::cout << " ]" << std::endl; } template diff --git a/thrust/examples/simple_moving_average.cu b/thrust/examples/simple_moving_average.cu index e52d8a0ae985..2d51a00153ca 100644 --- a/thrust/examples/simple_moving_average.cu +++ b/thrust/examples/simple_moving_average.cu @@ -70,9 +70,9 @@ int main() thrust::host_vector host_data(n); thrust::default_random_engine rng; thrust::uniform_int_distribution dist(0, 10); - for (size_t i = 0; i < host_data.size(); i++) + for (auto& e : host_data) { - host_data[i] = static_cast(dist(rng)); + e = static_cast(dist(rng)); } thrust::device_vector data = host_data; diff --git a/thrust/examples/sort.cu b/thrust/examples/sort.cu index 9beb0eb91112..d02e161ce731 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -4,6 +4,8 @@ #include #include +#include + #include #include #include @@ -15,9 +17,9 @@ void initialize(thrust::device_vector& v) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); thrust::host_vector host_data(v.size()); - for (size_t i = 0; i < host_data.size(); i++) + for (auto& e : host_data) { - host_data[i] = dist(rng); + e = dist(rng); } v = host_data; } @@ -27,23 +29,23 @@ void initialize(thrust::device_vector& v) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(2, 19); thrust::host_vector host_data(v.size()); - for (size_t i = 0; i < host_data.size(); i++) + for (auto& e : host_data) { - host_data[i] = dist(rng) / 2.0f; + e = dist(rng) / 2.0f; } v = host_data; } -void initialize(thrust::device_vector>& v) +void initialize(thrust::device_vector>& v) { thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(0, 9); - thrust::host_vector> host_data(v.size()); - for (size_t i = 0; i < host_data.size(); i++) + thrust::host_vector> host_data(v.size()); + for (auto& e : host_data) { - int a = dist(rng); - int b = dist(rng); - host_data[i] = thrust::make_pair(a, b); + int a = dist(rng); + int b = dist(rng); + e = cuda::std::make_pair(a, b); } v = host_data; } @@ -53,9 +55,9 @@ void initialize(thrust::device_vector& v1, thrust::device_vector& v2) thrust::default_random_engine rng(123456); thrust::uniform_int_distribution dist(10, 99); thrust::host_vector host_data(v1.size()); - for (size_t i = 0; i < host_data.size(); i++) + for (auto& e : host_data) { - host_data[i] = dist(rng); + e = dist(rng); } v1 = host_data; thrust::sequence(v2.begin(), v2.end(), 0); @@ -79,11 +81,11 @@ void print(const thrust::device_vector& v) std::cout << "\n"; } -void print(const thrust::device_vector>& v) +void print(const thrust::device_vector>& v) { for (const auto& p : v) { - thrust::pair local_p = p; + cuda::std::pair local_p = p; std::cout << " (" << local_p.first << "," << local_p.second << ")"; } std::cout << "\n"; @@ -161,7 +163,7 @@ int main() std::cout << "\nsorting pairs\n"; { - thrust::device_vector> keys(N); + thrust::device_vector> keys(N); initialize(keys); print(keys); thrust::sort(keys.begin(), keys.end()); diff --git a/thrust/examples/sorting_aos_vs_soa.cu b/thrust/examples/sorting_aos_vs_soa.cu index 7fdac79044c0..65bb84588c19 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -34,9 +34,9 @@ void initialize_keys(thrust::device_vector& keys) thrust::host_vector h_keys(keys.size()); - for (size_t i = 0; i < h_keys.size(); i++) + for (auto& e : h_keys) { - h_keys[i] = dist(rng); + e = dist(rng); } keys = h_keys; @@ -51,9 +51,7 @@ void initialize_keys(thrust::device_vector& structures) for (size_t i = 0; i < h_structures.size(); i++) { - MyStruct s; - s.key = dist(rng); - h_structures[i] = s; + h_structures[i].key = dist(rng); } structures = h_structures; diff --git a/thrust/examples/sum_rows.cu b/thrust/examples/sum_rows.cu index 1fa7fd2189eb..36d40bc3193d 100644 --- a/thrust/examples/sum_rows.cu +++ b/thrust/examples/sum_rows.cu @@ -32,9 +32,9 @@ int main() // initialize data thrust::host_vector host_array(R * C); - for (size_t i = 0; i < host_array.size(); i++) + for (auto& e : host_array) { - host_array[i] = dist(rng); + e = dist(rng); } thrust::device_vector array = host_array; @@ -53,10 +53,10 @@ int main() ::cuda::std::plus()); // print data - for (size_t i = 0; i < static_cast(R); i++) + for (int i = 0; i < R; i++) { std::cout << "[ "; - for (size_t j = 0; j < static_cast(C); j++) + for (int j = 0; j < C; j++) { std::cout << array[i * C + j] << " "; } From d171ad84ffe6e82392a8cc9b9931879971a18cc3 Mon Sep 17 00:00:00 2001 From: Sai Charan Date: Tue, 16 Sep 2025 00:14:19 -0500 Subject: [PATCH 19/19] reverting back --- thrust/examples/bounding_box.cu | 26 ++++++-------------------- thrust/examples/discrete_voronoi.cu | 2 +- 2 files changed, 7 insertions(+), 21 deletions(-) diff --git a/thrust/examples/bounding_box.cu b/thrust/examples/bounding_box.cu index f26d0e9c6fa1..60506f1eacc3 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -66,23 +66,6 @@ struct bbox_union } }; -// functor to generate random points -struct random_point_generator -{ - mutable thrust::default_random_engine rng; - mutable thrust::uniform_real_distribution u01; - - random_point_generator() - : rng(0) - , u01(0.0f, 1.0f) - {} - - __host__ __device__ point2d operator()() const - { - return point2d(u01(rng), u01(rng)); - } -}; - int main() { const size_t N = 40; @@ -91,10 +74,13 @@ int main() thrust::device_vector points(N); // generate some random points in the unit square - random_point_generator gen; - for (size_t i = 0; i < points.size(); ++i) + thrust::default_random_engine rng; + thrust::uniform_real_distribution u01(0.0f, 1.0f); + for (size_t i = 0; i < N; i++) { - points[i] = gen(); + float x = u01(rng); + float y = u01(rng); + points[i] = point2d(x, y); } // initial bounding box contains first point diff --git a/thrust/examples/discrete_voronoi.cu b/thrust/examples/discrete_voronoi.cu index 23756da90858..c23624a7ef02 100644 --- a/thrust/examples/discrete_voronoi.cu +++ b/thrust/examples/discrete_voronoi.cu @@ -139,7 +139,7 @@ void generate_random_sites(thrust::host_vector& t, int Nb, int m, int n) thrust::default_random_engine rng; thrust::uniform_int_distribution dist(0, m * n - 1); - for (size_t k = 0; k < static_cast(Nb); k++) + for (int k = 0; k < Nb; k++) { int index = dist(rng); t[index] = index + 1;