From c7d13db37ab80a8bdf86e24e5eb1321f1bcf9f59 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 19 Jun 2024 19:10:09 +0200 Subject: [PATCH] Drop facilities deprecated in a previous CUB major version (#1868) --- cub/cub/block/block_adjacent_difference.cuh | 456 -------------------- cub/cub/cub.cuh | 1 - cub/cub/iterator/tex_ref_input_iterator.cuh | 112 ----- cub/cub/util_type.cuh | 95 ---- cub/test/test_iterator_deprecated.cu | 294 ------------- 5 files changed, 958 deletions(-) delete mode 100644 cub/cub/iterator/tex_ref_input_iterator.cuh delete mode 100644 cub/test/test_iterator_deprecated.cu diff --git a/cub/cub/block/block_adjacent_difference.cuh b/cub/cub/block/block_adjacent_difference.cuh index 4a7fc7bda5..649f2a563f 100644 --- a/cub/cub/block/block_adjacent_difference.cuh +++ b/cub/cub/block/block_adjacent_difference.cuh @@ -959,462 +959,6 @@ public: } } } - - //! @} end member group - //! @name Head flag operations (deprecated) - //! @{ - -#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document - - /** - * @deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeads - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft instead. - * - * @param[out] output - * Calling thread's discontinuity head_flags - * - * @param[in] input - * Calling thread's input items - * - * @param[out] preds - * Calling thread's predecessor items - * - * @param[in] flag_op - * Binary boolean flag predicate - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void FlagHeads( - FlagT (&output)[ITEMS_PER_THREAD], T (&input)[ITEMS_PER_THREAD], T (&preds)[ITEMS_PER_THREAD], FlagOp flag_op) - { - // Share last item - temp_storage.last_items[linear_tid] = input[ITEMS_PER_THREAD - 1]; - - CTA_SYNC(); - - if (linear_tid == 0) - { - // Set flag for first thread-item (preds[0] is undefined) - output[0] = 1; - } - else - { - preds[0] = temp_storage.last_items[linear_tid - 1]; - output[0] = ApplyOp::FlagT(flag_op, preds[0], input[0], linear_tid * ITEMS_PER_THREAD); - } - - // Set output for remaining items - Iterate::FlagHeads(linear_tid, output, input, preds, flag_op); - } - - /** - * @deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeads - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft instead. - * - * @param[out] output - * Calling thread's discontinuity result - * - * @param[in] input - * Calling thread's input items - * - * @param[out] preds - * Calling thread's predecessor items - * - * @param[in] flag_op - * Binary boolean flag predicate - * - * @param[in] tile_predecessor_item - * [thread0 only] Item with which to compare the first tile item - * (input0 from thread0). - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void FlagHeads( - FlagT (&output)[ITEMS_PER_THREAD], - T (&input)[ITEMS_PER_THREAD], - T (&preds)[ITEMS_PER_THREAD], - FlagOp flag_op, - T tile_predecessor_item) - { - // Share last item - temp_storage.last_items[linear_tid] = input[ITEMS_PER_THREAD - 1]; - - CTA_SYNC(); - - // Set flag for first thread-item - preds[0] = (linear_tid == 0) ? tile_predecessor_item : // First thread - temp_storage.last_items[linear_tid - 1]; - - output[0] = ApplyOp::FlagT(flag_op, preds[0], input[0], linear_tid * ITEMS_PER_THREAD); - - // Set output for remaining items - Iterate::FlagHeads(linear_tid, output, input, preds, flag_op); - } - -#endif // DOXYGEN_SHOULD_SKIP_THIS - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeads - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft instead. - * - * @param[out] output - * Calling thread's discontinuity result - * - * @param[in] input - * Calling thread's input items - * - * @param[in] flag_op - * Binary boolean flag predicate - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void - FlagHeads(FlagT (&output)[ITEMS_PER_THREAD], T (&input)[ITEMS_PER_THREAD], FlagOp flag_op) - { - T preds[ITEMS_PER_THREAD]; - FlagHeads(output, input, preds, flag_op); - } - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeads - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft instead. - * - * @param[out] output - * Calling thread's discontinuity result - * - * @param[in] input - * Calling thread's input items - * - * @param[in] flag_op - * Binary boolean flag predicate - * - * @param[in] tile_predecessor_item - * [thread0 only] Item with which to compare the first tile item - * (input0 from thread0). - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void - FlagHeads(FlagT (&output)[ITEMS_PER_THREAD], T (&input)[ITEMS_PER_THREAD], FlagOp flag_op, T tile_predecessor_item) - { - T preds[ITEMS_PER_THREAD]; - FlagHeads(output, input, preds, flag_op, tile_predecessor_item); - } - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagTails - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractRight instead. - * - * @param output - * [out] Calling thread's discontinuity result - * - * @param input - * [in] Calling thread's input items - * - * @param flag_op - * [in] Binary boolean flag predicate - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void - FlagTails(FlagT (&output)[ITEMS_PER_THREAD], T (&input)[ITEMS_PER_THREAD], FlagOp flag_op) - { - // Share first item - temp_storage.first_items[linear_tid] = input[0]; - - CTA_SYNC(); - - // Set flag for last thread-item - output[ITEMS_PER_THREAD - 1] = - (linear_tid == BLOCK_THREADS - 1) ? 1 : // Last thread - ApplyOp::FlagT( - flag_op, - input[ITEMS_PER_THREAD - 1], - temp_storage.first_items[linear_tid + 1], - (linear_tid * ITEMS_PER_THREAD) + ITEMS_PER_THREAD); - - // Set output for remaining items - Iterate::FlagTails(linear_tid, output, input, flag_op); - } - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagTails - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractRight instead. - * - * @param[out] output - * Calling thread's discontinuity result - * - * @param[in] input - * Calling thread's input items - * - * @param[in] flag_op - * Binary boolean flag predicate - * - * @param[in] tile_successor_item - * [threadBLOCK_THREADS-1 only] Item with which to compare - * the last tile item (inputITEMS_PER_THREAD-1 from - * threadBLOCK_THREADS-1). - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void - FlagTails(FlagT (&output)[ITEMS_PER_THREAD], T (&input)[ITEMS_PER_THREAD], FlagOp flag_op, T tile_successor_item) - { - // Share first item - temp_storage.first_items[linear_tid] = input[0]; - - CTA_SYNC(); - - // Set flag for last thread-item - T successor_item = (linear_tid == BLOCK_THREADS - 1) ? tile_successor_item : // Last thread - temp_storage.first_items[linear_tid + 1]; - - output[ITEMS_PER_THREAD - 1] = ApplyOp::FlagT( - flag_op, input[ITEMS_PER_THREAD - 1], successor_item, (linear_tid * ITEMS_PER_THREAD) + ITEMS_PER_THREAD); - - // Set output for remaining items - Iterate::FlagTails(linear_tid, output, input, flag_op); - } - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeadsAndTails - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft or - * cub::BlockAdjacentDifference::SubtractRight instead. - * - * @param[out] head_flags - * Calling thread's discontinuity head_flags - * - * @param[out] tail_flags - * Calling thread's discontinuity tail_flags - * - * @param[in] input - * Calling thread's input items - * - * @param[in] flag_op - * Binary boolean flag predicate - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void FlagHeadsAndTails( - FlagT (&head_flags)[ITEMS_PER_THREAD], - FlagT (&tail_flags)[ITEMS_PER_THREAD], - T (&input)[ITEMS_PER_THREAD], - FlagOp flag_op) - { - // Share first and last items - temp_storage.first_items[linear_tid] = input[0]; - temp_storage.last_items[linear_tid] = input[ITEMS_PER_THREAD - 1]; - - CTA_SYNC(); - - T preds[ITEMS_PER_THREAD]; - - // Set flag for first thread-item - preds[0] = temp_storage.last_items[linear_tid - 1]; - if (linear_tid == 0) - { - head_flags[0] = 1; - } - else - { - head_flags[0] = ApplyOp::FlagT(flag_op, preds[0], input[0], linear_tid * ITEMS_PER_THREAD); - } - - // Set flag for last thread-item - tail_flags[ITEMS_PER_THREAD - 1] = - (linear_tid == BLOCK_THREADS - 1) ? 1 : // Last thread - ApplyOp::FlagT( - flag_op, - input[ITEMS_PER_THREAD - 1], - temp_storage.first_items[linear_tid + 1], - (linear_tid * ITEMS_PER_THREAD) + ITEMS_PER_THREAD); - - // Set head_flags for remaining items - Iterate::FlagHeads(linear_tid, head_flags, input, preds, flag_op); - - // Set tail_flags for remaining items - Iterate::FlagTails(linear_tid, tail_flags, input, flag_op); - } - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeadsAndTails - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft or - * cub::BlockAdjacentDifference::SubtractRight instead. - * - * @param[out] head_flags - * Calling thread's discontinuity head_flags - * - * @param[out] tail_flags - * Calling thread's discontinuity tail_flags - * - * @param[in] tile_successor_item - * [threadBLOCK_THREADS-1 only] Item with which to compare - * the last tile item (inputITEMS_PER_THREAD-1 from - * threadBLOCK_THREADS-1). - * - * @param[in] input - * Calling thread's input items - * - * @param[in] flag_op - * Binary boolean flag predicate - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void FlagHeadsAndTails( - FlagT (&head_flags)[ITEMS_PER_THREAD], - FlagT (&tail_flags)[ITEMS_PER_THREAD], - T tile_successor_item, - T (&input)[ITEMS_PER_THREAD], - FlagOp flag_op) - { - // Share first and last items - temp_storage.first_items[linear_tid] = input[0]; - temp_storage.last_items[linear_tid] = input[ITEMS_PER_THREAD - 1]; - - CTA_SYNC(); - - T preds[ITEMS_PER_THREAD]; - - // Set flag for first thread-item - if (linear_tid == 0) - { - head_flags[0] = 1; - } - else - { - preds[0] = temp_storage.last_items[linear_tid - 1]; - head_flags[0] = ApplyOp::FlagT(flag_op, preds[0], input[0], linear_tid * ITEMS_PER_THREAD); - } - - // Set flag for last thread-item - T successor_item = (linear_tid == BLOCK_THREADS - 1) ? tile_successor_item : // Last thread - temp_storage.first_items[linear_tid + 1]; - - tail_flags[ITEMS_PER_THREAD - 1] = ApplyOp::FlagT( - flag_op, input[ITEMS_PER_THREAD - 1], successor_item, (linear_tid * ITEMS_PER_THREAD) + ITEMS_PER_THREAD); - - // Set head_flags for remaining items - Iterate::FlagHeads(linear_tid, head_flags, input, preds, flag_op); - - // Set tail_flags for remaining items - Iterate::FlagTails(linear_tid, tail_flags, input, flag_op); - } - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeadsAndTails - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft or - * cub::BlockAdjacentDifference::SubtractRight instead. - * - * @param[out] head_flags - * Calling thread's discontinuity head_flags - * - * @param[in] tile_predecessor_item - * [thread0 only] Item with which to compare the first tile item - * (input0 from thread0). - * - * @param[out] tail_flags - * Calling thread's discontinuity tail_flags - * - * @param[in] input - * Calling thread's input items - * - * @param[in] flag_op - * Binary boolean flag predicate - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void FlagHeadsAndTails( - FlagT (&head_flags)[ITEMS_PER_THREAD], - T tile_predecessor_item, - FlagT (&tail_flags)[ITEMS_PER_THREAD], - T (&input)[ITEMS_PER_THREAD], - FlagOp flag_op) - { - // Share first and last items - temp_storage.first_items[linear_tid] = input[0]; - temp_storage.last_items[linear_tid] = input[ITEMS_PER_THREAD - 1]; - - CTA_SYNC(); - - T preds[ITEMS_PER_THREAD]; - - // Set flag for first thread-item - preds[0] = (linear_tid == 0) ? tile_predecessor_item : // First thread - temp_storage.last_items[linear_tid - 1]; - - head_flags[0] = ApplyOp::FlagT(flag_op, preds[0], input[0], linear_tid * ITEMS_PER_THREAD); - - // Set flag for last thread-item - tail_flags[ITEMS_PER_THREAD - 1] = - (linear_tid == BLOCK_THREADS - 1) ? 1 : // Last thread - ApplyOp::FlagT( - flag_op, - input[ITEMS_PER_THREAD - 1], - temp_storage.first_items[linear_tid + 1], - (linear_tid * ITEMS_PER_THREAD) + ITEMS_PER_THREAD); - - // Set head_flags for remaining items - Iterate::FlagHeads(linear_tid, head_flags, input, preds, flag_op); - - // Set tail_flags for remaining items - Iterate::FlagTails(linear_tid, tail_flags, input, flag_op); - } - - /** - * deprecated [Since 1.14.0] The cub::BlockAdjacentDifference::FlagHeadsAndTails - * APIs are deprecated. Use cub::BlockAdjacentDifference::SubtractLeft or - * cub::BlockAdjacentDifference::SubtractRight instead. - * - * @param head_flags - * [out] Calling thread's discontinuity head_flags - * - * @param tile_predecessor_item - * [in] [thread0 only] Item with which to compare the first tile - * item (input0 from thread0). - * - * @param tail_flags - * [out] Calling thread's discontinuity tail_flags - * - * @param tile_successor_item - * [in] [threadBLOCK_THREADS-1 only] Item with which to - * compare the last tile item (inputITEMS_PER_THREAD-1 from - * threadBLOCK_THREADS-1). - * - * @param input - * [in] Calling thread's input items - * - * @param flag_op - * [in] Binary boolean flag predicate - */ - template - CUB_DEPRECATED _CCCL_DEVICE _CCCL_FORCEINLINE void FlagHeadsAndTails( - FlagT (&head_flags)[ITEMS_PER_THREAD], - T tile_predecessor_item, - FlagT (&tail_flags)[ITEMS_PER_THREAD], - T tile_successor_item, - T (&input)[ITEMS_PER_THREAD], - FlagOp flag_op) - { - // Share first and last items - temp_storage.first_items[linear_tid] = input[0]; - temp_storage.last_items[linear_tid] = input[ITEMS_PER_THREAD - 1]; - - CTA_SYNC(); - - T preds[ITEMS_PER_THREAD]; - - // Set flag for first thread-item - preds[0] = (linear_tid == 0) ? tile_predecessor_item : // First thread - temp_storage.last_items[linear_tid - 1]; - - head_flags[0] = ApplyOp::FlagT(flag_op, preds[0], input[0], linear_tid * ITEMS_PER_THREAD); - - // Set flag for last thread-item - T successor_item = (linear_tid == BLOCK_THREADS - 1) ? tile_successor_item : // Last thread - temp_storage.first_items[linear_tid + 1]; - - tail_flags[ITEMS_PER_THREAD - 1] = ApplyOp::FlagT( - flag_op, input[ITEMS_PER_THREAD - 1], successor_item, (linear_tid * ITEMS_PER_THREAD) + ITEMS_PER_THREAD); - - // Set head_flags for remaining items - Iterate::FlagHeads(linear_tid, head_flags, input, preds, flag_op); - - // Set tail_flags for remaining items - Iterate::FlagTails(linear_tid, tail_flags, input, flag_op); - } - - //! @} end member group }; CUB_NAMESPACE_END diff --git a/cub/cub/cub.cuh b/cub/cub/cub.cuh index db05aabb49..ea199e7685 100644 --- a/cub/cub/cub.cuh +++ b/cub/cub/cub.cuh @@ -105,7 +105,6 @@ #include #include #include -#include #include // Util diff --git a/cub/cub/iterator/tex_ref_input_iterator.cuh b/cub/cub/iterator/tex_ref_input_iterator.cuh deleted file mode 100644 index e65f3eea18..0000000000 --- a/cub/cub/iterator/tex_ref_input_iterator.cuh +++ /dev/null @@ -1,112 +0,0 @@ -/****************************************************************************** - * Copyright (c) 2011, Duane Merrill. All rights reserved. - * Copyright (c) 2011-2018, 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: - * * Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * * Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution. - * * Neither the name of the NVIDIA CORPORATION nor the - * names of its contributors may be used to endorse or promote products - * derived from this software without specific prior written permission. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND - * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED - * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE - * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY - * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES - * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; - * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND - * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS - * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - * - ******************************************************************************/ - -/** - * @file - * Random-access iterator types - */ - -#pragma once - -#include - -#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) -# pragma GCC system_header -#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) -# pragma clang system_header -#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) -# pragma system_header -#endif // no system header - -#include - -#include - -CUB_NAMESPACE_BEGIN - -/** - * @brief A random-access input wrapper for dereferencing array values through texture cache. - * - * deprecated [Since 1.13.0] The CUDA texture management APIs used by - * TexRefInputIterator are deprecated. Use cub::TexObjInputIterator instead. - * - * @par Overview - * - TexRefInputIterator wraps a native device pointer of type ValueType*. References - * to elements are to be loaded through texture cache. - * - Can be used to load any data type from memory through texture cache. - * - Can be manipulated and exchanged within and between host and device - * functions, can only be constructed within host functions, and can only be - * dereferenced within device functions. - * - The @p UNIQUE_ID template parameter is used to statically name the underlying texture - * reference. Only one TexRefInputIterator instance can be bound at any given time for a - * specific combination of (1) data type @p T, (2) @p UNIQUE_ID, (3) host - * thread, and (4) compilation .o unit. - * - With regard to nested/dynamic parallelism, TexRefInputIterator iterators may only be - * created by the host thread and used by a top-level kernel (i.e. the one which is launched - * from the host). - * - Compatible with Thrust API v1.7 or newer. - * - * @par Snippet - * The code snippet below illustrates the use of @p TexRefInputIterator to - * dereference a device array of doubles through texture cache. - * @par - * @code - * #include // or equivalently - * - * // Declare, allocate, and initialize a device array - * int num_items; // e.g., 7 - * double *d_in; // e.g., [8.0, 6.0, 7.0, 5.0, 3.0, 0.0, 9.0] - * - * // Create an iterator wrapper - * cub::TexRefInputIterator itr; - * itr.BindTexture(d_in, sizeof(double) * num_items); - * ... - * - * // Within device code: - * printf("%f\n", itr[0]); // 8.0 - * printf("%f\n", itr[1]); // 6.0 - * printf("%f\n", itr[6]); // 9.0 - * - * ... - * itr.UnbindTexture(); - * - * @endcode - * - * @tparam T - * The value type of this iterator - * - * @tparam UNIQUE_ID - * A globally-unique identifier (within the compilation unit) to name the underlying texture reference - * - * @tparam OffsetT - * The difference type of this iterator (Default: @p ptrdiff_t) - */ -template -using TexRefInputIterator CUB_DEPRECATED = cub::TexObjInputIterator; - -CUB_NAMESPACE_END diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh index c4145c4801..8afdf18e03 100644 --- a/cub/cub/util_type.cuh +++ b/cub/cub/util_type.cuh @@ -126,35 +126,6 @@ template using non_void_value_t = typename non_void_value_impl::type; } // namespace detail -/** - * \brief Type selection (IF ? ThenType : ElseType) - * - * \deprecated [Since 1.16.0] The cub::If APIs are deprecated. - * Use cub::detail::conditional_t instead. - */ -template -struct CUB_DEPRECATED If -{ - using Type = cub::detail::conditional_t; -}; - -/****************************************************************************** - * Type equality - ******************************************************************************/ - -/** - * \brief Type equality test - * - * \deprecated [Since 1.16.0] The cub::Equals APIs are deprecated. - * Use std::is_same instead. - */ -template -struct CUB_DEPRECATED Equals -{ - static constexpr int VALUE = ::cuda::std::is_same::value ? 1 : 0; - static constexpr int NEGATE = VALUE ? 0 : 1; -}; - /****************************************************************************** * Static math ******************************************************************************/ @@ -203,56 +174,6 @@ struct PowerOfTwo }; }; -/****************************************************************************** - * Pointer vs. iterator detection - ******************************************************************************/ - -/** - * \brief Pointer vs. iterator - * - * \deprecated [Since 1.16.0] The cub::IsPointer APIs are deprecated. - * Use std::is_pointer instead. - */ -template -struct CUB_DEPRECATED IsPointer -{ - static constexpr int VALUE = ::cuda::std::is_pointer::value; -}; - -/****************************************************************************** - * Qualifier detection - ******************************************************************************/ - -/** - * \brief Volatile modifier test - * - * \deprecated [Since 1.16.0] The cub::IsVolatile APIs are deprecated. - * Use std::is_volatile instead. - */ -template -struct CUB_DEPRECATED IsVolatile -{ - static constexpr int VALUE = ::cuda::std::is_volatile::value; -}; - -/****************************************************************************** - * Qualifier removal - ******************************************************************************/ - -/** - * \brief Removes \p const and \p volatile qualifiers from type \p Tp. - * - * \deprecated [Since 1.16.0] The cub::RemoveQualifiers APIs are deprecated. - * Use std::remove_cv instead. - * - * For example: - * typename RemoveQualifiers::Type // int; - */ -template -struct CUB_DEPRECATED RemoveQualifiers -{ - using Type = typename ::cuda::std::remove_cv::type; -}; #endif // DOXYGEN_SHOULD_SKIP_THIS /****************************************************************************** @@ -956,22 +877,6 @@ struct DoubleBuffer }; \ }; -/****************************************************************************** - * Simple enable-if (similar to Boost) - ******************************************************************************/ - -/** - * \brief Simple enable-if (similar to Boost) - * - * \deprecated [Since 1.16.0] The cub::If APIs are deprecated. - * Use std::enable_if instead. - */ -template -struct CUB_DEPRECATED EnableIf -{ - using Type = typename ::cuda::std::enable_if::type; -}; - /****************************************************************************** * Typedef-detection ******************************************************************************/ diff --git a/cub/test/test_iterator_deprecated.cu b/cub/test/test_iterator_deprecated.cu deleted file mode 100644 index f6a8e3814a..0000000000 --- a/cub/test/test_iterator_deprecated.cu +++ /dev/null @@ -1,294 +0,0 @@ -/****************************************************************************** - * Copyright (c) 2011, Duane Merrill. All rights reserved. - * Copyright (c) 2011-2018, 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: - * * Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * * Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution. - * * Neither the name of the NVIDIA CORPORATION nor the - * names of its contributors may be used to endorse or promote products - * derived from this software without specific prior written permission. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND - * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED - * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE - * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY - * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES - * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; - * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND - * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS - * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - * - ******************************************************************************/ - -/****************************************************************************** - * Test of iterator utilities - ******************************************************************************/ - -// Ensure printing of CUDA runtime errors to console -#define CUB_STDERR - -// This file tests deprecated CUB APIs. Silence deprecation warnings: -#define CUB_IGNORE_DEPRECATED_API - -#include -#include -#include - -#include -#include -#include - -#include "test_util.h" - -using namespace cub; - -//--------------------------------------------------------------------- -// Globals, constants and typedefs -//--------------------------------------------------------------------- - -bool g_verbose = false; -CachingDeviceAllocator g_allocator(true); - -//--------------------------------------------------------------------- -// Test kernels -//--------------------------------------------------------------------- - -/** - * Test random access input iterator - */ -template -__global__ void Kernel(InputIteratorT d_in, T* d_out, InputIteratorT* d_itrs) -{ - d_out[0] = *d_in; // Value at offset 0 - d_out[1] = d_in[100]; // Value at offset 100 - d_out[2] = *(d_in + 1000); // Value at offset 1000 - d_out[3] = *(d_in + 10000); // Value at offset 10000 - - d_in++; - d_out[4] = d_in[0]; // Value at offset 1 - - d_in += 20; - d_out[5] = d_in[0]; // Value at offset 21 - d_itrs[0] = d_in; // Iterator at offset 21 - - d_in -= 10; - d_out[6] = d_in[0]; // Value at offset 11; - - d_in -= 11; - d_out[7] = d_in[0]; // Value at offset 0 - d_itrs[1] = d_in; // Iterator at offset 0 -} - -//--------------------------------------------------------------------- -// Host testing subroutines -//--------------------------------------------------------------------- - -/** - * Run iterator test on device - */ -template -void Test(InputIteratorT d_in, T (&h_reference)[TEST_VALUES]) -{ - // Allocate device arrays - T* d_out = nullptr; - InputIteratorT* d_itrs = nullptr; - CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(T) * TEST_VALUES)); - CubDebugExit(g_allocator.DeviceAllocate((void**) &d_itrs, sizeof(InputIteratorT) * 2)); - - int compare; - - // Run unguarded kernel - Kernel<<<1, 1>>>(d_in, d_out, d_itrs); - - CubDebugExit(cudaPeekAtLastError()); - CubDebugExit(cudaDeviceSynchronize()); - - // Check results - compare = CompareDeviceResults(h_reference, d_out, TEST_VALUES, g_verbose, g_verbose); - printf("\tValues: %s\n", (compare) ? "FAIL" : "PASS"); - AssertEquals(0, compare); - - // Check iterator at offset 21 - InputIteratorT h_itr = d_in + 21; - compare = CompareDeviceResults(&h_itr, d_itrs, 1, g_verbose, g_verbose); - printf("\tIterators: %s\n", (compare) ? "FAIL" : "PASS"); - AssertEquals(0, compare); - - // Check iterator at offset 0 - compare = CompareDeviceResults(&d_in, d_itrs + 1, 1, g_verbose, g_verbose); - printf("\tIterators: %s\n", (compare) ? "FAIL" : "PASS"); - AssertEquals(0, compare); - - // Cleanup - if (d_out) - { - CubDebugExit(g_allocator.DeviceFree(d_out)); - } - if (d_itrs) - { - CubDebugExit(g_allocator.DeviceFree(d_itrs)); - } -} - -/** - * Test tex-ref texture iterator - */ -template -void TestTexRef() -{ - printf("\nTesting tex-ref iterator on type %s\n", typeid(T).name()); - fflush(stdout); - - // - // Test iterator manipulation in kernel - // - - constexpr int TEST_VALUES = 11000; - constexpr unsigned int DUMMY_OFFSET = 500; - constexpr unsigned int DUMMY_TEST_VALUES = TEST_VALUES - DUMMY_OFFSET; - - T* h_data = new T[TEST_VALUES]; - for (int i = 0; i < TEST_VALUES; ++i) - { - RandomBits(h_data[i]); - } - - // Allocate device arrays - T* d_data = nullptr; - T* d_dummy = nullptr; - CubDebugExit(g_allocator.DeviceAllocate((void**) &d_data, sizeof(T) * TEST_VALUES)); - CubDebugExit(cudaMemcpy(d_data, h_data, sizeof(T) * TEST_VALUES, cudaMemcpyHostToDevice)); - - CubDebugExit(g_allocator.DeviceAllocate((void**) &d_dummy, sizeof(T) * DUMMY_TEST_VALUES)); - CubDebugExit(cudaMemcpy(d_dummy, h_data + DUMMY_OFFSET, sizeof(T) * DUMMY_TEST_VALUES, cudaMemcpyHostToDevice)); - - // Initialize reference data - T h_reference[8]; - h_reference[0] = h_data[0]; // Value at offset 0 - h_reference[1] = h_data[100]; // Value at offset 100 - h_reference[2] = h_data[1000]; // Value at offset 1000 - h_reference[3] = h_data[10000]; // Value at offset 10000 - h_reference[4] = h_data[1]; // Value at offset 1 - h_reference[5] = h_data[21]; // Value at offset 21 - h_reference[6] = h_data[11]; // Value at offset 11 - h_reference[7] = h_data[0]; // Value at offset 0; - - // Create and bind ref-based test iterator - TexRefInputIterator d_ref_itr; - CubDebugExit(d_ref_itr.BindTexture((CastT*) d_data, sizeof(T) * TEST_VALUES)); - - // Create and bind dummy iterator of same type to check with interferance - TexRefInputIterator d_ref_itr2; - CubDebugExit(d_ref_itr2.BindTexture((CastT*) d_dummy, sizeof(T) * DUMMY_TEST_VALUES)); - - Test(d_ref_itr, h_reference); - - CubDebugExit(d_ref_itr.UnbindTexture()); - CubDebugExit(d_ref_itr2.UnbindTexture()); - - if (h_data) - { - delete[] h_data; - } - if (d_data) - { - CubDebugExit(g_allocator.DeviceFree(d_data)); - } - if (d_dummy) - { - CubDebugExit(g_allocator.DeviceFree(d_dummy)); - } -} - -/** - * Run non-integer tests - */ -template -void Test() -{ - TestTexRef(); -} - -/** - * Run tests - */ -template -void Test() -{ - // Test non-const type - Test(); - - // Test non-const type - Test(); -} - -/** - * Main - */ -int main(int argc, char** argv) -{ - // Initialize command line - CommandLineArgs args(argc, argv); - g_verbose = args.CheckCmdLineFlag("v"); - - // Print usage - if (args.CheckCmdLineFlag("help")) - { - printf("%s " - "[--device=] " - "[--v] " - "\n", - argv[0]); - exit(0); - } - - // Initialize device - CubDebugExit(args.DeviceInit()); - - // Evaluate different data types - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - Test(); - - Test(); - Test(); - - printf("\nTest complete\n"); - fflush(stdout); - - return 0; -}