Replace CUB_RUNTIME_FUNCTION with THRUST_RUNTIME_FUNCTION.
Replace CUDA_VERSION (only available if <cuda.h> is included) with CUDART_VERSION.
MODIFY
dependencies/cub
dependencies/cub
-Subproject commit 464a90bb6d7ffec28a02ccfbc93c6f5c99e8fd6f
+Subproject commit 04d36e691fed3a765909e266e88fae563e07ffa9

MODIFY
testing/vector.cu
testing/vector.cu
@@ -616,7 +616,7 @@ void TestVectorResizing(void)
// reset the CUDA error
cudaGetLastError();
} // end catch
-#endif // defined(__CUDACC__) && CUDA_VERSION==3000
+#endif // defined(__CUDACC__) && CUDART_VERSION==3000
ASSERT_EQUAL(v.size(), 0lu);
}
@@ -647,7 +647,7 @@ void TestVectorReserving(void)
}
catch(std::length_error e) {}
catch(std::bad_alloc e) {}
-#endif // defined(__CUDACC__) && CUDA_VERSION==3000
+#endif // defined(__CUDACC__) && CUDART_VERSION==3000
ASSERT_EQUAL(v.capacity(), old_capacity);
}

MODIFY
thrust/detail/complex/c99math.h
thrust/detail/complex/c99math.h
@@ -124,7 +124,7 @@ __host__ __device__ inline int isfinite(double x){
// sometimes the CUDA toolkit provides these these names as macros,
// sometimes functions in the global scope
-# if (CUDA_VERSION >= 6500)
+# if (CUDART_VERSION >= 6500)
using ::isinf;
using ::isnan;
using ::signbit;
@@ -133,7 +133,7 @@ using ::isfinite;
# else
// these names are macros, we don't need to define them
-# endif // CUDA_VERSION
+# endif // CUDART_VERSION
# else
@@ -234,3 +234,4 @@ inline double hypot(double x, double y){
} // namespace detail
} // namespace thrust
+

MODIFY
thrust/detail/config/global_workarounds.h
thrust/detail/config/global_workarounds.h
// XXX workaround gcc 4.8+'s complaints about unused local typedefs by silencing them globally
#if defined(THRUST_GCC_VERSION) && (THRUST_GCC_VERSION >= 40800)
-# if defined(__NVCC__) && (CUDA_VERSION >= 6000)
+# if defined(__NVCC__) && (CUDART_VERSION >= 6000)
# pragma GCC diagnostic ignored "-Wunused-local-typedefs"
# endif // nvcc & cuda 6+
#endif // gcc 4.8

MODIFY
thrust/detail/malloc_and_free.h
thrust/detail/malloc_and_free.h
@@ -54,7 +54,7 @@ pointer<T,DerivedPolicy> malloc(const thrust::detail::execution_policy_base<Deri
// XXX WAR nvbug 992955
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#if CUDA_VERSION < 5000
+#if CUDART_VERSION < 5000
// cudafe generates unqualified calls to free(int *volatile)
// which get confused with thrust::free
@@ -65,7 +65,7 @@ void free(int *volatile ptr)
::free(ptr);
}
-#endif // CUDA_VERSION
+#endif // CUDART_VERSION
#endif // THRUST_DEVICE_COMPILER
__thrust_exec_check_disable__

