Skip to content

Commit

Permalink
Revert "patch thrust to fix intmax num elements limitation in scan_by…
Browse files Browse the repository at this point in the history
…_key" (#8263)

Reverts #8199

According to @allisonvacanti (NVIDIA/thrust#1424 (comment)) this patch will likely have adverse effect on performance. We should revert it until a better solution can be found.

Authors:
  - Christopher Harris (https://github.com/cwharris)

Approvers:
  - David Wendt (https://github.com/davidwendt)
  - Keith Kraus (https://github.com/kkraus14)
  - Elias Stehle (https://github.com/elstehle)

URL: #8263
  • Loading branch information
cwharris authored May 24, 2021
1 parent 936b02d commit 259d69b
Showing 1 changed file with 0 additions and 22 deletions.
22 changes: 0 additions & 22 deletions cpp/cmake/thrust.patch
Original file line number Diff line number Diff line change
Expand Up @@ -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<Size>(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<int>(thrust::distance(keys_first, keys_last));
+ size_t num_items = static_cast<size_t>(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;

0 comments on commit 259d69b

Please sign in to comment.