Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Update to Thrust 1.17.0 #11437

Merged
merged 11 commits into from
Aug 11, 2022
1 change: 1 addition & 0 deletions cpp/benchmarks/common/generate_input.cu
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@
#include <rmm/device_uvector.hpp>

#include <thrust/binary_search.h>
#include <thrust/copy.h>
#include <thrust/device_ptr.h>
#include <thrust/execution_policy.h>
#include <thrust/fill.h>
Expand Down
2 changes: 2 additions & 0 deletions cpp/benchmarks/io/orc/orc_writer_chunks.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,8 @@
#include <cudf/io/orc.hpp>
#include <cudf/table/table.hpp>

#include <thrust/iterator/transform_iterator.h>

// to enable, run cmake with -DBUILD_BENCHMARKS=ON

constexpr int64_t data_size = 512 << 20;
Expand Down
1 change: 1 addition & 0 deletions cpp/benchmarks/iterator/iterator.cu
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/pair.h>
#include <thrust/reduce.h>

#include <random>

Expand Down
1 change: 1 addition & 0 deletions cpp/benchmarks/join/join_common.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,7 @@

#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/random/linear_congruential_engine.h>
#include <thrust/random/uniform_int_distribution.h>

Expand Down
2 changes: 1 addition & 1 deletion cpp/cmake/thirdparty/get_thrust.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -80,6 +80,6 @@ function(find_and_configure_thrust VERSION)
endif()
endfunction()

set(CUDF_MIN_VERSION_Thrust 1.15.0)
set(CUDF_MIN_VERSION_Thrust 1.17.0)