MODIFY
thrust/system/cuda/detail/core/agent_launcher.h
thrust/system/cuda/detail/core/agent_launcher.h
@@ -397,7 +397,7 @@ namespace core {
MAX_SHMEM_PER_BLOCK> shm1;
template <class Size>
- CUB_RUNTIME_FUNCTION
+ THRUST_RUNTIME_FUNCTION
AgentLauncher(AgentPlan plan_,
Size count_,
cudaStream_t stream_,
@@ -417,7 +417,7 @@ namespace core {
}
template <class Size>
- CUB_RUNTIME_FUNCTION
+ THRUST_RUNTIME_FUNCTION
AgentLauncher(AgentPlan plan_,
Size count_,
cudaStream_t stream_,
@@ -437,7 +437,7 @@ namespace core {
assert(count > 0);
}
- CUB_RUNTIME_FUNCTION
+ THRUST_RUNTIME_FUNCTION
AgentLauncher(AgentPlan plan_,
cudaStream_t stream_,
char const* name_,
@@ -455,7 +455,7 @@ namespace core {
assert(plan.grid_size > 0);
}
- CUB_RUNTIME_FUNCTION
+ THRUST_RUNTIME_FUNCTION
AgentLauncher(AgentPlan plan_,
cudaStream_t stream_,
char* vshmem,
@@ -500,7 +500,7 @@ namespace core {
}
#endif
- CUB_RUNTIME_FUNCTION
+ THRUST_RUNTIME_FUNCTION
typename core::get_plan<Agent>::type static get_plan(cudaStream_t , void* d_ptr = 0)
{
THRUST_UNUSED_VAR(d_ptr);
@@ -514,7 +514,7 @@ namespace core {
return get_agent_plan<Agent>(lowest_supported_sm_arch::ver);
}
- CUB_RUNTIME_FUNCTION void sync() const
+ THRUST_RUNTIME_FUNCTION void sync() const
{
if (debug_sync)
{
@@ -709,7 +709,7 @@ namespace core {
// don't compile other kernel which accepts pointer
// and save on compilations
template <class... Args>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, Args... args) const
{
assert(has_shmem && vshmem == NULL);
@@ -727,7 +727,7 @@ namespace core {
// do actually have enough shared memory, the compilation time will double.
//
template <class... Args>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, Args... args) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -737,7 +737,7 @@ namespace core {
}
template <class... Args>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(Args... args) const
{
#if __THRUST__TEMPLATE_DEBUG
@@ -755,7 +755,7 @@ namespace core {
}
#else
template <class _0>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -765,7 +765,7 @@ namespace core {
.doit(ptr, vshmem, x0);
}
template <class _0, class _1>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -775,7 +775,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1);
}
template <class _0, class _1, class _2>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -785,7 +785,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2);
}
template <class _0, class _1, class _2, class _3>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -795,7 +795,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3);
}
template <class _0, class _1, class _2, class _3, class _4>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -805,7 +805,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4);
}
template <class _0, class _1, class _2, class _3, class _4, class _5>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -815,7 +815,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -825,7 +825,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -835,7 +835,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6, x7);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -844,7 +844,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6, x7, x8);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -854,7 +854,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9,_xA xA) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -864,7 +864,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9,_xA xA,_xB xB) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -874,7 +874,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9,_xA xA,_xB xB,_xC xC) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -884,7 +884,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC, class _xD>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9,_xA xA,_xB xB,_xC xC,_xD xD) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -894,7 +894,7 @@ namespace core {
.doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC, xD);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC, class _xD, class _xE>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9,_xA xA,_xB xB,_xC xC,_xD xD,_xE xE) const
{
assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0));
@@ -909,7 +909,7 @@ namespace core {
////////////////////////////////////////////////////////
template <class _0>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0) const
{
assert(has_shmem && vshmem == NULL);
@@ -919,7 +919,7 @@ namespace core {
.doit(ptr, x0);
}
template <class _0, class _1>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1) const
{
assert(has_shmem && vshmem == NULL);
@@ -929,7 +929,7 @@ namespace core {
.doit(ptr, x0, x1);
}
template <class _0, class _1, class _2>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2) const
{
assert(has_shmem && vshmem == NULL);
@@ -939,7 +939,7 @@ namespace core {
.doit(ptr, x0, x1, x2);
}
template <class _0, class _1, class _2, class _3>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3) const
{
assert(has_shmem && vshmem == NULL);
@@ -949,7 +949,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3);
}
template <class _0, class _1, class _2, class _3, class _4>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4) const
{
assert(has_shmem && vshmem == NULL);
@@ -959,7 +959,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4);
}
template <class _0, class _1, class _2, class _3, class _4, class _5>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5) const
{
assert(has_shmem && vshmem == NULL);
@@ -969,7 +969,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6) const
{
assert(has_shmem && vshmem == NULL);
@@ -979,7 +979,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7) const
{
assert(has_shmem && vshmem == NULL);
@@ -989,7 +989,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6, x7);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8) const
{
assert(has_shmem && vshmem == NULL);
@@ -999,7 +999,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6, x7, x8);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9) const
{
assert(has_shmem && vshmem == NULL);
@@ -1009,7 +1009,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA) const
{
assert(has_shmem && vshmem == NULL);
@@ -1019,7 +1019,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB) const
{
assert(has_shmem && vshmem == NULL);
@@ -1029,7 +1029,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB, _xC xC) const
{
assert(has_shmem && vshmem == NULL);
@@ -1039,7 +1039,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC, class _xD>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB, _xC xC, _xD xD) const
{
assert(has_shmem && vshmem == NULL);
@@ -1049,7 +1049,7 @@ namespace core {
.doit(ptr, x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC, xD);
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC, class _xD, class _xE>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB, _xC xC, _xD xD, _xE xE) const
{
assert(has_shmem && vshmem == NULL);
@@ -1064,105 +1064,105 @@ namespace core {
////////////////////////////////////////////////////////
template <class _0>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0) const
{
launch_impl(has_enough_shmem_t(), x0);
sync();
}
template <class _0, class _1>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1) const
{
launch_impl(has_enough_shmem_t(), x0, x1);
sync();
}
template <class _0, class _1, class _2>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2);
sync();
}
template <class _0, class _1, class _2, class _3>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3);
sync();
}
template <class _0, class _1, class _2, class _3, class _4>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7, x8);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7, x8, x9);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB, _xC xC) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC, class _xD>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB, _xC xC, _xD xD) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC, xD);
sync();
}
template <class _0, class _1, class _2, class _3, class _4, class _5, class _6, class _7, class _8, class _9, class _xA, class _xB, class _xC, class _xD, class _xE>
- void CUB_RUNTIME_FUNCTION
+ void THRUST_RUNTIME_FUNCTION
launch(_0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB, _xC xC, _xD xD, _xE xE) const
{
launch_impl(has_enough_shmem_t(), x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC, xD, xE);

MODIFY
thrust/system/cuda/detail/core/triple_chevron_launch.h
thrust/system/cuda/detail/core/triple_chevron_launch.h
@@ -45,7 +45,7 @@ namespace launcher {
Size const shared_mem;
cudaStream_t const stream;
- CUB_RUNTIME_FUNCTION
+ THRUST_RUNTIME_FUNCTION
triple_chevron(dim3 grid_,
dim3 block_,
Size shared_mem_ = 0,

MODIFY
thrust/system/cuda/detail/core/util.h
thrust/system/cuda/detail/core/util.h
@@ -451,8 +451,8 @@ namespace core {
return i32value;
}
- size_t CUB_RUNTIME_FUNCTION
- inline get_max_shared_memory_per_block()
+ size_t THRUST_RUNTIME_FUNCTION
+ get_max_shared_memory_per_block()
{
int dev_id;
cuda_cub::throw_on_error(cudaGetDevice(&dev_id),
@@ -471,8 +471,8 @@ namespace core {
return static_cast<size_t>(i32value);
}
- size_t CUB_RUNTIME_FUNCTION
- inline virtual_shmem_size(size_t shmem_per_block)
+ size_t THRUST_RUNTIME_FUNCTION
+ virtual_shmem_size(size_t shmem_per_block)
{
size_t max_shmem_per_block = core::get_max_shared_memory_per_block();
if (shmem_per_block > max_shmem_per_block)
@@ -481,8 +481,8 @@ namespace core {
return 0;
}
- size_t CUB_RUNTIME_FUNCTION
- inline vshmem_size(size_t shmem_per_block, size_t num_blocks)
+ size_t THRUST_RUNTIME_FUNCTION
+ vshmem_size(size_t shmem_per_block, size_t num_blocks)
{
size_t max_shmem_per_block = core::get_max_shared_memory_per_block();
if (shmem_per_block > max_shmem_per_block)
@@ -603,7 +603,7 @@ namespace core {
__host__ __device__ operator T const &() const { return value_; }
};
- inline cuda_optional<int> CUB_RUNTIME_FUNCTION
+ cuda_optional<int> THRUST_RUNTIME_FUNCTION
get_ptx_version()
{
int ptx_version = 0;
@@ -611,7 +611,7 @@ namespace core {
return cuda_optional<int>(ptx_version, status);
}
- inline cudaError_t CUB_RUNTIME_FUNCTION
+ cudaError_t THRUST_RUNTIME_FUNCTION
sync_stream(cudaStream_t stream)
{
return cub::SyncStream(stream);

MODIFY
thrust/system/cuda/detail/extrema.h
thrust/system/cuda/detail/extrema.h
@@ -376,7 +376,7 @@ namespace __extrema {
class Derived,
class ItemsIt,
class BinaryPred>
- ItemsIt CUB_RUNTIME_FUNCTION
+ ItemsIt THRUST_RUNTIME_FUNCTION
element(execution_policy<Derived> &policy,
ItemsIt first,
ItemsIt last,

MODIFY
thrust/system/cuda/detail/internal/copy_device_to_device.h
thrust/system/cuda/detail/internal/copy_device_to_device.h
@@ -42,7 +42,7 @@ namespace __copy {
template <class Derived,
class InputIt,
class OutputIt>
- OutputIt CUB_RUNTIME_FUNCTION
+ OutputIt THRUST_RUNTIME_FUNCTION
device_to_device(execution_policy<Derived>& policy,
InputIt first,
InputIt last,

MODIFY
thrust/system/cuda/detail/merge.h
thrust/system/cuda/detail/merge.h
@@ -673,7 +673,7 @@ namespace __merge {
class KeysOutputIt,
class ItemsOutputIt,
class CompareOp>
- cudaError_t CUB_RUNTIME_FUNCTION
+ cudaError_t THRUST_RUNTIME_FUNCTION
doit_step(void* d_temp_storage,
size_t& temp_storage_bytes,
KeysIt1 keys1,

Replace CUB_RUNTIME_FUNCTION with THRUST_RUNTIME_FUNCTION.
Replace CUDA_VERSION (only available if <cuda.h> is included) with CUDART_VERSION.