diff --git a/thrust/examples/CMakeLists.txt b/thrust/examples/CMakeLists.txt index aa374a3891f1..2b8f92679233 100644 --- a/thrust/examples/CMakeLists.txt +++ b/thrust/examples/CMakeLists.txt @@ -92,6 +92,12 @@ 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 + $<$:/FI include/host_device.h> + $<$:-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/arbitrary_transformation.cu b/thrust/examples/arbitrary_transformation.cu index 2389955cb84c..e16be0f731d6 100644 --- a/thrust/examples/arbitrary_transformation.cu +++ b/thrust/examples/arbitrary_transformation.cu @@ -7,8 +7,6 @@ #include -#include "include/host_device.h" - // 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. @@ -62,21 +60,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 + 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 +73,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 +86,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..f2cc84925f5b 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 holds 4 integers + 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..60506f1eacc3 100644 --- a/thrust/examples/bounding_box.cu +++ b/thrust/examples/bounding_box.cu @@ -4,8 +4,6 @@ #include #include -#include "include/host_device.h" - // This example shows how to compute a bounding box // for a set of points in two dimensions. @@ -54,7 +52,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 { __host__ __device__ bbox operator()(bbox a, bbox b) { @@ -71,13 +69,13 @@ struct bbox_reduction 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 + thrust::default_random_engine rng; + thrust::uniform_real_distribution u01(0.0f, 1.0f); for (size_t i = 0; i < N; i++) { float x = u01(rng); @@ -86,13 +84,10 @@ int main() } // initial bounding box contains first point - bbox init = bbox(points[0], points[0]); - - // binary reduction operation - bbox_reduction binary_op; + bbox init(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(), init, 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..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" - // 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..48c4ac90810c 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 " << 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 be013cba810d..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 " << (last - first) << " 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/discrete_voronoi.cu b/thrust/examples/discrete_voronoi.cu index e9887f6a9cfe..c23624a7ef02 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) @@ -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..710160e8e1a6 100644 --- a/thrust/examples/dot_products_with_zip.cu +++ b/thrust/examples/dot_products_with_zip.cu @@ -5,8 +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 // vector type (Float3) whose (x,y,z) components will be stored in @@ -33,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); - for (size_t i = 0; i < N; i++) + for (size_t i = 0; i < N; ++i) { temp[i] = u01(rng); } diff --git a/thrust/examples/expand.cu b/thrust/examples/expand.cu index 39cd750aa769..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); @@ -62,15 +62,10 @@ 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); + 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/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__ diff --git a/thrust/examples/lambda.cu b/thrust/examples/lambda.cu index 783f7244ad95..2b8fde4cf5de 100644 --- a/thrust/examples/lambda.cu +++ b/thrust/examples/lambda.cu @@ -4,8 +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 +39,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 +56,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 +75,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..6ff55cb3708b 100644 --- a/thrust/examples/lexicographical_sort.cu +++ b/thrust/examples/lexicographical_sort.cu @@ -59,7 +59,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 +82,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..24cd034f8097 100644 --- a/thrust/examples/max_abs_diff.cu +++ b/thrust/examples/max_abs_diff.cu @@ -2,10 +2,9 @@ #include #include -#include -#include +#include -#include "include/host_device.h" +#include // this example computes the maximum absolute difference // between the elements of two vectors @@ -21,21 +20,14 @@ 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; // binary operations - ::cuda::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); diff --git a/thrust/examples/minmax.cu b/thrust/examples/minmax.cu index 772257cf7af8..006fe866d3f5 100644 --- a/thrust/examples/minmax.cu +++ b/thrust/examples/minmax.cu @@ -5,8 +5,6 @@ #include #include -#include "include/host_device.h" - // compute minimum and maximum values in a single reduction // minmax_pair stores the minimum and maximum @@ -59,11 +57,12 @@ int main() thrust::uniform_int_distribution dist(10, 99); // initialize data on host - thrust::device_vector data(N); - for (size_t i = 0; i < data.size(); i++) + thrust::host_vector host_data(N); + for (auto& e : host_data) { - data[i] = dist(rng); + e = dist(rng); } + thrust::device_vector data = host_data; // setup arguments minmax_unary_op unary_op; @@ -77,9 +76,9 @@ int main() // print results std::cout << "[ "; - for (size_t i = 0; i < N; 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 6849c2c5fc70..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 < N; i++) + for (auto& e : h_data) { - h_data[i] = dist(rng); + e = dist(rng); } // transfer data to device @@ -70,7 +70,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[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/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..8fe446643627 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. 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..7121e2e54f59 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 @@ -50,7 +48,7 @@ int main() // 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 +67,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_by_key.cu b/thrust/examples/scan_by_key.cu index d275362da9eb..daba51044007 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 @@ -20,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 d19991553fdc..01047814f104 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. 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 c1ed4fc17171..2d51a00153ca 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 @@ -8,8 +9,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. // @@ -68,13 +67,14 @@ 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); - for (size_t i = 0; i < n; i++) + for (auto& e : host_data) { - data[i] = static_cast(dist(rng)); + e = static_cast(dist(rng)); } + thrust::device_vector data = host_data; // allocate storage for averages thrust::device_vector averages(data.size() - (w - 1)); @@ -84,9 +84,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..d02e161ce731 100644 --- a/thrust/examples/sort.cu +++ b/thrust/examples/sort.cu @@ -1,11 +1,14 @@ #include +#include #include +#include #include +#include + #include #include - -#include "include/host_device.h" +#include // Helper routines @@ -13,69 +16,77 @@ 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++) + thrust::host_vector host_data(v.size()); + for (auto& e : host_data) { - v[i] = dist(rng); + e = dist(rng); } + v = host_data; } 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++) + thrust::host_vector host_data(v.size()); + for (auto& e : host_data) { - v[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); - for (size_t i = 0; i < v.size(); i++) + thrust::host_vector> host_data(v.size()); + for (auto& e : host_data) { int a = dist(rng); int b = dist(rng); - v[i] = thrust::make_pair(a, b); + e = cuda::std::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); - for (size_t i = 0; i < v1.size(); i++) + thrust::host_vector host_data(v1.size()); + for (auto& e : host_data) { - v1[i] = dist(rng); - v2[i] = static_cast(i); + e = dist(rng); } + v1 = host_data; + thrust::sequence(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) +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 << ")"; + cuda::std::pair local_p = p; + std::cout << " (" << local_p.first << "," << local_p.second << ")"; } std::cout << "\n"; } @@ -152,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 4532d0c657e9..65bb84588c19 100644 --- a/thrust/examples/sorting_aos_vs_soa.cu +++ b/thrust/examples/sorting_aos_vs_soa.cu @@ -5,7 +5,6 @@ #include -#include "include/host_device.h" #include "include/timer.h" // This examples compares sorting performance using Array of Structures (AoS) @@ -35,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; 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/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 718355a22dfe..36d40bc3193d 100644 --- a/thrust/examples/sum_rows.cu +++ b/thrust/examples/sum_rows.cu @@ -7,8 +7,6 @@ #include -#include "include/host_device.h" - // convert a linear index to a row index template struct linear_index_to_row_index @@ -33,11 +31,12 @@ int main() thrust::uniform_int_distribution dist(10, 99); // initialize data - thrust::device_vector array(R * C); - for (size_t i = 0; i < array.size(); i++) + thrust::host_vector host_array(R * C); + for (auto& e : host_array) { - array[i] = dist(rng); + e = 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/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..3f6af6aa5558 100644 --- a/thrust/examples/transform_output_iterator.cu +++ b/thrust/examples/transform_output_iterator.cu @@ -1,11 +1,11 @@ #include #include #include +#include #include #include -#include "include/host_device.h" struct Functor { template @@ -19,15 +19,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 +36,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..08fc2fd32076 100644 --- a/thrust/examples/weld_vertices.cu +++ b/thrust/examples/weld_vertices.cu @@ -1,5 +1,6 @@ #include #include +#include #include #include #include @@ -41,15 +42,18 @@ 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..a41842bcab5f 100644 --- a/thrust/examples/word_count.cu +++ b/thrust/examples/word_count.cu @@ -1,12 +1,11 @@ #include #include +#include #include #include #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.