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

Commit

Permalink
Alias deprecated TexRefInputIterator to TexObjInputIterator.
Browse files Browse the repository at this point in the history 
Removes usage of deprecated CUDART texture reference APIs without
breaking CUB's API.

Other changes:
- Remove obsolete checks for CUDART_VERSION >= 5.5.
- Split `test_iterator.cu` to create a test that only handles the
  deprecated `TexRefInputIterator` API.
  - Reduces the scope of our deprecation suppressions.
  - Also removed the deprecation suppression from our CMake logic.
- Enable testing for `TexObjInputIterator` without `CUB_CDP`.
- Add a meaningful implementation for `ostream << TexObjInputIterator`.
  • Loading branch information
@alliepiper
alliepiper committed Apr 6, 2022
1 parent 998ff61 commit bc7a1cc
Show file tree
Hide file tree
Showing 5 changed files with 316 additions and 424 deletions.
10 changes: 0 additions & 10 deletions cmake/CubBuildCompilerTargets.cmake
View file
Original file line number Diff line number Diff line change
Expand Up @@ -57,11 +57,6 @@ function(cub_build_compiler_targets)

# This complains about functions in CUDA system headers when used with nvcc.
append_option_if_available("-Wno-unused-function" cxx_compile_options)

# CUB uses deprecated texture functions (cudaBindTexture, etc). These
# need to be replaced, but silence the warnings for now.
# This can be removed once NVIDIA/cub#191 is fixed.
append_option_if_available("-Wno-deprecated-declarations" cxx_compile_options)
endif()

if ("GNU" STREQUAL "${CMAKE_CXX_COMPILER_ID}")
Expand Down Expand Up @@ -115,10 +110,5 @@ function(cub_build_compiler_targets)
$<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcudafe=--promote_warnings>
# Don't complain about deprecated GPU targets.
$<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Wno-deprecated-gpu-targets>
# Suppress deprecation warnings in nvcc < 11.5.
# TexRefInputIterator uses deprecated CUDART APIs, see NVIDIA/cub#191.
# After 11.5, we will suppress these in-code via pragma, but for older nvcc
# we have to use the big hammer:
$<$<AND:$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>,$<VERSION_LESS:$<CUDA_COMPILER_VERSION>,11.5>>:-Wno-deprecated-declarations>
)
endfunction()
3 changes: 3 additions & 0 deletions cub/iterator/tex_obj_input_iterator.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -306,6 +306,9 @@ public:
/// ostream operator
friend std::ostream& operator<<(std::ostream& os, const self_type& itr)
{
os << "cub::TexObjInputIterator( ptr=" << itr.ptr
<< ", offset=" << itr.tex_offset
<< ", tex_obj=" << itr.tex_obj << " )";
return os;
}

Expand Down
340 changes: 6 additions & 334 deletions cub/iterator/tex_ref_input_iterator.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -33,146 +33,20 @@

#pragma once

#include <iterator>
#include <iostream>
#include <cub/config.cuh>
#include <cub/iterator/tex_obj_input_iterator.cuh>

#include "../thread/thread_load.cuh"
#include "../thread/thread_store.cuh"
#include "../util_device.cuh"
#include "../util_debug.cuh"
#include "../config.cuh"

#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
#include <thrust/iterator/iterator_facade.h>
#include <thrust/iterator/iterator_traits.h>
#endif // THRUST_VERSION
#include <cstddef>

CUB_NAMESPACE_BEGIN


/******************************************************************************
* Static file-scope Tesla/Fermi-style texture references
*****************************************************************************/

#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document

// Anonymous namespace
namespace {

/// Global texture reference specialized by type
template <typename T>
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

#ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
#pragma nv_diagnostic push
#pragma nv_diag_suppress 1215
#endif

/// And by unique ID
template <int UNIQUE_ID>
struct TexId
{
// Largest texture word we can use in device
typedef typename UnitWord<T>::DeviceWord DeviceWord;
typedef typename UnitWord<T>::TextureWord TextureWord;

// Number of texture words per T
enum {
DEVICE_MULTIPLE = sizeof(T) / sizeof(DeviceWord),
TEXTURE_MULTIPLE = sizeof(T) / sizeof(TextureWord)
};

// Texture reference type
typedef texture<TextureWord> TexRef;

// Texture reference
static TexRef ref;

/// Bind texture
static cudaError_t BindTexture(void *d_in, size_t &bytes, size_t &offset)
{
if (d_in)
{
cudaChannelFormatDesc tex_desc = cudaCreateChannelDesc<TextureWord>();
ref.channelDesc = tex_desc;
return (CubDebug(cudaBindTexture(&offset, ref, d_in, bytes)));
}

return cudaSuccess;
}

/// Unbind texture
static cudaError_t UnbindTexture()
{
return CubDebug(cudaUnbindTexture(ref));
}

/// Fetch element
template <typename Distance>
static __device__ __forceinline__ T Fetch(Distance tex_offset)
{
DeviceWord temp[DEVICE_MULTIPLE];
TextureWord *words = reinterpret_cast<TextureWord*>(temp);

#pragma unroll
for (int i = 0; i < TEXTURE_MULTIPLE; ++i)
{
words[i] = tex1Dfetch(ref, (tex_offset * TEXTURE_MULTIPLE) + i);
}

return reinterpret_cast<T&>(temp);
}
};
};

