diff --git a/CMakeLists.txt b/CMakeLists.txt index 020ebfa916..1d0cf765e8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -60,9 +60,19 @@ option(CUB_ENABLE_CPP_DIALECT_IN_NAMES "Include C++ dialect information in target/object/etc names." ON ) - mark_as_advanced(CUB_ENABLE_CPP_DIALECT_IN_NAMES) +# By default, suppress deprecation warnings when building our test suite, +# since we'll need to test deprecated APIs with `-Werror`. This option is only +# used when CUB is built stand-alone; otherwise the Thrust option has the same +# effect. +if (NOT CUB_IN_THRUST) + option(CUB_IGNORE_DEPRECATED_API + "Suppress warnings about deprecated Thrust/CUB API." + ON + ) +endif() + # Check if we're actually building anything before continuing. If not, no need # to search for deps, etc. This is a common approach for packagers that just # need the install rules. See GH issue NVIDIA/thrust#1211. diff --git a/cub/agent/agent_spmv_orig.cuh b/cub/agent/agent_spmv_orig.cuh index 810f893fbe..e377797a27 100644 --- a/cub/agent/agent_spmv_orig.cuh +++ b/cub/agent/agent_spmv_orig.cuh @@ -44,7 +44,7 @@ #include "../thread/thread_operators.cuh" #include "../iterator/cache_modified_input_iterator.cuh" #include "../iterator/counting_input_iterator.cuh" -#include "../iterator/tex_ref_input_iterator.cuh" +#include "../iterator/tex_obj_input_iterator.cuh" /// Optional outer namespace(s) CUB_NS_PREFIX @@ -109,7 +109,7 @@ struct SpmvParams ValueT alpha; ///< Alpha multiplicand ValueT beta; ///< Beta addend-multiplicand - TexRefInputIterator t_vector_x; + TexObjInputIterator t_vector_x; }; diff --git a/cub/cmake/cub-config.cmake b/cub/cmake/cub-config.cmake index 1512a564f0..3d7f64bd4d 100644 --- a/cub/cmake/cub-config.cmake +++ b/cub/cmake/cub-config.cmake @@ -38,6 +38,10 @@ set(_CUB_INCLUDE_DIR "${_CUB_VERSION_INCLUDE_DIR}" unset(_CUB_VERSION_INCLUDE_DIR CACHE) # Clear tmp variable from cache target_include_directories(_CUB_CUB INTERFACE "${_CUB_INCLUDE_DIR}") +if (CUB_IGNORE_DEPRECATED_API OR THRUST_IGNORE_DEPRECATED_API) + target_compile_definitions(_CUB_CUB INTERFACE "CUB_IGNORE_DEPRECATED_API") +endif() + if (CUB_IGNORE_DEPRECATED_CPP_DIALECT OR THRUST_IGNORE_DEPRECATED_CPP_DIALECT) target_compile_definitions(_CUB_CUB INTERFACE "CUB_IGNORE_DEPRECATED_CPP_DIALECT") diff --git a/cub/iterator/tex_obj_input_iterator.cuh b/cub/iterator/tex_obj_input_iterator.cuh index 2bd3a607e1..4c286813c3 100644 --- a/cub/iterator/tex_obj_input_iterator.cuh +++ b/cub/iterator/tex_obj_input_iterator.cuh @@ -77,7 +77,7 @@ namespace cub { * - Compatible with Thrust API v1.7 or newer. * * \par Snippet - * The code snippet below illustrates the use of \p TexRefInputIterator to + * The code snippet below illustrates the use of \p TexObjInputIterator to * dereference a device array of doubles through texture cache. * \par * \code @@ -208,6 +208,11 @@ public: #if CUB_INCLUDE_HOST_CODE // Simply dereference the pointer on the host return ptr[tex_offset]; + #else + // Never executed, just need a return value for this codepath. + // The `reference` type is actually just T, so we can fake this + // easily. + return reference{}; #endif } else { #if CUB_INCLUDE_DEVICE_CODE diff --git a/cub/iterator/tex_ref_input_iterator.cuh b/cub/iterator/tex_ref_input_iterator.cuh index e4bb7dbb74..8f23928ca4 100644 --- a/cub/iterator/tex_ref_input_iterator.cuh +++ b/cub/iterator/tex_ref_input_iterator.cuh @@ -42,18 +42,6 @@ #include "../util_debug.cuh" #include "../config.cuh" -// This class needs to go through a deprecation cycle and be removed, as the -// underlying cudaBindTexture / cudaUnbindTexture APIs are deprecated. -// See issue NVIDIA/cub#191. -// Turn off deprecation warnings when compiling this file until then. -#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC -#pragma warning(disable:4996) -#elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_GCC || \ - CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG -#pragma GCC diagnostic push -#pragma GCC diagnostic ignored "-Wdeprecated-declarations" -#endif - #if (CUDART_VERSION >= 5050) || defined(DOXYGEN_ACTIVE) // This iterator is compatible with CUDA 5.5 and newer #if (THRUST_VERSION >= 100700) // This iterator is compatible with Thrust API 1.7 and newer @@ -79,8 +67,21 @@ namespace { /// Global texture reference specialized by type template -struct IteratorTexRef +struct CUB_DEPRECATED IteratorTexRef { + +// This class uses the deprecated cudaBindTexture / cudaUnbindTexture APIs. +// See issue NVIDIA/cub#191. +// Turn off deprecation warnings when compiling class implementation in favor +// of deprecating TexRefInputIterator instead. +#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC +#pragma warning(disable:4996) +#elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_GCC || \ + CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wdeprecated-declarations" +#endif + /// And by unique ID template struct TexId @@ -143,6 +144,13 @@ template template typename IteratorTexRef::template TexId::TexRef IteratorTexRef::template TexId::ref = 0; +// Re-enable deprecation warnings: +#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC +#pragma warning(default:4996) +#elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_GCC || \ + CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG +#pragma GCC diagnostic pop +#endif } // Anonymous namespace @@ -161,6 +169,9 @@ typename IteratorTexRef::template TexId::TexRef IteratorTexRef: /** * \brief A random-access input wrapper for dereferencing array values through texture cache. Uses older Tesla/Fermi-style texture references. * + * \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. @@ -176,7 +187,6 @@ typename IteratorTexRef::template TexId::TexRef IteratorTexRef: * 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. - * - Compatible with CUDA toolkit v5.5 or newer. * * \par Snippet * The code snippet below illustrates the use of \p TexRefInputIterator to @@ -212,8 +222,21 @@ template < typename T, int UNIQUE_ID, typename OffsetT = ptrdiff_t> -class TexRefInputIterator +class CUB_DEPRECATED TexRefInputIterator { + +// This class uses the deprecated cudaBindTexture / cudaUnbindTexture APIs. +// See issue NVIDIA/cub#191. +// Turn off deprecation warnings when compiling class implementation in favor +// of deprecating TexRefInputIterator instead. +#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC +#pragma warning(disable:4996) +#elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_GCC || \ + CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wdeprecated-declarations" +#endif + public: // Required iterator traits @@ -379,6 +402,14 @@ public: return os; } +// Re-enable deprecation warnings: +#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC +#pragma warning(default:4996) +#elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_GCC || \ + CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG +#pragma GCC diagnostic pop +#endif + }; @@ -389,11 +420,3 @@ public: CUB_NS_POSTFIX // Optional outer namespace(s) #endif // CUDART_VERSION - -// Re-enable deprecation warnings: -#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC -#pragma warning(default:4996) -#elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_GCC || \ - CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG -#pragma GCC diagnostic pop -#endif diff --git a/cub/util_deprecated.cuh b/cub/util_deprecated.cuh index b2bf4658b5..6819f731f6 100644 --- a/cub/util_deprecated.cuh +++ b/cub/util_deprecated.cuh @@ -33,8 +33,17 @@ #pragma once #include "util_compiler.cuh" +#include "util_cpp_dialect.cuh" -#if CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC +#if defined(THRUST_IGNORE_DEPRECATED_API) && !defined(CUB_IGNORE_DEPRECATED_API) +# define CUB_IGNORE_DEPRECATED_API +#endif + +#ifdef CUB_IGNORE_DEPRECATED_API +# define CUB_DEPRECATED +#elif CUB_CPP_DIALECT >= 2014 +# define CUB_DEPRECATED [[deprecated]] +#elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_MSVC # define CUB_DEPRECATED __declspec(deprecated) #elif CUB_HOST_COMPILER == CUB_HOST_COMPILER_CLANG # define CUB_DEPRECATED __attribute__((deprecated)) @@ -43,4 +52,3 @@ #else # define CUB_DEPRECATED #endif - diff --git a/test/test_iterator.cu b/test/test_iterator.cu index 65e107a37d..8a36a1a755 100644 --- a/test/test_iterator.cu +++ b/test/test_iterator.cu @@ -629,7 +629,7 @@ void TestTexTransform() h_reference[7] = op(h_data[0]); // Value at offset 0; // Create and bind texture iterator - typedef TexRefInputIterator TextureIterator; + typedef TexObjInputIterator TextureIterator; TextureIterator d_tex_itr; CubDebugExit(d_tex_itr.BindTexture((CastT*) d_data, sizeof(T) * TEST_VALUES));