Merge pull request #320 from allisonvacanti/191_deprecate_texref

Deprecate TexRefInputIterator.

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.

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<ValueT, 66778899, OffsetT>  t_vector_x;
+    TexObjInputIterator<ValueT, OffsetT>  t_vector_x;
 };
 
 

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")

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

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 <typename T>
-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 <int UNIQUE_ID>
     struct TexId
@@ -143,6 +144,13 @@ template <typename  T>
 template <int       UNIQUE_ID>
 typename IteratorTexRef<T>::template TexId<UNIQUE_ID>::TexRef IteratorTexRef<T>::template TexId<UNIQUE_ID>::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<T>::template TexId<UNIQUE_ID>::TexRef IteratorTexRef<T>:
 /**
  * \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 <tt>ValueType*</tt>. References
  *   to elements are to be loaded through texture cache.
@@ -176,7 +187,6 @@ typename IteratorTexRef<T>::template TexId<UNIQUE_ID>::TexRef IteratorTexRef<T>:
  *   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

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
-

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<T, __LINE__> TextureIterator;
+    typedef TexObjInputIterator<T> TextureIterator;
 
     TextureIterator d_tex_itr;
     CubDebugExit(d_tex_itr.BindTexture((CastT*) d_data, sizeof(T) * TEST_VALUES));