Alien-XGBoost
view release on metacpan - search on metacpan
view release on metacpan or search on metacpan
xgboost/cub/cub/thread/thread_load.cuh view on Meta::CPAN
* * 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
* Thread utilities for reading memory using PTX cache modifiers.
*/
#pragma once
#include <cuda.h>
#include <iterator>
#include "../util_ptx.cuh"
#include "../util_type.cuh"
#include "../util_namespace.cuh"
/// Optional outer namespace(s)
CUB_NS_PREFIX
/// CUB namespace
namespace cub {
/**
* \addtogroup UtilIo
* @{
*/
//-----------------------------------------------------------------------------
// Tags and constants
//-----------------------------------------------------------------------------
/**
* \brief Enumeration of cache modifiers for memory load operations.
*/
enum CacheLoadModifier
{
LOAD_DEFAULT, ///< Default (no modifier)
LOAD_CA, ///< Cache at all levels
LOAD_CG, ///< Cache at global level
LOAD_CS, ///< Cache streaming (likely to be accessed once)
LOAD_CV, ///< Cache as volatile (including cached system lines)
LOAD_LDG, ///< Cache as texture
LOAD_VOLATILE, ///< Volatile (any memory space)
};
/**
* \name Thread I/O (cache modified)
* @{
*/
/**
* \brief Thread utility for reading memory using cub::CacheLoadModifier cache modifiers. Can be used to load any data type.
*
* \par Example
* \code
* #include <cub/cub.cuh> // or equivalently <cub/thread/thread_load.cuh>
*
* // 32-bit load using cache-global modifier:
* int *d_in;
* int val = cub::ThreadLoad<cub::LOAD_CA>(d_in + threadIdx.x);
*
* // 16-bit load using default modifier
* short *d_in;
* short val = cub::ThreadLoad<cub::LOAD_DEFAULT>(d_in + threadIdx.x);
*
* // 256-bit load using cache-volatile modifier
* double4 *d_in;
* double4 val = cub::ThreadLoad<cub::LOAD_CV>(d_in + threadIdx.x);
*
* // 96-bit load using cache-streaming modifier
* struct TestFoo { bool a; short b; };
* TestFoo *d_struct;
* TestFoo val = cub::ThreadLoad<cub::LOAD_CS>(d_in + threadIdx.x);
* \endcode
*
* \tparam MODIFIER <b>[inferred]</b> CacheLoadModifier enumeration
* \tparam InputIteratorT <b>[inferred]</b> Input iterator type \iterator
*/
template <
CacheLoadModifier MODIFIER,
typename InputIteratorT>
__device__ __forceinline__ typename std::iterator_traits<InputIteratorT>::value_type ThreadLoad(InputIteratorT itr);
//@} end member group
#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
/// Helper structure for templated load iteration (inductive case)
template <int COUNT, int MAX>
struct IterateThreadLoad
{
template <CacheLoadModifier MODIFIER, typename T>
static __device__ __forceinline__ void Load(T const *ptr, T *vals)
{
vals[COUNT] = ThreadLoad<MODIFIER>(ptr + COUNT);
IterateThreadLoad<COUNT + 1, MAX>::template Load<MODIFIER>(ptr, vals);
}
template <typename InputIteratorT, typename T>
static __device__ __forceinline__ void Dereference(InputIteratorT itr, T *vals)
{
vals[COUNT] = itr[COUNT];
IterateThreadLoad<COUNT + 1, MAX>::Dereference(itr, vals);
}
};
/// Helper structure for templated load iteration (termination case)
template <int MAX>
struct IterateThreadLoad<MAX, MAX>
{
template <CacheLoadModifier MODIFIER, typename T>
static __device__ __forceinline__ void Load(T const * /*ptr*/, T * /*vals*/) {}
template <typename InputIteratorT, typename T>
static __device__ __forceinline__ void Dereference(InputIteratorT /*itr*/, T * /*vals*/) {}
};
/**
* Define a uint4 (16B) ThreadLoad specialization for the given Cache load modifier
*/
#define _CUB_LOAD_16(cub_modifier, ptx_modifier) \
template<> \
__device__ __forceinline__ uint4 ThreadLoad<cub_modifier, uint4 const *>(uint4 const *ptr) \
{ \
uint4 retval; \
view all matches for this distributionview release on metacpan - search on metacpan
( run in 0.409 second using v1.00-cache-2.02-grep-82fe00e-cpan-dad7e4baca0 )