From ffaa54c703115011aa5f987d9b5b60d29381ffc0 Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Mon, 10 May 2021 15:19:49 -0500 Subject: [PATCH 1/2] patch thrust to fix > intmax num scan limitation --- cpp/cmake/thrust.patch | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/cpp/cmake/thrust.patch b/cpp/cmake/thrust.patch index 3f876f7ffb7..40e412ea53e 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 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 e61849831af09f667cd1eb4695d078113c9dba51 Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Tue, 11 May 2021 10:14:11 -0500 Subject: [PATCH 2/2] fix error where Type was not yet available. --- cpp/cmake/thrust.patch | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/cmake/thrust.patch b/cpp/cmake/thrust.patch index 40e412ea53e..3cedff8b80d 100644 --- a/cpp/cmake/thrust.patch +++ b/cpp/cmake/thrust.patch @@ -60,7 +60,7 @@ index fe4b321c..b3974c69 100644 AddInitToScan add_init_to_scan) { - int num_items = static_cast(thrust::distance(keys_first, keys_last)); -+ Size 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;