Modify CUDA backend to use `::cub` instead of `::thrust::cuda_cub::cub`.
MODIFY
.dependencies/cub
.dependencies/cub
-Subproject commit 2b5c0cde428f58d75915466c5b6704b6ebbb5b64
+Subproject commit 464a90bb6d7ffec28a02ccfbc93c6f5c99e8fd6f

MODIFY
thrust/system/cuda/detail/adjacent_difference.h
thrust/system/cuda/detail/adjacent_difference.h
#include <thrust/detail/cstdint.h>
#include <thrust/detail/temporary_array.h>
#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/cub/device/device_select.cuh>
-#include <thrust/system/cuda/detail/cub/block/block_adjacent_difference.cuh>
+#include <cub/device/device_select.cuh>
+#include <cub/block/block_adjacent_difference.cuh>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/par_to_seq.h>
#include <thrust/functional.h>

MODIFY
thrust/system/cuda/detail/async/reduce.h
thrust/system/cuda/detail/async/reduce.h
@@ -82,7 +82,7 @@ auto async_reduce_n(
size_t tmp_size = 0;
thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceReduce::Reduce(
+ cub::DeviceReduce::Reduce(
nullptr
, tmp_size
, first
@@ -164,7 +164,7 @@ auto async_reduce_n(
// Run reduction.
thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceReduce::Reduce(
+ cub::DeviceReduce::Reduce(
tmp_ptr
, tmp_size
, first
@@ -237,7 +237,7 @@ auto async_reduce_into_n(
size_t tmp_size = 0;
thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceReduce::Reduce(
+ cub::DeviceReduce::Reduce(
nullptr
, tmp_size
, first
@@ -301,7 +301,7 @@ auto async_reduce_into_n(
// Run reduction.
thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceReduce::Reduce(
+ cub::DeviceReduce::Reduce(
tmp_ptr
, tmp_size
, first

MODIFY
thrust/system/cuda/detail/async/sort.h
thrust/system/cuda/detail/async/sort.h
@@ -297,12 +297,12 @@ invoke_radix_sort(
cudaStream_t stream
, void* tmp_ptr
, std::size_t& tmp_size
-, thrust::cuda_cub::cub::DoubleBuffer<T>& keys
+, cub::DoubleBuffer<T>& keys
, Size& n
, StrictWeakOrdering
)
{
- return thrust::cuda_cub::cub::DeviceRadixSort::SortKeys(
+ return cub::DeviceRadixSort::SortKeys(
tmp_ptr
, tmp_size
, keys
@@ -324,12 +324,12 @@ invoke_radix_sort(
cudaStream_t stream
, void* tmp_ptr
, std::size_t& tmp_size
-, thrust::cuda_cub::cub::DoubleBuffer<T>& keys
+, cub::DoubleBuffer<T>& keys
, Size& n
, StrictWeakOrdering
)
{
- return thrust::cuda_cub::cub::DeviceRadixSort::SortKeysDescending(
+ return cub::DeviceRadixSort::SortKeysDescending(
tmp_ptr
, tmp_size
, keys
@@ -372,7 +372,7 @@ auto async_stable_sort_n(
unique_eager_event e;
- thrust::cuda_cub::cub::DoubleBuffer<T> keys(
+ cub::DoubleBuffer<T> keys(
raw_pointer_cast(&*first), nullptr
);

MODIFY
thrust/system/cuda/detail/copy_if.h
thrust/system/cuda/detail/copy_if.h
#include <thrust/detail/cstdint.h>
#include <thrust/detail/temporary_array.h>
#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/cub/device/device_select.cuh>
+#include <cub/device/device_select.cuh>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/core/util.h>
#include <thrust/system/cuda/detail/par_to_seq.h>

MODIFY
thrust/system/cuda/detail/core/util.h
thrust/system/cuda/detail/core/util.h
#include <thrust/type_traits/is_contiguous_iterator.h>
#include <thrust/detail/raw_pointer_cast.h>
#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/cub/block/block_load.cuh>
-#include <thrust/system/cuda/detail/cub/block/block_store.cuh>
-#include <thrust/system/cuda/detail/cub/block/block_scan.cuh>
+#include <cub/block/block_load.cuh>
+#include <cub/block/block_store.cuh>
+#include <cub/block/block_scan.cuh>
THRUST_BEGIN_NS
@@ -491,6 +491,51 @@ namespace core {
return 0;
}
+ template <class Kernel>
+ int CUB_RUNTIME_FUNCTION
+ get_max_block_size(Kernel k)
+ {
+ int devId;
+ cuda_cub::throw_on_error(cudaGetDevice(&devId),
+ "get_max_block_size :"
+ "failed to cudaGetDevice");
+
+ cudaOccDeviceProp occ_prop;
+ cuda_cub::throw_on_error(get_occ_device_properties(occ_prop, devId),
+ "get_max_block_size: "
+ "failed to cudaGetDeviceProperties");
+
+
+ cudaFuncAttributes attribs;
+ cuda_cub::throw_on_error(cudaFuncGetAttributes(&attribs, reinterpret_cast<void *>(k)),
+ "get_max_block_size: "
+ "failed to cudaFuncGetAttributes");
+ cudaOccFuncAttributes occ_attrib(attribs);
+
+
+ cudaFuncCache cacheConfig;
+ cuda_cub::throw_on_error(cudaDeviceGetCacheConfig(&cacheConfig),
+ "get_max_block_size: "
+ "failed to cudaDeviceGetCacheConfig");
+
+ cudaOccDeviceState occ_state;
+ occ_state.cacheConfig = (cudaOccCacheConfig)cacheConfig;
+ int block_size = 0;
+ int min_grid_size = 0;
+ cudaOccError occ_status = cudaOccMaxPotentialOccupancyBlockSize(&min_grid_size,
+ &block_size,
+ &occ_prop,
+ &occ_attrib,
+ &occ_state,
+ 0);
+ if (CUDA_OCC_SUCCESS != occ_status || block_size <= 0)
+ cuda_cub::throw_on_error(cudaErrorInvalidConfiguration,
+ "get_max_block_size: "
+ "failed to cudaOccMaxPotentialOccupancyBlockSize");
+
+ return block_size;
+ }
+
// LoadIterator
// ------------
// if trivial iterator is passed, wrap loads into LDG
@@ -623,7 +668,7 @@ namespace core {
}
#define CUDA_CUB_RET_IF_FAIL(e) \
- if (thrust::cuda_cub::cub::Debug((e), __FILE__, __LINE__)) return e;
+ if (cub::Debug((e), __FILE__, __LINE__)) return e;
// uninitialized
// -------

MODIFY
thrust/system/cuda/detail/malloc_and_free.h
thrust/system/cuda/detail/malloc_and_free.h
#include <thrust/memory.h>
#include <thrust/system/cuda/config.h>
#ifdef THRUST_CACHING_DEVICE_MALLOC
-#include <thrust/system/cuda/detail/cub/util_allocator.cuh>
+#include <cub/util_allocator.cuh>
#endif
#include <thrust/system/cuda/detail/util.h>
#include <thrust/system/detail/bad_alloc.h>

MODIFY
thrust/system/cuda/detail/partition.h
thrust/system/cuda/detail/partition.h
#include <thrust/system/cuda/detail/reverse.h>
#include <thrust/system/cuda/detail/find.h>
#include <thrust/system/cuda/detail/uninitialized_copy.h>
-#include <thrust/system/cuda/detail/cub/device/device_partition.cuh>
+#include <cub/device/device_partition.cuh>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/par_to_seq.h>
#include <thrust/partition.h>

MODIFY
thrust/system/cuda/detail/reduce.h
thrust/system/cuda/detail/reduce.h
#include <thrust/system/cuda/detail/util.h>
#include <thrust/detail/raw_reference_cast.h>
#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
-#include <thrust/system/cuda/detail/cub/device/device_reduce.cuh>
+#include <cub/device/device_reduce.cuh>
#include <thrust/system/cuda/detail/par_to_seq.h>
#include <thrust/system/cuda/detail/get_value.h>
#include <thrust/functional.h>

MODIFY
thrust/system/cuda/detail/reduce_by_key.h
thrust/system/cuda/detail/reduce_by_key.h
#include <thrust/system/cuda/detail/util.h>
#include <thrust/detail/raw_reference_cast.h>
#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
-#include <thrust/system/cuda/detail/cub/device/device_reduce.cuh>
+#include <cub/device/device_reduce.cuh>
#include <thrust/system/cuda/detail/par_to_seq.h>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/get_value.h>

MODIFY
thrust/system/cuda/detail/scan.h
thrust/system/cuda/detail/scan.h
#include <thrust/detail/cstdint.h>
#include <thrust/detail/temporary_array.h>
#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/cub/device/device_scan.cuh>
+#include <cub/device/device_scan.cuh>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/par_to_seq.h>
#include <thrust/detail/mpl/math.h>

MODIFY
thrust/system/cuda/detail/sort.h
thrust/system/cuda/detail/sort.h
#include <thrust/system/cuda/config.h>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/core/util.h>
-#include <thrust/system/cuda/detail/cub/device/device_radix_sort.cuh>
+#include <cub/device/device_radix_sort.cuh>
#include <thrust/system/cuda/detail/execution_policy.h>
#include <thrust/system/cuda/detail/par_to_seq.h>

MODIFY
thrust/system/cuda/detail/unique.h
thrust/system/cuda/detail/unique.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
#include <thrust/system/cuda/config.h>
-#include <thrust/system/cuda/detail/cub/device/device_select.cuh>
+#include <cub/device/device_select.cuh>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/par_to_seq.h>
#include <thrust/detail/cstdint.h>

MODIFY
thrust/system/cuda/detail/unique_by_key.h
thrust/system/cuda/detail/unique_by_key.h
#include <thrust/detail/cstdint.h>
#include <thrust/detail/temporary_array.h>
#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/cub/device/device_select.cuh>
+#include <cub/device/device_select.cuh>
#include <thrust/system/cuda/detail/core/agent_launcher.h>
#include <thrust/system/cuda/detail/get_value.h>
#include <thrust/system/cuda/detail/par_to_seq.h>

MODIFY
thrust/system/cuda/detail/util.h
thrust/system/cuda/detail/util.h
#include <cstdio>
#include <thrust/detail/config.h>
#include <thrust/iterator/iterator_traits.h>
-#include <thrust/system/cuda/detail/cub/util_arch.cuh>
+#include <cub/util_arch.cuh>
#include <thrust/system/cuda/detail/execution_policy.h>
#include <thrust/system_error.h>
#include <thrust/system/cuda/error.h>

Modify CUDA backend to use `::cub` instead of `::thrust::cuda_cub::cub`.