diff --git a/cmake/CubHeaderTesting.cmake b/cmake/CubHeaderTesting.cmake index 036c2fe72a..845b4c16db 100644 --- a/cmake/CubHeaderTesting.cmake +++ b/cmake/CubHeaderTesting.cmake @@ -26,6 +26,11 @@ foreach(cub_target IN LISTS CUB_TARGETS) set(headertest_target ${config_prefix}.headers) add_library(${headertest_target} OBJECT ${headertest_srcs}) target_link_libraries(${headertest_target} PUBLIC ${cub_target}) + # Wrap Thrust/CUB in a custom namespace to check proper use of ns macros: + target_compile_definitions(${headertest_target} PRIVATE + "THRUST_WRAPPED_NAMESPACE=wrapped_thrust" + "CUB_WRAPPED_NAMESPACE=wrapped_cub" + ) cub_clone_target_properties(${headertest_target} ${cub_target}) add_dependencies(cub.all.headers ${headertest_target}) diff --git a/cub/agent/agent_histogram.cuh b/cub/agent/agent_histogram.cuh index 7559bf126b..b1d2143aed 100644 --- a/cub/agent/agent_histogram.cuh +++ b/cub/agent/agent_histogram.cuh @@ -41,11 +41,7 @@ #include "../grid/grid_queue.cuh" #include "../iterator/cache_modified_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -779,9 +775,4 @@ struct AgentHistogram }; - - - -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) - +CUB_NAMESPACE_END diff --git a/cub/agent/agent_radix_sort_downsweep.cuh b/cub/agent/agent_radix_sort_downsweep.cuh index a90571d630..5621b2348e 100644 --- a/cub/agent/agent_radix_sort_downsweep.cuh +++ b/cub/agent/agent_radix_sort_downsweep.cuh @@ -46,11 +46,7 @@ #include "../util_type.cuh" #include "../iterator/cache_modified_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -777,6 +773,5 @@ struct AgentRadixSortDownsweep -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_radix_sort_histogram.cuh b/cub/agent/agent_radix_sort_histogram.cuh index c8178dd797..a16c1f6883 100644 --- a/cub/agent/agent_radix_sort_histogram.cuh +++ b/cub/agent/agent_radix_sort_histogram.cuh @@ -41,11 +41,7 @@ #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN template < int _BLOCK_THREADS, @@ -237,5 +233,4 @@ struct AgentRadixSortHistogram } }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_radix_sort_onesweep.cuh b/cub/agent/agent_radix_sort_onesweep.cuh index 641f35a708..6b5d115c58 100644 --- a/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/agent/agent_radix_sort_onesweep.cuh @@ -40,11 +40,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** \brief cub::RadixSortStoreAlgorithm enumerates different algorithms to write * partitioned elements (keys, values) stored in shared memory into global @@ -662,5 +658,4 @@ struct AgentRadixSortOnesweep } }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_radix_sort_upsweep.cuh b/cub/agent/agent_radix_sort_upsweep.cuh index 5865a60a2f..56dc218735 100644 --- a/cub/agent/agent_radix_sort_upsweep.cuh +++ b/cub/agent/agent_radix_sort_upsweep.cuh @@ -42,11 +42,7 @@ #include "../util_type.cuh" #include "../iterator/cache_modified_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** * Tuning policy types @@ -521,6 +517,5 @@ struct AgentRadixSortUpsweep }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_reduce.cuh b/cub/agent/agent_reduce.cuh index 8556835ac6..d43fe104c3 100644 --- a/cub/agent/agent_reduce.cuh +++ b/cub/agent/agent_reduce.cuh @@ -44,11 +44,7 @@ #include "../iterator/cache_modified_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -381,6 +377,5 @@ struct AgentReduce }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_reduce_by_key.cuh b/cub/agent/agent_reduce_by_key.cuh index 82714e3a8d..d0162ac1f4 100644 --- a/cub/agent/agent_reduce_by_key.cuh +++ b/cub/agent/agent_reduce_by_key.cuh @@ -44,11 +44,7 @@ #include "../iterator/cache_modified_input_iterator.cuh" #include "../iterator/constant_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -542,6 +538,5 @@ struct AgentReduceByKey }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_rle.cuh b/cub/agent/agent_rle.cuh index e21bfdb6ab..db32231d21 100644 --- a/cub/agent/agent_rle.cuh +++ b/cub/agent/agent_rle.cuh @@ -46,11 +46,7 @@ #include "../iterator/cache_modified_input_iterator.cuh" #include "../iterator/constant_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -832,6 +828,5 @@ struct AgentRle }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_scan.cuh b/cub/agent/agent_scan.cuh index 0abdb2b7d0..d5da15ae06 100644 --- a/cub/agent/agent_scan.cuh +++ b/cub/agent/agent_scan.cuh @@ -43,11 +43,7 @@ #include "../grid/grid_queue.cuh" #include "../iterator/cache_modified_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -485,6 +481,5 @@ struct AgentScan }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_segment_fixup.cuh b/cub/agent/agent_segment_fixup.cuh index 5c0251def4..c6eecf2d6f 100644 --- a/cub/agent/agent_segment_fixup.cuh +++ b/cub/agent/agent_segment_fixup.cuh @@ -44,11 +44,7 @@ #include "../iterator/cache_modified_input_iterator.cuh" #include "../iterator/constant_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -370,6 +366,5 @@ struct AgentSegmentFixup }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_select_if.cuh b/cub/agent/agent_select_if.cuh index 20d68fc0a8..3fc7061e0e 100644 --- a/cub/agent/agent_select_if.cuh +++ b/cub/agent/agent_select_if.cuh @@ -45,11 +45,7 @@ #include "../grid/grid_queue.cuh" #include "../iterator/cache_modified_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -698,6 +694,5 @@ struct AgentSelectIf -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/agent_spmv_orig.cuh b/cub/agent/agent_spmv_orig.cuh index e377797a27..d3ef413677 100644 --- a/cub/agent/agent_spmv_orig.cuh +++ b/cub/agent/agent_spmv_orig.cuh @@ -46,11 +46,7 @@ #include "../iterator/counting_input_iterator.cuh" #include "../iterator/tex_obj_input_iterator.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -665,6 +661,5 @@ struct AgentSpmv -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/agent/single_pass_scan_operators.cuh b/cub/agent/single_pass_scan_operators.cuh index 924ef2a7ac..03000b3981 100644 --- a/cub/agent/single_pass_scan_operators.cuh +++ b/cub/agent/single_pass_scan_operators.cuh @@ -41,11 +41,7 @@ #include "../config.cuh" #include "../util_device.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -809,6 +805,5 @@ struct TilePrefixCallbackOp }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_adjacent_difference.cuh b/cub/block/block_adjacent_difference.cuh index c8953756db..3e32d32275 100644 --- a/cub/block/block_adjacent_difference.cuh +++ b/cub/block/block_adjacent_difference.cuh @@ -37,11 +37,7 @@ #include "../util_type.cuh" #include "../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN template < typename T, @@ -592,5 +588,4 @@ public: }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_discontinuity.cuh b/cub/block/block_discontinuity.cuh index 8c72a2405a..23ba00e552 100644 --- a/cub/block/block_discontinuity.cuh +++ b/cub/block/block_discontinuity.cuh @@ -37,11 +37,7 @@ #include "../util_type.cuh" #include "../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief The BlockDiscontinuity class provides [collective](index.html#sec0) methods for flagging discontinuities within an ordered set of items partitioned across a CUDA thread block. ![](discont_logo.png) @@ -1144,5 +1140,4 @@ public: }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_exchange.cuh b/cub/block/block_exchange.cuh index fbe64afc19..56c90c07f8 100644 --- a/cub/block/block_exchange.cuh +++ b/cub/block/block_exchange.cuh @@ -37,11 +37,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief The BlockExchange class provides [collective](index.html#sec0) methods for rearranging data partitioned across a CUDA thread block. ![](transpose_logo.png) @@ -1241,6 +1237,5 @@ public: -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_histogram.cuh b/cub/block/block_histogram.cuh index 030209063b..cff9c206ef 100644 --- a/cub/block/block_histogram.cuh +++ b/cub/block/block_histogram.cuh @@ -38,11 +38,7 @@ #include "../config.cuh" #include "../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -409,6 +405,5 @@ public: }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_load.cuh b/cub/block/block_load.cuh index d689954e0b..cf4f5acc99 100644 --- a/cub/block/block_load.cuh +++ b/cub/block/block_load.cuh @@ -41,11 +41,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIo @@ -1289,6 +1285,5 @@ public: }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_radix_rank.cuh b/cub/block/block_radix_rank.cuh index ffcc7280b0..95678feadd 100644 --- a/cub/block/block_radix_rank.cuh +++ b/cub/block/block_radix_rank.cuh @@ -44,11 +44,7 @@ #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -1120,7 +1116,6 @@ struct BlockRadixRankMatchEarlyCounts }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_radix_sort.cuh b/cub/block/block_radix_sort.cuh index 1fa7184c11..dc08cf10a9 100644 --- a/cub/block/block_radix_sort.cuh +++ b/cub/block/block_radix_sort.cuh @@ -41,11 +41,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief The BlockRadixSort class provides [collective](index.html#sec0) methods for sorting items partitioned across a CUDA thread block using a radix sorting method. ![](sorting_logo.png) @@ -898,6 +894,5 @@ public: * \example example_block_radix_sort.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_raking_layout.cuh b/cub/block/block_raking_layout.cuh index bbacdf3e02..8bd1be13ad 100644 --- a/cub/block/block_raking_layout.cuh +++ b/cub/block/block_raking_layout.cuh @@ -37,11 +37,7 @@ #include "../config.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief BlockRakingLayout provides a conflict-free shared memory layout abstraction for 1D raking across thread block data. ![](raking.png) @@ -145,6 +141,5 @@ struct BlockRakingLayout } }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_reduce.cuh b/cub/block/block_reduce.cuh index 1bf971f0f4..2dd85e6591 100644 --- a/cub/block/block_reduce.cuh +++ b/cub/block/block_reduce.cuh @@ -41,11 +41,7 @@ #include "../util_type.cuh" #include "../thread/thread_operators.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN @@ -602,6 +598,5 @@ public: * \example example_block_reduce.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_scan.cuh b/cub/block/block_scan.cuh index 513ef358bd..279b6c1a21 100644 --- a/cub/block/block_scan.cuh +++ b/cub/block/block_scan.cuh @@ -39,11 +39,7 @@ #include "../util_type.cuh" #include "../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -2136,6 +2132,5 @@ public: * \example example_block_scan.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_shuffle.cuh b/cub/block/block_shuffle.cuh index 39b0d2c116..ad4bd0c378 100644 --- a/cub/block/block_shuffle.cuh +++ b/cub/block/block_shuffle.cuh @@ -37,11 +37,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief The BlockShuffle class provides [collective](index.html#sec0) methods for shuffling data partitioned across a CUDA thread block. @@ -296,6 +292,5 @@ public: }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/block_store.cuh b/cub/block/block_store.cuh index cb00ec7287..dafccfed99 100644 --- a/cub/block/block_store.cuh +++ b/cub/block/block_store.cuh @@ -40,11 +40,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIo @@ -1051,6 +1047,5 @@ public: }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/radix_rank_sort_operations.cuh b/cub/block/radix_rank_sort_operations.cuh index 17a399efb6..da921d2d1a 100644 --- a/cub/block/radix_rank_sort_operations.cuh +++ b/cub/block/radix_rank_sort_operations.cuh @@ -38,11 +38,7 @@ #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** \brief Twiddling keys for radix sort. */ template @@ -144,5 +140,4 @@ struct ShiftDigitExtractor : BaseDigitExtractor } }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_histogram_atomic.cuh b/cub/block/specializations/block_histogram_atomic.cuh index 3be0a3dfa6..93299fa719 100644 --- a/cub/block/specializations/block_histogram_atomic.cuh +++ b/cub/block/specializations/block_histogram_atomic.cuh @@ -35,11 +35,7 @@ #include "../../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -77,6 +73,5 @@ struct BlockHistogramAtomic }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_histogram_sort.cuh b/cub/block/specializations/block_histogram_sort.cuh index 878d27cd0a..5bd2a80b33 100644 --- a/cub/block/specializations/block_histogram_sort.cuh +++ b/cub/block/specializations/block_histogram_sort.cuh @@ -38,11 +38,7 @@ #include "../../config.cuh" #include "../../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN @@ -221,6 +217,5 @@ struct BlockHistogramSort }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_reduce_raking.cuh b/cub/block/specializations/block_reduce_raking.cuh index 3ba5f05418..d828484915 100644 --- a/cub/block/specializations/block_reduce_raking.cuh +++ b/cub/block/specializations/block_reduce_raking.cuh @@ -39,11 +39,7 @@ #include "../../config.cuh" #include "../../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -221,6 +217,5 @@ struct BlockReduceRaking }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_reduce_raking_commutative_only.cuh b/cub/block/specializations/block_reduce_raking_commutative_only.cuh index df118031ab..4dd8a13050 100644 --- a/cub/block/specializations/block_reduce_raking_commutative_only.cuh +++ b/cub/block/specializations/block_reduce_raking_commutative_only.cuh @@ -39,11 +39,7 @@ #include "../../config.cuh" #include "../../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -195,6 +191,5 @@ struct BlockReduceRakingCommutativeOnly }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_reduce_warp_reductions.cuh b/cub/block/specializations/block_reduce_warp_reductions.cuh index 4dd3451b88..6b440e5a9c 100644 --- a/cub/block/specializations/block_reduce_warp_reductions.cuh +++ b/cub/block/specializations/block_reduce_warp_reductions.cuh @@ -37,11 +37,7 @@ #include "../../config.cuh" #include "../../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -212,6 +208,5 @@ struct BlockReduceWarpReductions }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_scan_raking.cuh b/cub/block/specializations/block_scan_raking.cuh index 8f20818bfc..863939e1fe 100644 --- a/cub/block/specializations/block_scan_raking.cuh +++ b/cub/block/specializations/block_scan_raking.cuh @@ -41,11 +41,7 @@ #include "../../thread/thread_scan.cuh" #include "../../warp/warp_scan.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -660,6 +656,5 @@ struct BlockScanRaking }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_scan_warp_scans.cuh b/cub/block/specializations/block_scan_warp_scans.cuh index 3835e484e1..8273e43bf4 100644 --- a/cub/block/specializations/block_scan_warp_scans.cuh +++ b/cub/block/specializations/block_scan_warp_scans.cuh @@ -37,11 +37,7 @@ #include "../../util_ptx.cuh" #include "../../warp/warp_scan.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief BlockScanWarpScans provides warpscan-based variants of parallel prefix scan across a CUDA thread block. @@ -386,6 +382,5 @@ struct BlockScanWarpScans }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_scan_warp_scans2.cuh b/cub/block/specializations/block_scan_warp_scans2.cuh index 6617160d1b..a485356bc3 100644 --- a/cub/block/specializations/block_scan_warp_scans2.cuh +++ b/cub/block/specializations/block_scan_warp_scans2.cuh @@ -37,11 +37,7 @@ #include "../../util_ptx.cuh" #include "../../warp/warp_scan.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief BlockScanWarpScans provides warpscan-based variants of parallel prefix scan across a CUDA thread block. @@ -430,6 +426,5 @@ struct BlockScanWarpScans }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/block/specializations/block_scan_warp_scans3.cuh b/cub/block/specializations/block_scan_warp_scans3.cuh index a8279d5765..dad06fd298 100644 --- a/cub/block/specializations/block_scan_warp_scans3.cuh +++ b/cub/block/specializations/block_scan_warp_scans3.cuh @@ -37,11 +37,7 @@ #include "../../util_ptx.cuh" #include "../../warp/warp_scan.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief BlockScanWarpScans provides warpscan-based variants of parallel prefix scan across a CUDA thread block. @@ -412,6 +408,5 @@ struct BlockScanWarpScans }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_histogram.cuh b/cub/device/device_histogram.cuh index 684bd44256..39e14ad5b5 100644 --- a/cub/device/device_histogram.cuh +++ b/cub/device/device_histogram.cuh @@ -41,11 +41,7 @@ #include "dispatch/dispatch_histogram.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -860,7 +856,6 @@ struct DeviceHistogram //@} end member group }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_partition.cuh b/cub/device/device_partition.cuh index 5760eadfcc..17b741944c 100644 --- a/cub/device/device_partition.cuh +++ b/cub/device/device_partition.cuh @@ -40,11 +40,7 @@ #include "dispatch/dispatch_select_if.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -267,7 +263,6 @@ struct DevicePartition * \example example_device_partition_if.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_radix_sort.cuh b/cub/device/device_radix_sort.cuh index bfe914ac2f..4d540568a1 100644 --- a/cub/device/device_radix_sort.cuh +++ b/cub/device/device_radix_sort.cuh @@ -40,11 +40,7 @@ #include "dispatch/dispatch_radix_sort.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -904,7 +900,6 @@ struct DeviceRadixSort * \example example_device_radix_sort.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_reduce.cuh b/cub/device/device_reduce.cuh index 4f01c2446a..4a217de7e9 100644 --- a/cub/device/device_reduce.cuh +++ b/cub/device/device_reduce.cuh @@ -43,11 +43,7 @@ #include "dispatch/dispatch_reduce_by_key.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -728,7 +724,6 @@ struct DeviceReduce * \example example_device_reduce.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_run_length_encode.cuh b/cub/device/device_run_length_encode.cuh index e31ebf0142..ccdfe39331 100644 --- a/cub/device/device_run_length_encode.cuh +++ b/cub/device/device_run_length_encode.cuh @@ -41,11 +41,7 @@ #include "dispatch/dispatch_rle.cuh" #include "dispatch/dispatch_reduce_by_key.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -272,7 +268,6 @@ struct DeviceRunLengthEncode }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_scan.cuh b/cub/device/device_scan.cuh index 6928c22025..c71d145d6c 100644 --- a/cub/device/device_scan.cuh +++ b/cub/device/device_scan.cuh @@ -40,11 +40,7 @@ #include "../config.cuh" #include "dispatch/dispatch_scan.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -436,7 +432,6 @@ struct DeviceScan * \example example_device_scan.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_segmented_radix_sort.cuh b/cub/device/device_segmented_radix_sort.cuh index 401bc1fd4c..30d3028875 100644 --- a/cub/device/device_segmented_radix_sort.cuh +++ b/cub/device/device_segmented_radix_sort.cuh @@ -40,11 +40,7 @@ #include "../config.cuh" #include "dispatch/dispatch_radix_sort.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -881,7 +877,6 @@ struct DeviceSegmentedRadixSort }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_segmented_reduce.cuh b/cub/device/device_segmented_reduce.cuh index 7179173b23..d64c2cb88b 100644 --- a/cub/device/device_segmented_reduce.cuh +++ b/cub/device/device_segmented_reduce.cuh @@ -43,11 +43,7 @@ #include "../config.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -630,7 +626,6 @@ struct DeviceSegmentedReduce }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_select.cuh b/cub/device/device_select.cuh index 136d26044a..95e145e399 100644 --- a/cub/device/device_select.cuh +++ b/cub/device/device_select.cuh @@ -40,11 +40,7 @@ #include "dispatch/dispatch_select_if.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -363,7 +359,6 @@ struct DeviceSelect * \example example_device_select_unique.cu */ -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/device_spmv.cuh b/cub/device/device_spmv.cuh index 0be0c20e7b..d4bac08468 100644 --- a/cub/device/device_spmv.cuh +++ b/cub/device/device_spmv.cuh @@ -41,11 +41,7 @@ #include "dispatch/dispatch_spmv_orig.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -168,7 +164,6 @@ struct DeviceSpmv -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/dispatch/dispatch_histogram.cuh b/cub/device/dispatch/dispatch_histogram.cuh index b62af21a82..e103ad7886 100644 --- a/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/device/dispatch/dispatch_histogram.cuh @@ -48,11 +48,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN @@ -590,7 +586,7 @@ struct DispatchHistogram histogram_init_grid_dims, histogram_init_block_threads, (long long) stream); // Invoke histogram_init_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( histogram_init_grid_dims, histogram_init_block_threads, 0, stream ).doit(histogram_init_kernel, @@ -608,7 +604,7 @@ struct DispatchHistogram histogram_sweep_config.block_threads, (long long) stream, histogram_sweep_config.pixels_per_thread, histogram_sweep_sm_occupancy); // Invoke histogram_sweep_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( sweep_grid_dims, histogram_sweep_config.block_threads, 0, stream ).doit(histogram_sweep_kernel, d_samples, @@ -1015,7 +1011,6 @@ struct DispatchHistogram }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/device/dispatch/dispatch_radix_sort.cuh index 0f97152193..f48371e7c8 100644 --- a/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/device/dispatch/dispatch_radix_sort.cuh @@ -59,11 +59,7 @@ # pragma clang diagnostic ignored "-Wpass-failed" #endif -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -1065,7 +1061,7 @@ struct DispatchRadixSort : ActivePolicyT::SingleTilePolicy::ITEMS_PER_THREAD, 1, begin_bit, ActivePolicyT::SingleTilePolicy::RADIX_BITS); // Invoke upsweep_kernel with same grid size as downsweep_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( 1, ActivePolicyT::SingleTilePolicy::BLOCK_THREADS, 0, stream ).doit(single_tile_kernel, d_keys.Current(), @@ -1128,7 +1124,7 @@ struct DispatchRadixSort : int pass_spine_length = pass_config.even_share.grid_size * pass_config.radix_digits; // Invoke upsweep_kernel with same grid size as downsweep_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( pass_config.even_share.grid_size, pass_config.upsweep_config.block_threads, 0, stream ).doit(pass_config.upsweep_kernel, @@ -1150,7 +1146,7 @@ struct DispatchRadixSort : 1, pass_config.scan_config.block_threads, (long long) stream, pass_config.scan_config.items_per_thread); // Invoke scan_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( 1, pass_config.scan_config.block_threads, 0, stream ).doit(pass_config.scan_kernel, d_spine, @@ -1168,7 +1164,7 @@ struct DispatchRadixSort : pass_config.downsweep_config.items_per_thread, pass_config.downsweep_config.sm_occupancy); // Invoke downsweep_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( pass_config.even_share.grid_size, pass_config.downsweep_config.block_threads, 0, stream ).doit(pass_config.downsweep_kernel, @@ -1750,7 +1746,7 @@ struct DispatchSegmentedRadixSort : pass_bits); } - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( num_segments, pass_config.segmented_config.block_threads, 0, stream ).doit(pass_config.segmented_kernel, @@ -1963,8 +1959,7 @@ struct DispatchSegmentedRadixSort : }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END #if defined(__clang__) diff --git a/cub/device/dispatch/dispatch_reduce.cuh b/cub/device/dispatch/dispatch_reduce.cuh index 8fc9907281..9bdf89f63f 100644 --- a/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/device/dispatch/dispatch_reduce.cuh @@ -48,11 +48,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** * Kernel entry points @@ -415,7 +411,7 @@ struct DispatchReduce : ActivePolicyT::SingleTilePolicy::ITEMS_PER_THREAD); // Invoke single_reduce_sweep_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( 1, ActivePolicyT::SingleTilePolicy::BLOCK_THREADS, 0, stream ).doit(single_tile_kernel, d_in, @@ -511,7 +507,7 @@ struct DispatchReduce : reduce_config.sm_occupancy); // Invoke DeviceReduceKernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( reduce_grid_size, ActivePolicyT::ReducePolicy::BLOCK_THREADS, 0, stream ).doit(reduce_kernel, @@ -534,7 +530,7 @@ struct DispatchReduce : ActivePolicyT::SingleTilePolicy::ITEMS_PER_THREAD); // Invoke DeviceReduceSingleTileKernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( 1, ActivePolicyT::SingleTilePolicy::BLOCK_THREADS, 0, stream ).doit(single_tile_kernel, d_block_reductions, @@ -751,7 +747,7 @@ struct DispatchSegmentedReduce : segmented_reduce_config.sm_occupancy); // Invoke DeviceReduceKernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( num_segments, ActivePolicyT::SegmentedReducePolicy::BLOCK_THREADS, 0, stream ).doit(segmented_reduce_kernel, @@ -843,7 +839,6 @@ struct DispatchSegmentedReduce : -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/device/dispatch/dispatch_reduce_by_key.cuh index b22fb78323..e3064fd7d0 100644 --- a/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -47,11 +47,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** * Kernel entry points @@ -334,7 +330,7 @@ struct DispatchReduceByKey if (debug_synchronous) _CubLog("Invoking init_kernel<<<%d, %d, 0, %lld>>>()\n", init_grid_size, INIT_KERNEL_THREADS, (long long) stream); // Invoke init_kernel to initialize tile descriptors - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( init_grid_size, INIT_KERNEL_THREADS, 0, stream ).doit(init_kernel, tile_state, @@ -371,7 +367,7 @@ struct DispatchReduceByKey start_tile, scan_grid_size, reduce_by_key_config.block_threads, (long long) stream, reduce_by_key_config.items_per_thread, reduce_by_key_sm_occupancy); // Invoke reduce_by_key_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( scan_grid_size, reduce_by_key_config.block_threads, 0, stream ).doit(reduce_by_key_kernel, @@ -455,7 +451,6 @@ struct DispatchReduceByKey } }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/dispatch/dispatch_rle.cuh b/cub/device/dispatch/dispatch_rle.cuh index 25bdb7abc1..3835abf0fb 100644 --- a/cub/device/dispatch/dispatch_rle.cuh +++ b/cub/device/dispatch/dispatch_rle.cuh @@ -47,11 +47,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -316,7 +312,7 @@ struct DeviceRleDispatch if (debug_synchronous) _CubLog("Invoking device_scan_init_kernel<<<%d, %d, 0, %lld>>>()\n", init_grid_size, INIT_KERNEL_THREADS, (long long) stream); // Invoke device_scan_init_kernel to initialize tile descriptors and queue descriptors - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( init_grid_size, INIT_KERNEL_THREADS, 0, stream ).doit(device_scan_init_kernel, tile_status, @@ -355,7 +351,7 @@ struct DeviceRleDispatch scan_grid_size.x, scan_grid_size.y, scan_grid_size.z, device_rle_config.block_threads, (long long) stream, device_rle_config.items_per_thread, device_rle_kernel_sm_occupancy); // Invoke device_rle_sweep_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( scan_grid_size, device_rle_config.block_threads, 0, stream ).doit(device_rle_sweep_kernel, d_in, @@ -433,7 +429,6 @@ struct DeviceRleDispatch }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/dispatch/dispatch_scan.cuh b/cub/device/dispatch/dispatch_scan.cuh index 2946d1a2e7..7908f8f0b4 100644 --- a/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/device/dispatch/dispatch_scan.cuh @@ -47,11 +47,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -327,7 +323,7 @@ struct DispatchScan: if (debug_synchronous) _CubLog("Invoking init_kernel<<<%d, %d, 0, %lld>>>()\n", init_grid_size, INIT_KERNEL_THREADS, (long long) stream); // Invoke init_kernel to initialize tile descriptors - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( init_grid_size, INIT_KERNEL_THREADS, 0, stream ).doit(init_kernel, tile_state, @@ -360,7 +356,7 @@ struct DispatchScan: start_tile, scan_grid_size, Policy::BLOCK_THREADS, (long long) stream, Policy::ITEMS_PER_THREAD, scan_sm_occupancy); // Invoke scan_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( scan_grid_size, Policy::BLOCK_THREADS, 0, stream ).doit(scan_kernel, d_in, @@ -447,5 +443,4 @@ struct DispatchScan: -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/dispatch/dispatch_select_if.cuh b/cub/device/dispatch/dispatch_select_if.cuh index ebb6b5b98c..73c2bd9522 100644 --- a/cub/device/dispatch/dispatch_select_if.cuh +++ b/cub/device/dispatch/dispatch_select_if.cuh @@ -47,11 +47,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** * Kernel entry points @@ -322,7 +318,7 @@ struct DispatchSelectIf if (debug_synchronous) _CubLog("Invoking scan_init_kernel<<<%d, %d, 0, %lld>>>()\n", init_grid_size, INIT_KERNEL_THREADS, (long long) stream); // Invoke scan_init_kernel to initialize tile descriptors - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( init_grid_size, INIT_KERNEL_THREADS, 0, stream ).doit(scan_init_kernel, tile_status, @@ -361,7 +357,7 @@ struct DispatchSelectIf scan_grid_size.x, scan_grid_size.y, scan_grid_size.z, select_if_config.block_threads, (long long) stream, select_if_config.items_per_thread, range_select_sm_occupancy); // Invoke select_if_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( scan_grid_size, select_if_config.block_threads, 0, stream ).doit(select_if_kernel, d_in, @@ -441,7 +437,6 @@ struct DispatchSelectIf }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/device/dispatch/dispatch_spmv_orig.cuh b/cub/device/dispatch/dispatch_spmv_orig.cuh index e821ff23e3..fc920fd047 100644 --- a/cub/device/dispatch/dispatch_spmv_orig.cuh +++ b/cub/device/dispatch/dispatch_spmv_orig.cuh @@ -50,11 +50,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -518,7 +514,7 @@ struct DispatchSpmv degen_col_kernel_grid_size, degen_col_kernel_block_size, (long long) stream); // Invoke spmv_search_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( degen_col_kernel_grid_size, degen_col_kernel_block_size, 0, stream ).doit(spmv_1col_kernel, @@ -630,7 +626,7 @@ struct DispatchSpmv search_grid_size, search_block_size, (long long) stream); // Invoke spmv_search_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( search_grid_size, search_block_size, 0, stream ).doit(spmv_search_kernel, num_merge_tiles, @@ -649,7 +645,7 @@ struct DispatchSpmv spmv_grid_size.x, spmv_grid_size.y, spmv_grid_size.z, spmv_config.block_threads, (long long) stream, spmv_config.items_per_thread, spmv_sm_occupancy); // Invoke spmv_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( spmv_grid_size, spmv_config.block_threads, 0, stream ).doit(spmv_kernel, spmv_params, @@ -673,7 +669,7 @@ struct DispatchSpmv segment_fixup_grid_size.x, segment_fixup_grid_size.y, segment_fixup_grid_size.z, segment_fixup_config.block_threads, (long long) stream, segment_fixup_config.items_per_thread, segment_fixup_sm_occupancy); // Invoke segment_fixup_kernel - thrust::cuda_cub::launcher::triple_chevron( + THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron( segment_fixup_grid_size, segment_fixup_config.block_threads, 0, stream ).doit(segment_fixup_kernel, @@ -744,7 +740,6 @@ struct DispatchSpmv }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/grid/grid_barrier.cuh b/cub/grid/grid_barrier.cuh index 1bcb533ee4..ca946f3cc8 100644 --- a/cub/grid/grid_barrier.cuh +++ b/cub/grid/grid_barrier.cuh @@ -37,11 +37,7 @@ #include "../config.cuh" #include "../thread/thread_load.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -206,6 +202,5 @@ public: /** @} */ // end group GridModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/grid/grid_even_share.cuh b/cub/grid/grid_even_share.cuh index badbfd62b7..736ed12577 100644 --- a/cub/grid/grid_even_share.cuh +++ b/cub/grid/grid_even_share.cuh @@ -41,11 +41,7 @@ #include "../util_type.cuh" #include "grid_mapping.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -222,5 +218,4 @@ public: /** @} */ // end group GridModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/grid/grid_mapping.cuh b/cub/grid/grid_mapping.cuh index 889a94c96e..b57f193deb 100644 --- a/cub/grid/grid_mapping.cuh +++ b/cub/grid/grid_mapping.cuh @@ -35,11 +35,7 @@ #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -108,6 +104,5 @@ enum GridMappingStrategy /** @} */ // end group GridModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/grid/grid_queue.cuh b/cub/grid/grid_queue.cuh index 6b5f676b03..ebb82e4dc0 100644 --- a/cub/grid/grid_queue.cuh +++ b/cub/grid/grid_queue.cuh @@ -36,11 +36,7 @@ #include "../config.cuh" #include "../util_debug.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -238,7 +234,6 @@ __global__ void FillAndResetDrainKernel( /** @} */ // end group GridModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/host/mutex.cuh b/cub/host/mutex.cuh index 9880dee57c..a9c3dc7cc3 100644 --- a/cub/host/mutex.cuh +++ b/cub/host/mutex.cuh @@ -58,11 +58,7 @@ #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -162,6 +158,5 @@ struct Mutex -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/arg_index_input_iterator.cuh b/cub/iterator/arg_index_input_iterator.cuh index f16fab8c26..fc6e02224e 100644 --- a/cub/iterator/arg_index_input_iterator.cuh +++ b/cub/iterator/arg_index_input_iterator.cuh @@ -49,11 +49,7 @@ #include #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIterator @@ -123,9 +119,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::any_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::any_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -255,5 +251,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/cache_modified_input_iterator.cuh b/cub/iterator/cache_modified_input_iterator.cuh index 7a41a5d31a..8c60cd8d58 100644 --- a/cub/iterator/cache_modified_input_iterator.cuh +++ b/cub/iterator/cache_modified_input_iterator.cuh @@ -48,11 +48,7 @@ #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN @@ -117,9 +113,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::device_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::device_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -236,5 +232,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/cache_modified_output_iterator.cuh b/cub/iterator/cache_modified_output_iterator.cuh index e1697013c4..2536981e67 100644 --- a/cub/iterator/cache_modified_output_iterator.cuh +++ b/cub/iterator/cache_modified_output_iterator.cuh @@ -48,11 +48,7 @@ #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -137,9 +133,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::device_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::device_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -250,5 +246,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/constant_input_iterator.cuh b/cub/iterator/constant_input_iterator.cuh index 44fb56c920..dc2d72e897 100644 --- a/cub/iterator/constant_input_iterator.cuh +++ b/cub/iterator/constant_input_iterator.cuh @@ -47,11 +47,7 @@ #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -106,9 +102,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::any_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::any_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -231,5 +227,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/counting_input_iterator.cuh b/cub/iterator/counting_input_iterator.cuh index c7167a7066..e81f1d9680 100644 --- a/cub/iterator/counting_input_iterator.cuh +++ b/cub/iterator/counting_input_iterator.cuh @@ -48,11 +48,7 @@ #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIterator @@ -104,9 +100,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::any_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::any_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -224,5 +220,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/discard_output_iterator.cuh b/cub/iterator/discard_output_iterator.cuh index 2db03219ae..fe2ccca7e9 100644 --- a/cub/iterator/discard_output_iterator.cuh +++ b/cub/iterator/discard_output_iterator.cuh @@ -45,11 +45,7 @@ #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -75,9 +71,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::any_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::any_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -209,5 +205,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/tex_obj_input_iterator.cuh b/cub/iterator/tex_obj_input_iterator.cuh index 4c286813c3..255e6da0b9 100644 --- a/cub/iterator/tex_obj_input_iterator.cuh +++ b/cub/iterator/tex_obj_input_iterator.cuh @@ -49,11 +49,7 @@ #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIterator @@ -121,9 +117,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::device_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::device_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -319,5 +315,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/iterator/tex_ref_input_iterator.cuh b/cub/iterator/tex_ref_input_iterator.cuh index 8f23928ca4..f02dc431be 100644 --- a/cub/iterator/tex_ref_input_iterator.cuh +++ b/cub/iterator/tex_ref_input_iterator.cuh @@ -49,11 +49,7 @@ #include #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /****************************************************************************** @@ -248,9 +244,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::device_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::device_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -416,7 +412,6 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END #endif // CUDART_VERSION diff --git a/cub/iterator/transform_input_iterator.cuh b/cub/iterator/transform_input_iterator.cuh index bce8b817d7..b92821ceb9 100644 --- a/cub/iterator/transform_input_iterator.cuh +++ b/cub/iterator/transform_input_iterator.cuh @@ -48,11 +48,7 @@ #endif // THRUST_VERSION -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIterator @@ -127,9 +123,9 @@ public: #if (THRUST_VERSION >= 100700) // Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods - typedef typename thrust::detail::iterator_facade_category< - thrust::any_system_tag, - thrust::random_access_traversal_tag, + typedef typename THRUST_NS_QUALIFIER::detail::iterator_facade_category< + THRUST_NS_QUALIFIER::any_system_tag, + THRUST_NS_QUALIFIER::random_access_traversal_tag, value_type, reference >::type iterator_category; ///< The iterator category @@ -248,5 +244,4 @@ public: /** @} */ // end group UtilIterator -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/thread/thread_load.cuh b/cub/thread/thread_load.cuh index 31e759602f..5402f80f22 100644 --- a/cub/thread/thread_load.cuh +++ b/cub/thread/thread_load.cuh @@ -39,11 +39,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIo @@ -423,5 +419,4 @@ __device__ __forceinline__ typename std::iterator_traits::value_ /** @} */ // end group UtilIo -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/thread/thread_operators.cuh b/cub/thread/thread_operators.cuh index 6a3192bca3..3a75a79015 100644 --- a/cub/thread/thread_operators.cuh +++ b/cub/thread/thread_operators.cuh @@ -40,11 +40,7 @@ #include "../config.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -312,5 +308,4 @@ struct ReduceByKeyOp /** @} */ // end group UtilModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/thread/thread_reduce.cuh b/cub/thread/thread_reduce.cuh index 41063f9714..98fb2faab3 100644 --- a/cub/thread/thread_reduce.cuh +++ b/cub/thread/thread_reduce.cuh @@ -36,11 +36,7 @@ #include "../thread/thread_operators.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /// Internal namespace (to prevent ADL mishaps between static functions when mixing different CUB installations) namespace internal { @@ -148,5 +144,4 @@ __device__ __forceinline__ T ThreadReduce( } // internal namespace -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/thread/thread_scan.cuh b/cub/thread/thread_scan.cuh index fd907fcae1..b5e42710fc 100644 --- a/cub/thread/thread_scan.cuh +++ b/cub/thread/thread_scan.cuh @@ -36,11 +36,7 @@ #include "../config.cuh" #include "../thread/thread_operators.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /// Internal namespace (to prevent ADL mishaps between static functions when mixing different CUB installations) namespace internal { @@ -264,5 +260,4 @@ __device__ __forceinline__ T ThreadScanInclusive( } // internal namespace -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/thread/thread_search.cuh b/cub/thread/thread_search.cuh index 96b9e65a56..0182896d6d 100644 --- a/cub/thread/thread_search.cuh +++ b/cub/thread/thread_search.cuh @@ -37,11 +37,7 @@ #include "../util_namespace.cuh" #include "../config.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -152,5 +148,4 @@ __device__ __forceinline__ OffsetT UpperBound( -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/thread/thread_store.cuh b/cub/thread/thread_store.cuh index 47d6c6145f..7fc3af15f0 100644 --- a/cub/thread/thread_store.cuh +++ b/cub/thread/thread_store.cuh @@ -37,11 +37,7 @@ #include "../util_ptx.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup UtilIo @@ -416,5 +412,4 @@ __device__ __forceinline__ void ThreadStore(OutputIteratorT itr, T val) /** @} */ // end group UtilIo -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_allocator.cuh b/cub/util_allocator.cuh index dcdbb5cd89..b6d2015382 100644 --- a/cub/util_allocator.cuh +++ b/cub/util_allocator.cuh @@ -42,11 +42,7 @@ #include "host/mutex.cuh" #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -718,5 +714,4 @@ struct CachingDeviceAllocator /** @} */ // end group UtilMgmt -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_arch.cuh b/cub/util_arch.cuh index 82acfefbcb..ccdc5092fe 100644 --- a/cub/util_arch.cuh +++ b/cub/util_arch.cuh @@ -37,11 +37,7 @@ #include "util_namespace.cuh" #include "util_macro.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document @@ -183,5 +179,4 @@ struct MemBoundScaling #endif // Do not document -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_debug.cuh b/cub/util_debug.cuh index 8413f7bd4e..402dcd78d6 100644 --- a/cub/util_debug.cuh +++ b/cub/util_debug.cuh @@ -40,11 +40,7 @@ #include "util_namespace.cuh" #include "util_arch.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -102,7 +98,7 @@ __host__ __device__ __forceinline__ cudaError_t Debug( * \brief Debug macro */ #ifndef CubDebug - #define CubDebug(e) cub::Debug((cudaError_t) (e), __FILE__, __LINE__) + #define CubDebug(e) CUB_NS_QUALIFIER::Debug((cudaError_t) (e), __FILE__, __LINE__) #endif @@ -110,7 +106,7 @@ __host__ __device__ __forceinline__ cudaError_t Debug( * \brief Debug macro with exit */ #ifndef CubDebugExit - #define CubDebugExit(e) if (cub::Debug((cudaError_t) (e), __FILE__, __LINE__)) { exit(1); } + #define CubDebugExit(e) if (CUB_NS_QUALIFIER::Debug((cudaError_t) (e), __FILE__, __LINE__)) { exit(1); } #endif @@ -146,9 +142,9 @@ __host__ __device__ __forceinline__ cudaError_t Debug( #endif } #ifndef __CUDA_ARCH__ - #define _CubLog(format, ...) cub::va_printf(format,__VA_ARGS__); + #define _CubLog(format, ...) CUB_NS_QUALIFIER::va_printf(format,__VA_ARGS__); #else - #define _CubLog(format, ...) cub::va_printf("[block (%d,%d,%d), thread (%d,%d,%d)]: " format, __VA_ARGS__); + #define _CubLog(format, ...) CUB_NS_QUALIFIER::va_printf("[block (%d,%d,%d), thread (%d,%d,%d)]: " format, __VA_ARGS__); #endif #endif #endif @@ -158,5 +154,4 @@ __host__ __device__ __forceinline__ cudaError_t Debug( /** @} */ // end group UtilMgmt -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_device.cuh b/cub/util_device.cuh index 7cceb59453..e20bed97cd 100644 --- a/cub/util_device.cuh +++ b/cub/util_device.cuh @@ -46,11 +46,7 @@ #include #endif -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -402,6 +398,7 @@ CUB_RUNTIME_FUNCTION inline cudaError_t PtxVersionUncached(int& ptx_version) __host__ inline cudaError_t PtxVersionUncached(int& ptx_version, int device) { SwitchDevice sd(device); + (void)sd; return PtxVersionUncached(ptx_version); } @@ -712,5 +709,4 @@ struct ChainedPolicy /** @} */ // end group UtilMgmt -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_macro.cuh b/cub/util_macro.cuh index ff86365422..4050b8d7c5 100644 --- a/cub/util_macro.cuh +++ b/cub/util_macro.cuh @@ -34,11 +34,7 @@ #include "util_namespace.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -99,5 +95,4 @@ namespace cub { /** @} */ // end group UtilModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_math.cuh b/cub/util_math.cuh index 21bf843e12..78a045f48b 100644 --- a/cub/util_math.cuh +++ b/cub/util_math.cuh @@ -36,12 +36,7 @@ #include "util_namespace.cuh" -// Optional outer namespace(s) -CUB_NS_PREFIX - -// CUB namespace -namespace cub -{ +CUB_NAMESPACE_BEGIN namespace detail { @@ -71,5 +66,4 @@ DivideAndRoundUp(NumeratorT n, DenominatorT d) return static_cast(n / d + (n % d != 0 ? 1 : 0)); } -} // namespace cub -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_namespace.cuh b/cub/util_namespace.cuh index 5f322d5d65..b94f05e255 100644 --- a/cub/util_namespace.cuh +++ b/cub/util_namespace.cuh @@ -1,6 +1,6 @@ /****************************************************************************** * Copyright (c) 2011, Duane Merrill. All rights reserved. - * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2011-2021, NVIDIA CORPORATION. All rights reserved. * * Redistribution and use in source and binary forms, with or without * modification, are permitted provided that the following conditions are met: @@ -27,26 +27,102 @@ ******************************************************************************/ /** - * \file - * Place-holder for prefixing the cub namespace + * \file util_namespace.cuh + * \brief Utilities that allow `cub::` to be placed inside an + * application-specific namespace. */ + #pragma once +// This is not used by this file; this is a hack so that we can detect the +// CUB version from Thrust on older versions of CUB that did not have +// version.cuh. #include "version.cuh" -// For example: -//#define CUB_NS_PREFIX namespace thrust{ namespace detail { -//#define CUB_NS_POSTFIX } } +/** + * \def THRUST_CUB_WRAPPED_NAMESPACE + * If defined, this value will be used as the name of a namespace that wraps the + * `thrust::` and `cub::` namespaces. + * This macro should not be used with any other CUB namespace macros. + */ +#ifdef THRUST_CUB_WRAPPED_NAMESPACE +#define CUB_WRAPPED_NAMESPACE THRUST_CUB_WRAPPED_NAMESPACE +#endif + +/** + * \def CUB_WRAPPED_NAMESPACE + * If defined, this value will be used as the name of a namespace that wraps the + * `cub::` namespace. + * If THRUST_CUB_WRAPPED_NAMESPACE is set, this will inherit that macro's value. + * This macro should not be used with any other CUB namespace macros. + */ +#ifdef CUB_WRAPPED_NAMESPACE +#define CUB_NS_PREFIX \ + namespace CUB_WRAPPED_NAMESPACE \ + { +#define CUB_NS_POSTFIX } + +#define CUB_NS_QUALIFIER ::CUB_WRAPPED_NAMESPACE::cub +#endif + +/** + * \def CUB_NS_PREFIX + * This macro is inserted prior to all `namespace cub { ... }` blocks. It is + * derived from CUB_WRAPPED_NAMESPACE, if set, and will be empty otherwise. + * It may be defined by users, in which case CUB_NS_PREFIX, + * CUB_NS_POSTFIX, and CUB_NS_QUALIFIER must all be set consistently. + */ #ifndef CUB_NS_PREFIX #define CUB_NS_PREFIX #endif +/** + * \def CUB_NS_POSTFIX + * This macro is inserted following the closing braces of all + * `namespace cub { ... }` block. It is defined appropriately when + * CUB_WRAPPED_NAMESPACE is set, and will be empty otherwise. It may be + * defined by users, in which case CUB_NS_PREFIX, CUB_NS_POSTFIX, and + * CUB_NS_QUALIFIER must all be set consistently. + */ #ifndef CUB_NS_POSTFIX #define CUB_NS_POSTFIX #endif +/** + * \def CUB_NS_QUALIFIER + * This macro is used to qualify members of cub:: when accessing them from + * outside of their namespace. By default, this is just `::cub`, and will be + * set appropriately when CUB_WRAPPED_NAMESPACE is defined. This macro may be + * defined by users, in which case CUB_NS_PREFIX, CUB_NS_POSTFIX, and + * CUB_NS_QUALIFIER must all be set consistently. + */ +#ifndef CUB_NS_QUALIFIER +#define CUB_NS_QUALIFIER ::cub +#endif + +/** + * \def CUB_NAMESPACE_BEGIN + * This macro is used to open a `cub::` namespace block, along with any + * enclosing namespaces requested by CUB_WRAPPED_NAMESPACE, etc. + * This macro is defined by CUB and may not be overridden. + */ +#define CUB_NAMESPACE_BEGIN \ + CUB_NS_PREFIX \ + namespace cub \ + { + +/** + * \def CUB_NAMESPACE_END + * This macro is used to close a `cub::` namespace block, along with any + * enclosing namespaces requested by CUB_WRAPPED_NAMESPACE, etc. + * This macro is defined by CUB and may not be overridden. + */ +#define CUB_NAMESPACE_END \ + } /* end namespace cub */ \ + CUB_NS_POSTFIX + // Declare these namespaces here for the purpose of Doxygenating them CUB_NS_PREFIX @@ -56,7 +132,6 @@ CUB_NS_PREFIX */ namespace cub { - } CUB_NS_POSTFIX diff --git a/cub/util_ptx.cuh b/cub/util_ptx.cuh index 7b3ce7af21..ea63d7c853 100644 --- a/cub/util_ptx.cuh +++ b/cub/util_ptx.cuh @@ -40,11 +40,7 @@ #include "util_debug.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -736,5 +732,4 @@ inline __device__ unsigned int MatchAny(unsigned int label) } -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/util_type.cuh b/cub/util_type.cuh index c38051078e..5c57172b1b 100644 --- a/cub/util_type.cuh +++ b/cub/util_type.cuh @@ -50,11 +50,7 @@ -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -1184,5 +1180,4 @@ struct Traits : NumericTraits::Type> {}; /** @} */ // end group UtilModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/warp/specializations/warp_reduce_shfl.cuh b/cub/warp/specializations/warp_reduce_shfl.cuh index dbc56ec1bc..deb10c3fce 100644 --- a/cub/warp/specializations/warp_reduce_shfl.cuh +++ b/cub/warp/specializations/warp_reduce_shfl.cuh @@ -40,11 +40,7 @@ #include -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -538,5 +534,4 @@ struct WarpReduceShfl }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/warp/specializations/warp_reduce_smem.cuh b/cub/warp/specializations/warp_reduce_smem.cuh index 2442a8c4f2..4b0855e534 100644 --- a/cub/warp/specializations/warp_reduce_smem.cuh +++ b/cub/warp/specializations/warp_reduce_smem.cuh @@ -39,11 +39,7 @@ #include "../../thread/thread_store.cuh" #include "../../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief WarpReduceSmem provides smem-based variants of parallel reduction of items partitioned across a CUDA thread warp. @@ -368,5 +364,4 @@ struct WarpReduceSmem }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/warp/specializations/warp_scan_shfl.cuh b/cub/warp/specializations/warp_scan_shfl.cuh index 18b46dd998..c85f27ac45 100644 --- a/cub/warp/specializations/warp_scan_shfl.cuh +++ b/cub/warp/specializations/warp_scan_shfl.cuh @@ -38,11 +38,7 @@ #include "../../util_type.cuh" #include "../../util_ptx.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief WarpScanShfl provides SHFL-based variants of parallel prefix scan of items partitioned across a CUDA thread warp. @@ -628,5 +624,4 @@ struct WarpScanShfl }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/warp/specializations/warp_scan_smem.cuh b/cub/warp/specializations/warp_scan_smem.cuh index ccd1de30f6..cebe54ff4f 100644 --- a/cub/warp/specializations/warp_scan_smem.cuh +++ b/cub/warp/specializations/warp_scan_smem.cuh @@ -39,11 +39,7 @@ #include "../../thread/thread_store.cuh" #include "../../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \brief WarpScanSmem provides smem-based variants of parallel prefix scan of items partitioned across a CUDA thread warp. @@ -393,5 +389,4 @@ struct WarpScanSmem }; -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/warp/warp_reduce.cuh b/cub/warp/warp_reduce.cuh index 50ee7056c3..c49ae14cfe 100644 --- a/cub/warp/warp_reduce.cuh +++ b/cub/warp/warp_reduce.cuh @@ -39,11 +39,7 @@ #include "../thread/thread_operators.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** @@ -607,5 +603,4 @@ public: /** @} */ // end group WarpModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/cub/warp/warp_scan.cuh b/cub/warp/warp_scan.cuh index e9e95008a3..bc7ad40f7d 100644 --- a/cub/warp/warp_scan.cuh +++ b/cub/warp/warp_scan.cuh @@ -39,11 +39,7 @@ #include "../thread/thread_operators.cuh" #include "../util_type.cuh" -/// Optional outer namespace(s) -CUB_NS_PREFIX - -/// CUB namespace -namespace cub { +CUB_NAMESPACE_BEGIN /** * \addtogroup WarpModule @@ -931,5 +927,4 @@ public: /** @} */ // end group WarpModule -} // CUB namespace -CUB_NS_POSTFIX // Optional outer namespace(s) +CUB_NAMESPACE_END diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 87ed7532fa..444e0de6e1 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -135,3 +135,5 @@ foreach (test_src IN LISTS test_srcs) endif() endforeach() endforeach() + +add_subdirectory(cmake) diff --git a/test/bfloat16.h b/test/bfloat16.h index 4413f9145d..4fdd2c417f 100644 --- a/test/bfloat16.h +++ b/test/bfloat16.h @@ -224,15 +224,18 @@ std::ostream& operator<<(std::ostream &out, const __nv_bfloat16 &x) ******************************************************************************/ template <> -struct cub::FpLimits +struct CUB_NS_QUALIFIER::FpLimits { static __host__ __device__ __forceinline__ bfloat16_t Max() { return bfloat16_t::max(); } static __host__ __device__ __forceinline__ bfloat16_t Lowest() { return bfloat16_t::lowest(); } }; -template <> struct cub::NumericTraits : cub::BaseTraits {}; - +template <> +struct CUB_NS_QUALIFIER::NumericTraits + : CUB_NS_QUALIFIER:: + BaseTraits +{}; #ifdef __GNUC__ #pragma GCC diagnostic pop diff --git a/test/cmake/CMakeLists.txt b/test/cmake/CMakeLists.txt index a2d9af4f82..4062ec8349 100644 --- a/test/cmake/CMakeLists.txt +++ b/test/cmake/CMakeLists.txt @@ -13,3 +13,12 @@ if (NOT CUB_IN_THRUST) # Thrust has its own checks for this: -D "CMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE}" ) endif() + +# Check that namespace macros are used correctly: +add_test( + NAME cub.test.cmake.check_namespace + COMMAND + "${CMAKE_COMMAND}" + -D "CUB_SOURCE_DIR=${CUB_SOURCE_DIR}" + -P "${CMAKE_CURRENT_LIST_DIR}/check_namespace.cmake" +) diff --git a/test/cmake/check_namespace.cmake b/test/cmake/check_namespace.cmake new file mode 100644 index 0000000000..70d671beb2 --- /dev/null +++ b/test/cmake/check_namespace.cmake @@ -0,0 +1,89 @@ +# Check all files in thrust to make sure that they use +# CUB_NAMESPACE_BEGIN/END instead of bare `namespace cub {}` declarations. +# +# This is run as a ctest test named `cub.test.cmake.check_namespace`, or +# manually with: +# cmake -D "CUB_SOURCE_DIR=" -P check_namespace.cmake + +cmake_minimum_required(VERSION 3.15) + +set(exclusions + # This defines the macros and must have bare namespace declarations: + cub/util_namespace.cuh +) + +function(count_substrings input search_regex output_var) + string(REGEX MATCHALL "${search_regex}" matches "${input}") + list(LENGTH matches num_matches) + set(${output_var} ${num_matches} PARENT_SCOPE) +endfunction() + +set(bare_ns_regex "namespace[ \n\r\t]+cub[ \n\r\t]*\\{") + +# Validation check for the above regex: +count_substrings([=[ +namespace cub{ +namespace cub { +namespace cub { + namespace cub { +namespace cub +{ +namespace +cub +{ +]=] + ${bare_ns_regex} valid_count) +if (NOT valid_count EQUAL 6) + message(FATAL_ERROR "Validation of bare namespace regex failed: " + "Matched ${valid_count} times, expected 6.") +endif() + +set(found_errors 0) +file(GLOB_RECURSE cub_srcs + RELATIVE "${CUB_SOURCE_DIR}" + "${CUB_SOURCE_DIR}/*.cuh" + "${CUB_SOURCE_DIR}/*.cu" + "${CUB_SOURCE_DIR}/*.h" + "${CUB_SOURCE_DIR}/*.cpp" +) + +foreach(src ${cub_srcs}) + if (${src} IN_LIST exclusions) + continue() + endif() + + file(READ "${CUB_SOURCE_DIR}/${src}" src_contents) + + count_substrings("${src_contents}" "${bare_ns_regex}" bare_ns_count) + count_substrings("${src_contents}" CUB_NS_PREFIX prefix_count) + count_substrings("${src_contents}" CUB_NS_POSTFIX postfix_count) + count_substrings("${src_contents}" CUB_NAMESPACE_BEGIN begin_count) + count_substrings("${src_contents}" CUB_NAMESPACE_END end_count) + + if (NOT bare_ns_count EQUAL 0) + message("'${src}' contains 'namespace cub {...}'. Replace with CUB_NAMESPACE macros.") + set(found_errors 1) + endif() + + if (NOT prefix_count EQUAL 0) + message("'${src}' contains 'CUB_NS_PREFIX'. Replace with CUB_NAMESPACE macros.") + set(found_errors 1) + endif() + + if (NOT postfix_count EQUAL 0) + message("'${src}' contains 'CUB_NS_POSTFIX'. Replace with CUB_NAMESPACE macros.") + set(found_errors 1) + endif() + + if (NOT begin_count EQUAL end_count) + message("'${src}' namespace macros are unbalanced:") + message(" - CUB_NAMESPACE_BEGIN occurs ${begin_count} times.") + message(" - CUB_NAMESPACE_END occurs ${end_count} times.") + set(found_errors 1) + endif() + +endforeach() + +if (NOT found_errors EQUAL 0) + message(FATAL_ERROR "Errors detected.") +endif() diff --git a/test/half.h b/test/half.h index 335eaa6b37..0f9c9660d9 100644 --- a/test/half.h +++ b/test/half.h @@ -300,15 +300,18 @@ std::ostream& operator<<(std::ostream &out, const __half &x) ******************************************************************************/ template <> -struct cub::FpLimits +struct CUB_NS_QUALIFIER::FpLimits { static __host__ __device__ __forceinline__ half_t Max() { return half_t::max(); } static __host__ __device__ __forceinline__ half_t Lowest() { return half_t::lowest(); } }; -template <> struct cub::NumericTraits : cub::BaseTraits {}; - +template <> +struct CUB_NS_QUALIFIER::NumericTraits + : CUB_NS_QUALIFIER:: + BaseTraits +{}; #ifdef __GNUC__ #pragma GCC diagnostic pop diff --git a/test/test_device_radix_sort.cu b/test/test_device_radix_sort.cu index 16780f6587..74bf47f9aa 100644 --- a/test/test_device_radix_sort.cu +++ b/test/test_device_radix_sort.cu @@ -417,11 +417,11 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_keys_wrapper(d_keys.Current()); + THRUST_NS_QUALIFIER::device_ptr d_keys_wrapper(d_keys.Current()); - if (IS_DESCENDING) thrust::reverse(d_keys_wrapper, d_keys_wrapper + num_items); - thrust::sort(d_keys_wrapper, d_keys_wrapper + num_items); - if (IS_DESCENDING) thrust::reverse(d_keys_wrapper, d_keys_wrapper + num_items); + if (IS_DESCENDING) THRUST_NS_QUALIFIER::reverse(d_keys_wrapper, d_keys_wrapper + num_items); + THRUST_NS_QUALIFIER::sort(d_keys_wrapper, d_keys_wrapper + num_items); + if (IS_DESCENDING) THRUST_NS_QUALIFIER::reverse(d_keys_wrapper, d_keys_wrapper + num_items); } return cudaSuccess; @@ -459,19 +459,19 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_keys_wrapper(d_keys.Current()); - thrust::device_ptr d_values_wrapper(d_values.Current()); + THRUST_NS_QUALIFIER::device_ptr d_keys_wrapper(d_keys.Current()); + THRUST_NS_QUALIFIER::device_ptr d_values_wrapper(d_values.Current()); if (IS_DESCENDING) { - thrust::reverse(d_keys_wrapper, d_keys_wrapper + num_items); - thrust::reverse(d_values_wrapper, d_values_wrapper + num_items); + THRUST_NS_QUALIFIER::reverse(d_keys_wrapper, d_keys_wrapper + num_items); + THRUST_NS_QUALIFIER::reverse(d_values_wrapper, d_values_wrapper + num_items); } - thrust::sort_by_key(d_keys_wrapper, d_keys_wrapper + num_items, d_values_wrapper); + THRUST_NS_QUALIFIER::sort_by_key(d_keys_wrapper, d_keys_wrapper + num_items, d_values_wrapper); if (IS_DESCENDING) { - thrust::reverse(d_keys_wrapper, d_keys_wrapper + num_items); - thrust::reverse(d_values_wrapper, d_values_wrapper + num_items); + THRUST_NS_QUALIFIER::reverse(d_keys_wrapper, d_keys_wrapper + num_items); + THRUST_NS_QUALIFIER::reverse(d_values_wrapper, d_values_wrapper + num_items); } } diff --git a/test/test_device_reduce.cu b/test/test_device_reduce.cu index 7090e2f130..fb467249df 100644 --- a/test/test_device_reduce.cu +++ b/test/test_device_reduce.cu @@ -564,11 +564,11 @@ cudaError_t Dispatch( OutputT init; CubDebugExit(cudaMemcpy(&init, d_in + 0, sizeof(OutputT), cudaMemcpyDeviceToHost)); - thrust::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); OutputT retval; for (int i = 0; i < timing_iterations; ++i) { - retval = thrust::reduce(d_in_wrapper, d_in_wrapper + num_items, init, reduction_op); + retval = THRUST_NS_QUALIFIER::reduce(d_in_wrapper, d_in_wrapper + num_items, init, reduction_op); } if (!Equals >::VALUE) @@ -611,11 +611,11 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); OutputT retval; for (int i = 0; i < timing_iterations; ++i) { - retval = thrust::reduce(d_in_wrapper, d_in_wrapper + num_items); + retval = THRUST_NS_QUALIFIER::reduce(d_in_wrapper, d_in_wrapper + num_items); } if (!Equals >::VALUE) diff --git a/test/test_device_reduce_by_key.cu b/test/test_device_reduce_by_key.cu index a49a94bbf3..c08b228885 100644 --- a/test/test_device_reduce_by_key.cu +++ b/test/test_device_reduce_by_key.cu @@ -182,17 +182,17 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_keys_in_wrapper(d_keys_in); - thrust::device_ptr d_keys_out_wrapper(d_keys_out); + THRUST_NS_QUALIFIER::device_ptr d_keys_in_wrapper(d_keys_in); + THRUST_NS_QUALIFIER::device_ptr d_keys_out_wrapper(d_keys_out); - thrust::device_ptr d_values_in_wrapper(d_values_in); - thrust::device_ptr d_values_out_wrapper(d_values_out); + THRUST_NS_QUALIFIER::device_ptr d_values_in_wrapper(d_values_in); + THRUST_NS_QUALIFIER::device_ptr d_values_out_wrapper(d_values_out); - thrust::pair, thrust::device_ptr > d_out_ends; + THRUST_NS_QUALIFIER::pair, THRUST_NS_QUALIFIER::device_ptr > d_out_ends; for (int i = 0; i < timing_timing_iterations; ++i) { - d_out_ends = thrust::reduce_by_key( + d_out_ends = THRUST_NS_QUALIFIER::reduce_by_key( d_keys_in_wrapper, d_keys_in_wrapper + num_items, d_values_in_wrapper, diff --git a/test/test_device_run_length_encode.cu b/test/test_device_run_length_encode.cu index 6ae394a4c2..e90ee63089 100644 --- a/test/test_device_run_length_encode.cu +++ b/test/test_device_run_length_encode.cu @@ -231,19 +231,19 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_unique_out_wrapper(d_unique_out); - thrust::device_ptr d_lengths_out_wrapper(d_lengths_out); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_unique_out_wrapper(d_unique_out); + THRUST_NS_QUALIFIER::device_ptr d_lengths_out_wrapper(d_lengths_out); - thrust::pair, thrust::device_ptr > d_out_ends; + THRUST_NS_QUALIFIER::pair, THRUST_NS_QUALIFIER::device_ptr > d_out_ends; LengthT one_val; InitValue(INTEGER_SEED, one_val, 1); - thrust::constant_iterator constant_one(one_val); + THRUST_NS_QUALIFIER::constant_iterator constant_one(one_val); for (int i = 0; i < timing_timing_iterations; ++i) { - d_out_ends = thrust::reduce_by_key( + d_out_ends = THRUST_NS_QUALIFIER::reduce_by_key( d_in_wrapper, d_in_wrapper + num_items, constant_one, diff --git a/test/test_device_scan.cu b/test/test_device_scan.cu index 2a6feb42ae..0773a8d952 100644 --- a/test/test_device_scan.cu +++ b/test/test_device_scan.cu @@ -252,11 +252,11 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); for (int i = 0; i < timing_timing_iterations; ++i) { - thrust::exclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper, initial_value, scan_op); + THRUST_NS_QUALIFIER::exclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper, initial_value, scan_op); } } @@ -299,11 +299,11 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); for (int i = 0; i < timing_timing_iterations; ++i) { - thrust::exclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper); + THRUST_NS_QUALIFIER::exclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper); } } @@ -346,11 +346,11 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); for (int i = 0; i < timing_timing_iterations; ++i) { - thrust::inclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper, scan_op); + THRUST_NS_QUALIFIER::inclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper, scan_op); } } @@ -393,11 +393,11 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); for (int i = 0; i < timing_timing_iterations; ++i) { - thrust::inclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper); + THRUST_NS_QUALIFIER::inclusive_scan(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper); } } diff --git a/test/test_device_select_if.cu b/test/test_device_select_if.cu index ea12417868..fb77072487 100644 --- a/test/test_device_select_if.cu +++ b/test/test_device_select_if.cu @@ -264,13 +264,15 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_out_wrapper_end; - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper_end; + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); for (int i = 0; i < timing_timing_iterations; ++i) { - d_out_wrapper_end = thrust::copy_if(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper, select_op); + d_out_wrapper_end = THRUST_NS_QUALIFIER::copy_if(d_in_wrapper, + d_in_wrapper + num_items, + d_out_wrapper, select_op); } OffsetT num_selected = OffsetT(d_out_wrapper_end - d_out_wrapper); @@ -313,7 +315,7 @@ cudaError_t Dispatch( typename std::iterator_traits::value_type, // ... then the input iterator's value type, typename std::iterator_traits::value_type>::Type OutputT; // ... else the output iterator's value type - typedef thrust::reverse_iterator > ReverseOutputIteratorT; + typedef THRUST_NS_QUALIFIER::reverse_iterator > ReverseOutputIteratorT; if (d_temp_storage == 0) { @@ -321,16 +323,16 @@ cudaError_t Dispatch( } else { - thrust::pair, ReverseOutputIteratorT> d_out_wrapper_end; + THRUST_NS_QUALIFIER::pair, ReverseOutputIteratorT> d_out_wrapper_end; - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); ReverseOutputIteratorT d_out_unselected(d_out_wrapper + num_items); for (int i = 0; i < timing_timing_iterations; ++i) { - d_out_wrapper_end = thrust::partition_copy( + d_out_wrapper_end = THRUST_NS_QUALIFIER::partition_copy( d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper, @@ -387,14 +389,14 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_out_wrapper_end; - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); - thrust::device_ptr d_flags_wrapper(d_flags); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper_end; + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_flags_wrapper(d_flags); for (int i = 0; i < timing_timing_iterations; ++i) { - d_out_wrapper_end = thrust::copy_if(d_in_wrapper, d_in_wrapper + num_items, d_flags_wrapper, d_out_wrapper, CastOp()); + d_out_wrapper_end = THRUST_NS_QUALIFIER::copy_if(d_in_wrapper, d_in_wrapper + num_items, d_flags_wrapper, d_out_wrapper, CastOp()); } OffsetT num_selected = OffsetT(d_out_wrapper_end - d_out_wrapper); @@ -440,7 +442,7 @@ cudaError_t Dispatch( typename std::iterator_traits::value_type, // ... then the input iterator's value type, typename std::iterator_traits::value_type>::Type OutputT; // ... else the output iterator's value type - typedef thrust::reverse_iterator > ReverseOutputIteratorT; + typedef THRUST_NS_QUALIFIER::reverse_iterator > ReverseOutputIteratorT; if (d_temp_storage == 0) { @@ -448,16 +450,16 @@ cudaError_t Dispatch( } else { - thrust::pair, ReverseOutputIteratorT> d_out_wrapper_end; + THRUST_NS_QUALIFIER::pair, ReverseOutputIteratorT> d_out_wrapper_end; - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); - thrust::device_ptr d_flags_wrapper(d_flags); + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_flags_wrapper(d_flags); ReverseOutputIteratorT d_out_unselected(d_out_wrapper + num_items); for (int i = 0; i < timing_timing_iterations; ++i) { - d_out_wrapper_end = thrust::partition_copy( + d_out_wrapper_end = THRUST_NS_QUALIFIER::partition_copy( d_in_wrapper, d_in_wrapper + num_items, d_flags_wrapper, diff --git a/test/test_device_select_unique.cu b/test/test_device_select_unique.cu index 179974a069..a25fea8ead 100644 --- a/test/test_device_select_unique.cu +++ b/test/test_device_select_unique.cu @@ -143,12 +143,12 @@ cudaError_t Dispatch( } else { - thrust::device_ptr d_out_wrapper_end; - thrust::device_ptr d_in_wrapper(d_in); - thrust::device_ptr d_out_wrapper(d_out); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper_end; + THRUST_NS_QUALIFIER::device_ptr d_in_wrapper(d_in); + THRUST_NS_QUALIFIER::device_ptr d_out_wrapper(d_out); for (int i = 0; i < timing_timing_iterations; ++i) { - d_out_wrapper_end = thrust::unique_copy(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper); + d_out_wrapper_end = THRUST_NS_QUALIFIER::unique_copy(d_in_wrapper, d_in_wrapper + num_items, d_out_wrapper); } OffsetT num_selected = OffsetT(d_out_wrapper_end - d_out_wrapper); diff --git a/test/test_iterator.cu b/test/test_iterator.cu index 8a36a1a755..bd7efe126a 100644 --- a/test/test_iterator.cu +++ b/test/test_iterator.cu @@ -214,9 +214,9 @@ void TestConstant(T base) h_copy[i] = d_itr[i]; CubDebugExit(g_allocator.DeviceAllocate((void**)&d_copy, sizeof(T) * copy_items)); - thrust::device_ptr d_copy_wrapper(d_copy); + THRUST_NS_QUALIFIER::device_ptr d_copy_wrapper(d_copy); - thrust::copy_if(d_itr, d_itr + copy_items, d_copy_wrapper, SelectOp()); + THRUST_NS_QUALIFIER::copy_if(d_itr, d_itr + copy_items, d_copy_wrapper, SelectOp()); int compare = CompareDeviceResults(h_copy, d_copy, copy_items, g_verbose, g_verbose); printf("\tthrust::copy_if(): %s\n", (compare) ? "FAIL" : "PASS"); @@ -270,8 +270,8 @@ void TestCounting(T base) h_copy[i] = d_itr[i]; CubDebugExit(g_allocator.DeviceAllocate((void**)&d_copy, sizeof(T) * copy_items)); - thrust::device_ptr d_copy_wrapper(d_copy); - thrust::copy_if(d_itr, d_itr + copy_items, d_copy_wrapper, SelectOp()); + THRUST_NS_QUALIFIER::device_ptr d_copy_wrapper(d_copy); + THRUST_NS_QUALIFIER::copy_if(d_itr, d_itr + copy_items, d_copy_wrapper, SelectOp()); int compare = CompareDeviceResults(h_copy, d_copy, copy_items, g_verbose, g_verbose); printf("\tthrust::copy_if(): %s\n", (compare) ? "FAIL" : "PASS"); @@ -340,7 +340,7 @@ void TestModified() CacheModifiedInputIterator d_in_itr((CastT*) d_data); CacheModifiedOutputIterator d_out_itr((CastT*) d_copy); - thrust::copy_if(d_in_itr, d_in_itr + TEST_VALUES, d_out_itr, SelectOp()); + THRUST_NS_QUALIFIER::copy_if(d_in_itr, d_in_itr + TEST_VALUES, d_out_itr, SelectOp()); int compare = CompareDeviceResults(h_data, d_copy, TEST_VALUES, g_verbose, g_verbose); printf("\tthrust::copy_if(): %s\n", (compare) ? "FAIL" : "PASS"); @@ -409,9 +409,9 @@ void TestTransform() T *d_copy = NULL; CubDebugExit(g_allocator.DeviceAllocate((void**)&d_copy, sizeof(T) * TEST_VALUES)); - thrust::device_ptr d_copy_wrapper(d_copy); + THRUST_NS_QUALIFIER::device_ptr d_copy_wrapper(d_copy); - thrust::copy_if(d_itr, d_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); + THRUST_NS_QUALIFIER::copy_if(d_itr, d_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); int compare = CompareDeviceResults(h_copy, d_copy, TEST_VALUES, g_verbose, g_verbose); printf("\tthrust::copy_if(): %s\n", (compare) ? "FAIL" : "PASS"); @@ -484,10 +484,10 @@ void TestTexObj() T *d_copy = NULL; CubDebugExit(g_allocator.DeviceAllocate((void**)&d_copy, sizeof(T) * TEST_VALUES)); - thrust::device_ptr d_copy_wrapper(d_copy); + THRUST_NS_QUALIFIER::device_ptr d_copy_wrapper(d_copy); CubDebugExit(cudaMemset(d_copy, 0, sizeof(T) * TEST_VALUES)); - thrust::copy_if(d_obj_itr, d_obj_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); + THRUST_NS_QUALIFIER::copy_if(d_obj_itr, d_obj_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); int compare = CompareDeviceResults(h_data, d_copy, TEST_VALUES, g_verbose, g_verbose); printf("\tthrust::copy_if(): %s\n", (compare) ? "FAIL" : "PASS"); @@ -568,10 +568,10 @@ void TestTexRef() T *d_copy = NULL; CubDebugExit(g_allocator.DeviceAllocate((void**)&d_copy, sizeof(T) * TEST_VALUES)); - thrust::device_ptr d_copy_wrapper(d_copy); + THRUST_NS_QUALIFIER::device_ptr d_copy_wrapper(d_copy); CubDebugExit(cudaMemset(d_copy, 0, sizeof(T) * TEST_VALUES)); - thrust::copy_if(d_ref_itr, d_ref_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); + THRUST_NS_QUALIFIER::copy_if(d_ref_itr, d_ref_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); int compare = CompareDeviceResults(h_data, d_copy, TEST_VALUES, g_verbose, g_verbose); printf("\tthrust::copy_if(): %s\n", (compare) ? "FAIL" : "PASS"); @@ -651,9 +651,9 @@ void TestTexTransform() T *d_copy = NULL; CubDebugExit(g_allocator.DeviceAllocate((void**)&d_copy, sizeof(T) * TEST_VALUES)); - thrust::device_ptr d_copy_wrapper(d_copy); + THRUST_NS_QUALIFIER::device_ptr d_copy_wrapper(d_copy); - thrust::copy_if(xform_itr, xform_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); + THRUST_NS_QUALIFIER::copy_if(xform_itr, xform_itr + TEST_VALUES, d_copy_wrapper, SelectOp()); int compare = CompareDeviceResults(h_copy, d_copy, TEST_VALUES, g_verbose, g_verbose); printf("\tthrust::copy_if(): %s\n", (compare) ? "FAIL" : "PASS"); diff --git a/test/test_namespace_wrapped.cu b/test/test_namespace_wrapped.cu new file mode 100644 index 0000000000..58e16449c0 --- /dev/null +++ b/test/test_namespace_wrapped.cu @@ -0,0 +1,76 @@ +// Wrap thrust and cub in different enclosing namespaces +// (In practice, you probably want these to be the same, in which case just +// set THRUST_CUB_WRAPPED_NAMESPACE to set both). +#define THRUST_WRAPPED_NAMESPACE wrap_thrust +#define CUB_WRAPPED_NAMESPACE wrap_cub + +// Enable error checking: +#define CUB_STDERR + +#include +#include +#include + +#include +#include + +#include "test_util.h" + +#include +#include + +// Test that we can use a few common utilities and algorithms from wrapped +// Thrust/CUB namespaces at runtime. More extensive testing is performed by the +// header tests and the check_namespace.cmake test. +int main(int argc, char **argv) +{ + CommandLineArgs args(argc, argv); + CubDebugExit(args.DeviceInit()); + + const std::size_t n = 2048; + + // Fill a vector with random data: + ::wrap_thrust::thrust::host_vector h_input(n); + for (auto &val : h_input) + { + RandomBits(val); + } + + // Test the qualifier macro: + THRUST_NS_QUALIFIER::device_vector d_input(h_input); + THRUST_NS_QUALIFIER::device_vector d_output(n); + + std::size_t temp_storage_bytes{}; + + // Sort with DeviceRadixSort: + auto error = ::wrap_cub::cub::DeviceRadixSort::SortKeys( + nullptr, + temp_storage_bytes, + ::wrap_thrust::thrust::raw_pointer_cast(d_input.data()), + ::wrap_thrust::thrust::raw_pointer_cast(d_output.data()), + static_cast(n)); + + CubDebugExit(error); + + ::wrap_thrust::thrust::device_vector temp_storage( + temp_storage_bytes); + + // Test the CUB qualifier macro: + error = CUB_NS_QUALIFIER::DeviceRadixSort::SortKeys( + ::wrap_thrust::thrust::raw_pointer_cast(temp_storage.data()), + temp_storage_bytes, + ::wrap_thrust::thrust::raw_pointer_cast(d_input.data()), + ::wrap_thrust::thrust::raw_pointer_cast(d_output.data()), + static_cast(n)); + + CubDebugExit(error); + + // Verify output: + if (!::wrap_thrust::thrust::is_sorted(d_output.cbegin(), d_output.cend())) + { + std::cerr << "Output is not sorted!\n"; + return EXIT_FAILURE; + } + + return EXIT_SUCCESS; +} diff --git a/test/test_util.h b/test/test_util.h index ced27686db..fa856852ff 100644 --- a/test/test_util.h +++ b/test/test_util.h @@ -289,7 +289,7 @@ struct CommandLineArgs CubDebugExit(cudaMemGetInfo(&device_free_physmem, &device_total_physmem)); int ptx_version = 0; - error = CubDebug(cub::PtxVersion(ptx_version)); + error = CubDebug(CUB_NS_QUALIFIER::PtxVersion(ptx_version)); if (error) break; error = CubDebug(cudaGetDeviceProperties(&deviceProp, dev)); @@ -557,7 +557,7 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, T &value, i case RANDOM_MINUS_PLUS_ZERO: _CubLog("%s\n", "cub::InitValue cannot generate random numbers on device."); - cub::ThreadTrap(); + CUB_NS_QUALIFIER::ThreadTrap(); break; case UNIFORM: value = 2; @@ -587,7 +587,7 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, T &value, i case RANDOM_MINUS_PLUS_ZERO: { // Replace roughly 1/128 of values with -0.0 or +0.0, and generate the rest randomly - typedef typename cub::Traits::UnsignedBits UnsignedBits; + typedef typename CUB_NS_QUALIFIER::Traits::UnsignedBits UnsignedBits; char c; RandomBits(c); if (c == 0) @@ -638,7 +638,7 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, bool &value case RANDOM_MINUS_PLUS_ZERO: _CubLog("%s\n", "cub::InitValue cannot generate random numbers on device."); - cub::ThreadTrap(); + CUB_NS_QUALIFIER::ThreadTrap(); break; case UNIFORM: value = true; @@ -679,7 +679,7 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, bool &value * cub::NullType test initialization */ __host__ __device__ __forceinline__ void InitValue(GenMode /* gen_mode */, - cub::NullType &/* value */, + CUB_NS_QUALIFIER::NullType &/* value */, int /* index */ = 0) {} @@ -691,7 +691,7 @@ __host__ __device__ __forceinline__ void InitValue(GenMode /* gen_mode */, template __host__ __device__ __forceinline__ void InitValue( GenMode gen_mode, - cub::KeyValuePair& value, + CUB_NS_QUALIFIER::KeyValuePair& value, int index = 0) { InitValue(gen_mode, value.value, index); @@ -704,7 +704,7 @@ __host__ __device__ __forceinline__ void InitValue( #if CUB_INCLUDE_DEVICE_CODE _CubLog("%s\n", "cub::InitValue cannot generate random numbers on device."); - cub::ThreadTrap(); + CUB_NS_QUALIFIER::ThreadTrap(); #endif // CUB_INCLUDE_DEVICE_CODE } else @@ -727,7 +727,7 @@ __host__ __device__ __forceinline__ void InitValue( * KeyValuePair ostream operator */ template -std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair &val) +std::ostream& operator<<(std::ostream& os, const CUB_NS_QUALIFIER::KeyValuePair &val) { os << '(' << CoutCast(val.key) << ',' << CoutCast(val.value) << ')'; return os; @@ -791,7 +791,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & T retval = make_##T(a.x + b.x); \ return retval; \ } \ - namespace cub { \ + CUB_NAMESPACE_BEGIN \ template<> \ struct NumericTraits \ { \ @@ -813,7 +813,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & return retval; \ } \ }; \ - } /* namespace std */ + CUB_NAMESPACE_END @@ -879,7 +879,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & a.y + b.y); \ return retval; \ } \ - namespace cub { \ + CUB_NAMESPACE_BEGIN \ template<> \ struct NumericTraits \ { \ @@ -903,7 +903,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & return retval; \ } \ }; \ - } /* namespace cub */ + CUB_NAMESPACE_END @@ -976,7 +976,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & a.z + b.z); \ return retval; \ } \ - namespace cub { \ + CUB_NAMESPACE_BEGIN \ template<> \ struct NumericTraits \ { \ @@ -1002,7 +1002,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & return retval; \ } \ }; \ - } /* namespace cub */ + CUB_NAMESPACE_END /** @@ -1081,7 +1081,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & a.w + b.w); \ return retval; \ } \ - namespace cub { \ + CUB_NAMESPACE_BEGIN \ template<> \ struct NumericTraits \ { \ @@ -1109,7 +1109,7 @@ std::ostream& operator<<(std::ostream& os, const cub::KeyValuePair & return retval; \ } \ }; \ - } /* namespace cub */ + CUB_NAMESPACE_END /** * All vector overloads @@ -1233,7 +1233,7 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, TestFoo &va /// numeric_limits specialization -namespace cub { +CUB_NAMESPACE_BEGIN template<> struct NumericTraits { @@ -1260,7 +1260,7 @@ struct NumericTraits NumericTraits::Lowest()); } }; -} // namespace cub +CUB_NAMESPACE_END //--------------------------------------------------------------------- @@ -1349,7 +1349,7 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, TestBar &va } /// numeric_limits specialization -namespace cub { +CUB_NAMESPACE_BEGIN template<> struct NumericTraits { @@ -1372,7 +1372,7 @@ struct NumericTraits NumericTraits::Lowest()); } }; -} // namespace cub +CUB_NAMESPACE_END /****************************************************************************** @@ -1430,7 +1430,7 @@ int CompareResults(float* computed, float* reference, OffsetT len, bool verbose * Compares the equivalence of two arrays */ template -int CompareResults(cub::NullType* computed, cub::NullType* reference, OffsetT len, bool verbose = true) +int CompareResults(CUB_NS_QUALIFIER::NullType* computed, CUB_NS_QUALIFIER::NullType* reference, OffsetT len, bool verbose = true) { return 0; } @@ -1466,8 +1466,8 @@ int CompareResults(double* computed, double* reference, OffsetT len, bool verbos * of a host array */ int CompareDeviceResults( - cub::NullType */* h_reference */, - cub::NullType */* d_data */, + CUB_NS_QUALIFIER::NullType */* h_reference */, + CUB_NS_QUALIFIER::NullType */* d_data */, std::size_t /* num_items */, bool /* verbose */ = true, bool /* display_data */ = false) @@ -1482,7 +1482,7 @@ int CompareDeviceResults( template int CompareDeviceResults( S *h_reference, - cub::DiscardOutputIterator d_data, + CUB_NS_QUALIFIER::DiscardOutputIterator d_data, std::size_t num_items, bool verbose = true, bool display_data = false) @@ -1584,7 +1584,7 @@ int CompareDeviceDeviceResults( * Print the contents of a host array */ void DisplayResults( - cub::NullType */* h_data */, + CUB_NS_QUALIFIER::NullType */* h_data */, std::size_t /* num_items */) {} @@ -1643,7 +1643,7 @@ void InitializeSegments( if (num_segments <= 0) return; - unsigned int expected_segment_length = cub::DivideAndRoundUp(num_items, num_segments); + unsigned int expected_segment_length = CUB_NS_QUALIFIER::DivideAndRoundUp(num_items, num_segments); int offset = 0; for (int i = 0; i < num_segments; ++i) {