Skip to content
This repository was archived by the owner on Mar 21, 2024. It is now read-only.

Deprecate TexRefInputIterator. #320

Merged
merged 3 commits into from
Jun 14, 2021
Merged

Deprecate TexRefInputIterator. #320

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
9acf82a
Update deprecation mechanism (opt-out, cmake, c++14).
alliepiper Jun 9, 2021
4ef1a8e
Add missing return path.
alliepiper Jun 11, 2021
9c9fb93
Deprecate TexRefInputIterator (Replacement: TexObjInputIterator).
alliepiper Jun 9, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 11 additions & 1 deletion CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
4 changes: 2 additions & 2 deletions cub/agent/agent_spmv_orig.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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;
};


Expand Down
4 changes: 4 additions & 0 deletions cub/cmake/cub-config.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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")
Expand Down
7 changes: 6 additions & 1 deletion cub/iterator/tex_obj_input_iterator.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down
69 changes: 46 additions & 23 deletions cub/iterator/tex_ref_input_iterator.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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.
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't understand this comment, why "in favor of deprecating TexRefInputIterator"? Don't we disable warnings here so that they don't trigger in our own code, just on external usage?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The internal usages of deprecated CUDART APIs are always suppressed, and instead we deprecate the CUB utilities that call them. So instead of emitting deprecation warnings for cudaBindTexture, etc, we deprecate cub::TexRefInputIterator directly.

#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
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We do this so that we don't trigger deprecation warnings in our own code, only when users use the deprecated facility, correct?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Right.


/// And by unique ID
template <int UNIQUE_ID>
struct TexId
Expand Down Expand Up @@ -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

Expand All @@ -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.
Expand All @@ -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
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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

};


Expand All @@ -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
12 changes: 10 additions & 2 deletions cub/util_deprecated.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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))
Expand All @@ -43,4 +52,3 @@
#else
# define CUB_DEPRECATED
#endif

2 changes: 1 addition & 1 deletion test/test_iterator.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Expand Down