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

Commit

Permalink
Merge pull request #451 from allisonvacanti/depr_tex_ref
Browse files Browse the repository at this point in the history 
Alias deprecated TexRefInputIterator to TexObjInputIterator.
  • Loading branch information
@alliepiper
alliepiper authored Apr 11, 2022
2 parents 998ff61 + 5cd095b commit ab5ee72
Show file tree
Hide file tree
Showing 5 changed files with 317 additions and 425 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
342 changes: 7 additions & 335 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,
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

};


int /*UNIQUE_ID*/,
typename OffsetT = std::ptrdiff_t>
using TexRefInputIterator CUB_DEPRECATED = cub::TexObjInputIterator<T, OffsetT>;

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

CUB_NAMESPACE_END

#endif // CUDART_VERSION
Loading

0 comments on commit ab5ee72

Please sign in to comment.