Skip to content

Commit

Permalink
patch thrust to fix intmax num elements limitation in scan_by_key (#8199
Browse files Browse the repository at this point in the history
)

same fix seen here, but via patch: NVIDIA/thrust#1424

Also fixes rapidsai/cuspatial#393

Alternatively, we could wait and update our thrust version, rather than patching the existing one.

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

Approvers:
  - Mark Harris (https://github.com/harrism)
  - Paul Taylor (https://github.com/trxcllnt)

URL: #8199
  • Loading branch information
cwharris authored May 11, 2021
1 parent 2c70f1d commit ae08422
Showing 1 changed file with 22 additions and 0 deletions.
22 changes: 22 additions & 0 deletions cpp/cmake/thrust.patch
Original file line number Diff line number Diff line change
Expand Up @@ -42,3 +42,25 @@ 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<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 ae08422

Please sign in to comment.