find_and_configure_thrust(${CUDF_MIN_VERSION_Thrust})
112 changes: 58 additions & 54 deletions cpp/cmake/thrust.patch
Original file line number Diff line number Diff line change
@@ -1,83 +1,87 @@
diff --git a/thrust/system/cuda/detail/sort.h b/thrust/system/cuda/detail/sort.h
index 1ffeef0..5e80800 100644
--- a/thrust/system/cuda/detail/sort.h
+++ b/thrust/system/cuda/detail/sort.h
@@ -108,7 +108,7 @@ namespace __merge_sort {
key_type key2 = keys_shared[keys2_beg];

diff --git a/cub/block/block_merge_sort.cuh b/cub/block/block_merge_sort.cuh
index 4769df36..d86d6342 100644
--- a/cub/block/block_merge_sort.cuh
+++ b/cub/block/block_merge_sort.cuh
@@ -91,7 +91,7 @@ __device__ __forceinline__ void SerialMerge(KeyT *keys_shared,
KeyT key1 = keys_shared[keys1_beg];
KeyT key2 = keys_shared[keys2_beg];

-#pragma unroll
+#pragma unroll 1
for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
{
bool p = (keys2_beg < keys2_end) &&
@@ -311,10 +311,10 @@ namespace __merge_sort {
void stable_odd_even_sort(key_type (&keys)[ITEMS_PER_THREAD],
item_type (&items)[ITEMS_PER_THREAD])
for (int item = 0; item < ITEMS_PER_THREAD; ++item)
{
bool p = (keys2_beg < keys2_end) &&
@@ -383,7 +383,7 @@ public:
//
KeyT max_key = oob_default;

- #pragma unroll
+ #pragma unroll 1
for (int item = 1; item < ITEMS_PER_THREAD; ++item)
{
-#pragma unroll
+#pragma unroll 1
for (int i = 0; i < ITEMS_PER_THREAD; ++i)
{
-#pragma unroll
+#pragma unroll 1
for (int j = 1 & i; j < ITEMS_PER_THREAD - 1; j += 2)
{
if (compare_op(keys[j + 1], keys[j]))
@@ -350,7 +350,7 @@ namespace __merge_sort {
// each thread has sorted keys_loc
// merge sort keys_loc in shared memory
//
-#pragma unroll
+#pragma unroll 1
for (int coop = 2; coop <= BLOCK_THREADS; coop *= 2)
{
sync_threadblock();
@@ -479,7 +479,7 @@ namespace __merge_sort {
// and fill the remainig keys with it
//
key_type max_key = keys_loc[0];
-#pragma unroll
+#pragma unroll 1
for (int ITEM = 1; ITEM < ITEMS_PER_THREAD; ++ITEM)
{
if (ITEMS_PER_THREAD * tid + ITEM < num_remaining)
diff a/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/device/dispatch/dispatch_radix_sort.cuh
index 41eb1d2..f2893b4 100644
if (ITEMS_PER_THREAD * linear_tid + item < valid_items)
@@ -407,7 +407,7 @@ public:
// each thread has sorted keys
// merge sort keys in shared memory
//
- #pragma unroll
+ #pragma unroll 1
for (int target_merged_threads_number = 2;
target_merged_threads_number <= NUM_THREADS;
target_merged_threads_number *= 2)
diff --git a/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/device/dispatch/dispatch_radix_sort.cuh
index b188c75f..3f36656f 100644
--- a/cub/device/dispatch/dispatch_radix_sort.cuh
+++ b/cub/device/dispatch/dispatch_radix_sort.cuh
@@ -723,7 +723,7 @@ struct DeviceRadixSortPolicy
@@ -736,7 +736,7 @@ struct DeviceRadixSortPolicy


/// SM60 (GP100)
- struct Policy600 : ChainedPolicy<600, Policy600, Policy500>
+ struct Policy600 : ChainedPolicy<600, Policy600, Policy600>
{
enum {
PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5, // 6.9B 32b keys/s (Quadro P100)
diff a/cub/device/dispatch/dispatch_reduce.cuh b/cub/device/dispatch/dispatch_reduce.cuh
index f6aee45..dd64301 100644
diff --git a/cub/device/dispatch/dispatch_reduce.cuh b/cub/device/dispatch/dispatch_reduce.cuh
index e0470ccb..6a0c2ed6 100644
--- a/cub/device/dispatch/dispatch_reduce.cuh
+++ b/cub/device/dispatch/dispatch_reduce.cuh
@@ -284,7 +284,7 @@ struct DeviceReducePolicy
@@ -280,7 +280,7 @@ struct DeviceReducePolicy
};

/// SM60
- struct Policy600 : ChainedPolicy<600, Policy600, Policy350>
+ struct Policy600 : ChainedPolicy<600, Policy600, Policy600>
{
// ReducePolicy (P100: 591 GB/s @ 64M 4B items; 583 GB/s @ 256M 1B items)
typedef AgentReducePolicy<
diff a/cub/device/dispatch/dispatch_scan.cuh b/cub/device/dispatch/dispatch_scan.cuh
index c0c6d59..937ee31 100644
diff --git a/cub/device/dispatch/dispatch_scan.cuh b/cub/device/dispatch/dispatch_scan.cuh
index c2d04588..ac2d10e0 100644
--- a/cub/device/dispatch/dispatch_scan.cuh
+++ b/cub/device/dispatch/dispatch_scan.cuh
@@ -178,7 +178,7 @@ struct DeviceScanPolicy
@@ -177,7 +177,7 @@ struct DeviceScanPolicy
};

/// SM600
- struct Policy600 : ChainedPolicy<600, Policy600, Policy520>
+ struct Policy600 : ChainedPolicy<600, Policy600, Policy600>
{
typedef AgentScanPolicy<
128, 15, ///< Threads per block, items per thread
diff --git a/cub/thread/thread_sort.cuh b/cub/thread/thread_sort.cuh
index 5d486789..b42fb5f0 100644
--- a/cub/thread/thread_sort.cuh
+++ b/cub/thread/thread_sort.cuh
@@ -83,10 +83,10 @@ StableOddEvenSort(KeyT (&keys)[ITEMS_PER_THREAD],
{
constexpr bool KEYS_ONLY = std::is_same<ValueT, NullType>::value;

- #pragma unroll
+ #pragma unroll 1
for (int i = 0; i < ITEMS_PER_THREAD; ++i)
{
- #pragma unroll
+ #pragma unroll 1
for (int j = 1 & i; j < ITEMS_PER_THREAD - 1; j += 2)
{
if (compare_op(keys[j + 1], keys[j]))
2 changes: 2 additions & 0 deletions cpp/include/cudf/detail/copy.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,8 @@
#include <cudf/detail/copy.hpp>
#include <cudf/detail/gather.cuh>

#include <thrust/iterator/counting_iterator.h>

namespace cudf::detail {

/**
Expand Down
1 change: 1 addition & 0 deletions cpp/include/cudf/detail/indexalator.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@
#include <thrust/execution_policy.h>
#include <thrust/functional.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/optional.h>
#include <thrust/pair.h>

Expand Down
2 changes: 2 additions & 0 deletions cpp/include/cudf/detail/labeling/label_segments.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -23,8 +23,10 @@
#include <thrust/distance.h>
#include <thrust/for_each.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/reduce.h>
#include <thrust/scan.h>
#include <thrust/scatter.h>
#include <thrust/uninitialized_fill.h>

namespace cudf::detail {
Expand Down
1 change: 1 addition & 0 deletions cpp/include/cudf/detail/utilities/hash_functions.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@

#include <thrust/distance.h>
#include <thrust/execution_policy.h>
#include <thrust/find.h>
#include <thrust/iterator/reverse_iterator.h>
#include <thrust/pair.h>
#include <thrust/reverse.h>
Expand Down
1 change: 1 addition & 0 deletions cpp/include/cudf/lists/list_device_view.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@
#include <cudf/utilities/type_dispatcher.hpp>

#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/pair.h>

namespace cudf {
Expand Down
1 change: 1 addition & 0 deletions cpp/include/cudf/strings/detail/gather.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -26,6 +26,7 @@
#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>

#include <thrust/advance.h>
#include <thrust/binary_search.h>
#include <thrust/distance.h>
#include <thrust/execution_policy.h>
Expand Down
2 changes: 2 additions & 0 deletions cpp/include/cudf/strings/detail/scatter.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,8 @@
#include <rmm/device_uvector.hpp>
#include <rmm/exec_policy.hpp>

#include <thrust/distance.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/scatter.h>

namespace cudf {
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/pair.h>
#include <thrust/transform.h>
#include <thrust/transform_reduce.h>
Expand Down
6 changes: 6 additions & 0 deletions cpp/include/cudf/table/experimental/row_operators.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -32,9 +32,15 @@
#include <cudf/utilities/traits.hpp>
#include <cudf/utilities/type_dispatcher.hpp>

#include <thrust/detail/use_default.h>
#include <thrust/equal.h>
#include <thrust/execution_policy.h>
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/iterator_adaptor.h>
#include <thrust/iterator/iterator_categories.h>
#include <thrust/iterator/iterator_facade.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/logical.h>
#include <thrust/swap.h>
#include <thrust/transform_reduce.h>
Expand Down
2 changes: 2 additions & 0 deletions cpp/src/binaryop/compiled/struct_binary_ops.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,8 @@
#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>

#include <thrust/tabulate.h>

namespace cudf::binops::compiled::detail {
template <class T, class... Ts>
inline constexpr bool is_any_v = std::disjunction<std::is_same<T, Ts>...>::value;
Expand Down
1 change: 1 addition & 0 deletions cpp/src/copying/purge_nonempty_nulls.cu
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@
#include <cudf/detail/copy.cuh>
#include <cudf/utilities/default_stream.hpp>

#include <thrust/count.h>
#include <thrust/iterator/counting_iterator.h>

namespace cudf {
Expand Down
1 change: 1 addition & 0 deletions cpp/src/groupby/sort/group_collect.cu
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@
#include <rmm/cuda_stream_view.hpp>

#include <thrust/copy.h>
#include <thrust/count.h>
#include <thrust/execution_policy.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/transform.h>
Expand Down
2 changes: 2 additions & 0 deletions cpp/src/groupby/sort/group_correlation.cu
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,8 @@

#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/reduce.h>
#include <thrust/transform.h>
#include <thrust/tuple.h>
Expand Down
1 change: 1 addition & 0 deletions cpp/src/groupby/sort/group_quantiles.cu
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,7 @@
#include <thrust/execution_policy.h>
#include <thrust/for_each.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/transform.h>

namespace cudf {
namespace groupby {
Expand Down
1 change: 1 addition & 0 deletions cpp/src/groupby/sort/group_rank_scan.cu
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@

#include <thrust/functional.h>
#include <thrust/iterator/reverse_iterator.h>
#include <thrust/pair.h>
#include <thrust/scan.h>
#include <thrust/tabulate.h>
#include <thrust/transform.h>
Expand Down
1 change: 1 addition & 0 deletions cpp/src/hash/spark_murmur_hash.cu
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@
#include <rmm/cuda_stream_view.hpp>
#include <rmm/exec_policy.hpp>

#include <thrust/iterator/counting_iterator.h>
#include <thrust/tabulate.h>

namespace cudf {
Expand Down
4 changes: 4 additions & 0 deletions cpp/src/io/comp/nvcomp_adapter.cu
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,10 @@

#include <rmm/exec_policy.hpp>

#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform.h>
#include <thrust/tuple.h>

namespace cudf::io::nvcomp {

batched_args create_batched_nvcomp_args(device_span<device_span<uint8_t const> const> inputs,
Expand Down
1 change: 1 addition & 0 deletions cpp/src/io/csv/datetime.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@
#include <cudf/fixed_point/fixed_point.hpp>

#include <thrust/execution_policy.h>
#include <thrust/find.h>
#include <thrust/reduce.h>

namespace cudf {
Expand Down
1 change: 1 addition & 0 deletions cpp/src/io/json/json_gpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@
#include <rmm/device_uvector.hpp>
#include <rmm/exec_policy.hpp>

#include <thrust/advance.h>
#include <thrust/detail/copy.h>
#include <thrust/execution_policy.h>
#include <thrust/find.h>
Expand Down
4 changes: 4 additions & 0 deletions cpp/src/io/orc/writer_impl.cu
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,11 @@
#include <thrust/host_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/reverse_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/optional.h>
#include <thrust/pair.h>
#include <thrust/reduce.h>
#include <thrust/scan.h>
#include <thrust/tabulate.h>
#include <thrust/transform.h>

Expand Down
1 change: 1 addition & 0 deletions cpp/src/io/parquet/page_enc.cu
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,7 @@
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/reverse_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/merge.h>
#include <thrust/scan.h>
#include <thrust/scatter.h>
Expand Down
1 change: 1 addition & 0 deletions cpp/src/io/parquet/reader_impl.cu
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,7 @@
#include <rmm/device_uvector.hpp>
#include <rmm/exec_policy.hpp>

#include <thrust/fill.h>
#include <thrust/for_each.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/logical.h>
Expand Down
1 change: 1 addition & 0 deletions cpp/src/io/parquet/writer_impl.cu
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,7 @@

#include <thrust/binary_search.h>
#include <thrust/for_each.h>
#include <thrust/host_vector.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform.h>
#include <thrust/tuple.h>
Expand Down
Loading