From f537e429e1a71b17b8da0fa946cedb8ce43b459d Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Thu, 6 May 2021 14:14:06 -0500 Subject: [PATCH 1/3] fix > intmax num inputs for scan_by_key --- thrust/system/cuda/detail/scan_by_key.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/thrust/system/cuda/detail/scan_by_key.h b/thrust/system/cuda/detail/scan_by_key.h index fe4b321c0..ed6d7db56 100644 --- a/thrust/system/cuda/detail/scan_by_key.h +++ b/thrust/system/cuda/detail/scan_by_key.h @@ -512,7 +512,7 @@ namespace __scan_by_key { inequality_op(equality_op_), scan_op(scan_op_) { - int tile_idx = blockIdx.x; + Size tile_idx = blockIdx.x; Size tile_base = ITEMS_PER_TILE * tile_idx; Size num_remaining = num_items - tile_base; @@ -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 33a33d3112aa404cb50a328c0a3eac6058f1dbcb Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Mon, 10 May 2021 14:26:06 -0500 Subject: [PATCH 2/3] revert scan_by_key tile_idx type to int, static_cast on multiplication instead. --- thrust/system/cuda/detail/scan_by_key.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/thrust/system/cuda/detail/scan_by_key.h b/thrust/system/cuda/detail/scan_by_key.h index ed6d7db56..17b925fe5 100644 --- a/thrust/system/cuda/detail/scan_by_key.h +++ b/thrust/system/cuda/detail/scan_by_key.h @@ -512,8 +512,8 @@ namespace __scan_by_key { inequality_op(equality_op_), scan_op(scan_op_) { - Size tile_idx = blockIdx.x; - Size tile_base = ITEMS_PER_TILE * tile_idx; + int tile_idx = blockIdx.x; + Size tile_base = ITEMS_PER_TILE * static_cast(tile_idx); Size num_remaining = num_items - tile_base; if (num_remaining > ITEMS_PER_TILE) From 949fc338ed64faedbc319104610c5d18aa45fbf9 Mon Sep 17 00:00:00 2001 From: Christopher Harris Date: Tue, 11 May 2021 10:11:26 -0500 Subject: [PATCH 3/3] use size_t where Size was unavailable --- thrust/system/cuda/detail/scan_by_key.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/system/cuda/detail/scan_by_key.h b/thrust/system/cuda/detail/scan_by_key.h index 17b925fe5..edf579695 100644 --- a/thrust/system/cuda/detail/scan_by_key.h +++ b/thrust/system/cuda/detail/scan_by_key.h @@ -734,7 +734,7 @@ namespace __scan_by_key { ScanOp scan_op, AddInitToScan add_init_to_scan) { - 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;