From c42502eb85dd1654e55c2ede6a24f6109be431a3 Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Wed, 12 May 2021 13:17:17 -0500 Subject: [PATCH 1/2] Revert "patch thrust to fix intmax num elements limitation in scan_by_key (#8199)" This reverts commit ae08422ac7b7b9f693579370652cf6de0891ba71. --- cpp/cmake/thrust.patch | 22 ---------------------- 1 file changed, 22 deletions(-) diff --git a/cpp/cmake/thrust.patch b/cpp/cmake/thrust.patch index 3cedff8b80d..3f876f7ffb7 100644 --- a/cpp/cmake/thrust.patch +++ b/cpp/cmake/thrust.patch @@ -42,25 +42,3 @@ index 1ffeef0..5e80800 100644 for (int ITEM = 1; ITEM < ITEMS_PER_THREAD; ++ITEM) { if (ITEMS_PER_THREAD * tid + ITEM < num_remaining) -diff --git a/thrust/system/cuda/detail/scan_by_key.h b/thrust/system/cuda/detail/scan_by_key.h -index fe4b321c..b3974c69 100644 ---- a/thrust/system/cuda/detail/scan_by_key.h -+++ b/thrust/system/cuda/detail/scan_by_key.h -@@ -513,7 +513,7 @@ namespace __scan_by_key { - scan_op(scan_op_) - { - int tile_idx = blockIdx.x; -- Size tile_base = ITEMS_PER_TILE * tile_idx; -+ Size tile_base = ITEMS_PER_TILE * static_cast(tile_idx); - Size num_remaining = num_items - tile_base; - - if (num_remaining > ITEMS_PER_TILE) -@@ -734,7 +734,7 @@ namespace __scan_by_key { - ScanOp scan_op, - AddInitToScan add_init_to_scan) - { -- int num_items = static_cast(thrust::distance(keys_first, keys_last)); -+ size_t num_items = static_cast(thrust::distance(keys_first, keys_last)); - size_t storage_size = 0; - cudaStream_t stream = cuda_cub::stream(policy); - bool debug_sync = THRUST_DEBUG_SYNC_FLAG; From eec348ecec37b5d71db1853a73d1cc80de85ceef Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Mon, 17 May 2021 13:31:46 -0500 Subject: [PATCH 2/2] . --- cpp/cmake/thrust.patch | 22 ---------------------- 1 file changed, 22 deletions(-) diff --git a/cpp/cmake/thrust.patch b/cpp/cmake/thrust.patch index c14b8cdafe5..2f9201d8ab4 100644 --- a/cpp/cmake/thrust.patch +++ b/cpp/cmake/thrust.patch @@ -81,25 +81,3 @@ index c0c6d59..937ee31 100644 { typedef AgentScanPolicy< 128, 15, ///< Threads per block, items per thread -diff --git a/thrust/system/cuda/detail/scan_by_key.h b/thrust/system/cuda/detail/scan_by_key.h -index fe4b321c..b3974c69 100644 ---- a/thrust/system/cuda/detail/scan_by_key.h -+++ b/thrust/system/cuda/detail/scan_by_key.h -@@ -513,7 +513,7 @@ namespace __scan_by_key { - scan_op(scan_op_) - { - int tile_idx = blockIdx.x; -- Size tile_base = ITEMS_PER_TILE * tile_idx; -+ Size tile_base = ITEMS_PER_TILE * static_cast(tile_idx); - Size num_remaining = num_items - tile_base; - - if (num_remaining > ITEMS_PER_TILE) -@@ -734,7 +734,7 @@ namespace __scan_by_key { - ScanOp scan_op, - AddInitToScan add_init_to_scan) - { -- int num_items = static_cast(thrust::distance(keys_first, keys_last)); -+ size_t num_items = static_cast(thrust::distance(keys_first, keys_last)); - size_t storage_size = 0; - cudaStream_t stream = cuda_cub::stream(policy); - bool debug_sync = THRUST_DEBUG_SYNC_FLAG;