/****************************************************************************** * Copyright (c) 2011, Duane Merrill. All rights reserved. * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved. * * Redistribution and use in source and binary forms, with or without * modification, are permitted provided that the following conditions are met: * * Redistributions of source code must retain the above copyright * notice, this list of conditions and the following disclaimer. * * Redistributions in binary form must reproduce the above copyright * notice, this list of conditions and the following disclaimer in the * documentation and/or other materials provided with the distribution. * * Neither the name of the NVIDIA CORPORATION nor the * names of its contributors may be used to endorse or promote products * derived from this software without specific prior written permission. * * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. * ******************************************************************************/ /** * \file * Random-access iterator types */ #pragma once #include #include #include "../util_namespace.cuh" #include "../util_macro.cuh" #if (THRUST_VERSION >= 100700) // This iterator is compatible with Thrust API 1.7 and newer #include #include #endif // THRUST_VERSION /// Optional outer namespace(s) CUB_NS_PREFIX /// CUB namespace namespace cub { /** * \addtogroup UtilIterator * @{ */ /** * \brief A discard iterator */ template class DiscardOutputIterator { public: // Required iterator traits typedef DiscardOutputIterator self_type; ///< My own type typedef OffsetT difference_type; ///< Type to express the result of subtracting one iterator from another typedef void value_type; ///< The type of the element the iterator can point to typedef void pointer; ///< The type of a pointer to an element the iterator can point to typedef void 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::detail::iterator_facade_category< thrust::any_system_tag, thrust::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: OffsetT offset; #if defined(_WIN32) || !defined(_WIN64) // Workaround for win32 parameter-passing bug (ulonglong2 argmin DeviceReduce) OffsetT pad[CUB_MAX(1, (16 / sizeof(OffsetT) - 1))]; #endif public: /// Constructor __host__ __device__ __forceinline__ DiscardOutputIterator( OffsetT offset = 0) ///< Base offset : offset(offset) {} /// Postfix increment __host__ __device__ __forceinline__ self_type operator++(int) { self_type retval = *this; offset++; return retval; } /// Prefix increment __host__ __device__ __forceinline__ self_type operator++() { offset++; return *this; } /// Indirection __host__ __device__ __forceinline__ self_type& operator*() { // return self reference, which can be assigned to anything return *this; } /// Addition template __host__ __device__ __forceinline__ self_type operator+(Distance n) const { self_type retval(offset + n); return retval; } /// Addition assignment template __host__ __device__ __forceinline__ self_type& operator+=(Distance n) { offset += n; return *this; } /// Subtraction template __host__ __device__ __forceinline__ self_type operator-(Distance n) const { self_type retval(offset - n); return retval; } /// Subtraction assignment template __host__ __device__ __forceinline__ self_type& operator-=(Distance n) { offset -= n; return *this; } /// Distance __host__ __device__ __forceinline__ difference_type operator-(self_type other) const { return offset - other.offset; } /// Array subscript template __host__ __device__ __forceinline__ self_type& operator[](Distance n) { // return self reference, which can be assigned to anything return *this; } /// Structure dereference __host__ __device__ __forceinline__ pointer operator->() { return; } /// Assignment to self (no-op) __host__ __device__ __forceinline__ void operator=(self_type const& other) { offset = other.offset; } /// Assignment to anything else (no-op) template __host__ __device__ __forceinline__ void operator=(T const&) {} /// Cast to void* operator __host__ __device__ __forceinline__ operator void*() const { return NULL; } /// Equal to __host__ __device__ __forceinline__ bool operator==(const self_type& rhs) { return (offset == rhs.offset); } /// Not equal to __host__ __device__ __forceinline__ bool operator!=(const self_type& rhs) { return (offset != rhs.offset); } /// ostream operator friend std::ostream& operator<<(std::ostream& os, const self_type& itr) { os << "[" << itr.offset << "]"; return os; } }; /** @} */ // end group UtilIterator } // CUB namespace CUB_NS_POSTFIX // Optional outer namespace(s)