From ae08422ac7b7b9f693579370652cf6de0891ba71 Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Tue, 11 May 2021 18:00:54 -0500 Subject: [PATCH] patch thrust to fix intmax num elements limitation in scan_by_key (#8199) same fix seen here, but via patch: https://github.com/NVIDIA/thrust/pull/1424 Also fixes https://github.com/rapidsai/cuspatial/issues/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: https://github.com/rapidsai/cudf/pull/8199 --- cpp/cmake/thrust.patch | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/cpp/cmake/thrust.patch b/cpp/cmake/thrust.patch index 3f876f7ffb7..3cedff8b80d 100644 --- a/cpp/cmake/thrust.patch +++ b/cpp/cmake/thrust.patch @@ -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(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;