// Texture reference definitions
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:
#ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
#pragma nv_diagnostic pop
#endif

#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


#endif // DOXYGEN_SHOULD_SKIP_THIS



/**
* \addtogroup UtilIterator
* @{
*/



/**
* \brief A random-access input wrapper for dereferencing array values through texture cache. Uses older Tesla/Fermi-style texture references.
* \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.
Expand Down Expand Up @@ -225,212 +99,10 @@ typename IteratorTexRef<T>::template TexId<UNIQUE_ID>::TexRef IteratorTexRef<T>:
*/
template <
typename T,
int UNIQUE_ID,
int /*UNIQUE_ID*/,
typename OffsetT = ptrdiff_t>
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

#ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
#pragma nv_diagnostic push
#pragma nv_diag_suppress 1215
#endif

public:

// Required iterator traits
typedef TexRefInputIterator self_type; ///< My own type
typedef OffsetT difference_type; ///< Type to express the result of subtracting one iterator from another
typedef T value_type; ///< The type of the element the iterator can point to
typedef T* pointer; ///< The type of a pointer to an element the iterator can point to
typedef T reference; ///< The type of a reference to an element the iterator can point to

#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_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
#else
typedef std::random_access_iterator_tag iterator_category; ///< The iterator category
#endif // THRUST_VERSION

private:

T* ptr;
difference_type tex_offset;

// Texture reference wrapper (old Tesla/Fermi-style textures)
typedef typename IteratorTexRef<T>::template TexId<UNIQUE_ID> TexId;

public:
/*
/// Constructor
__host__ __device__ __forceinline__ TexRefInputIterator()
:
ptr(NULL),
tex_offset(0)
{}
*/
/// Use this iterator to bind \p ptr with a texture reference
template <typename QualifiedT>
cudaError_t BindTexture(
QualifiedT *ptr, ///< Native pointer to wrap that is aligned to cudaDeviceProp::textureAlignment
size_t bytes, ///< Number of bytes in the range
size_t tex_offset = 0) ///< OffsetT (in items) from \p ptr denoting the position of the iterator
{
this->ptr = const_cast<typename std::remove_cv<QualifiedT>::type *>(ptr);
size_t offset;
cudaError_t retval = TexId::BindTexture(this->ptr + tex_offset, bytes, offset);
this->tex_offset = (difference_type) (offset / sizeof(QualifiedT));
return retval;
}

/// Unbind this iterator from its texture reference
cudaError_t UnbindTexture()
{
return TexId::UnbindTexture();
}

/// Postfix increment
__host__ __device__ __forceinline__ self_type operator++(int)
{
self_type retval = *this;
tex_offset++;
return retval;
}

/// Prefix increment
__host__ __device__ __forceinline__ self_type operator++()
{
tex_offset++;
return *this;
}

/// Indirection
__host__ __device__ __forceinline__ reference operator*() const
{
if (CUB_IS_HOST_CODE) {
// Simply dereference the pointer on the host
return ptr[tex_offset];
} else {
#if CUB_INCLUDE_DEVICE_CODE
// Use the texture reference
return TexId::Fetch(tex_offset);
#else
// This is dead code that will never be executed. It is here
// only to avoid warnings about missing returns.
return ptr[tex_offset];
#endif
}
}

/// Addition
template <typename Distance>
__host__ __device__ __forceinline__ self_type operator+(Distance n) const
{
self_type retval;
retval.ptr = ptr;
retval.tex_offset = tex_offset + n;
return retval;
}

/// Addition assignment
template <typename Distance>
__host__ __device__ __forceinline__ self_type& operator+=(Distance n)
{
tex_offset += n;
return *this;
}

/// Subtraction
template <typename Distance>
__host__ __device__ __forceinline__ self_type operator-(Distance n) const
{
self_type retval;
retval.ptr = ptr;
retval.tex_offset = tex_offset - n;
return retval;
}

/// Subtraction assignment
template <typename Distance>
__host__ __device__ __forceinline__ self_type& operator-=(Distance n)
{
tex_offset -= n;
return *this;
}

/// Distance
__host__ __device__ __forceinline__ difference_type operator-(self_type other) const
{
return tex_offset - other.tex_offset;
}

/// Array subscript
template <typename Distance>
__host__ __device__ __forceinline__ reference operator[](Distance n) const
{
self_type offset = (*this) + n;
return *offset;
}

/// Structure dereference
__host__ __device__ __forceinline__ pointer operator->()
{
return &(*(*this));
}

/// Equal to
__host__ __device__ __forceinline__ bool operator==(const self_type& rhs)
{
return ((ptr == rhs.ptr) && (tex_offset == rhs.tex_offset));
}

/// Not equal to
__host__ __device__ __forceinline__ bool operator!=(const self_type& rhs)
{
return ((ptr != rhs.ptr) || (tex_offset != rhs.tex_offset));
}

/// ostream operator
friend std::ostream& operator<<(std::ostream& os, const self_type& /*itr*/)
{
return os;
}

// Re-enable deprecation warnings:

#ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
#pragma nv_diagnostic pop
#endif

#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

};


using TexRefInputIterator CUB_DEPRECATED = cub::TexObjInputIterator<T, OffsetT>;

/** @} */ // end group UtilIterator

CUB_NAMESPACE_END

#endif // CUDART_VERSION
Loading

0 comments on commit bc7a1cc

Please sign in to comment.