Data Structures
struct	cub::Equality
	Default equality functor. More...

struct	cub::Inequality
	Default inequality functor. More...

struct	cub::InequalityWrapper< EqualityOp >
	Inequality functor (wraps equality functor) More...

struct	cub::Sum
	Default sum functor. More...

struct	cub::Max
	Default max functor. More...

struct	cub::ArgMax
	Arg max functor (keeps the value and offset of the first occurrence of the larger item) More...

struct	cub::Min
	Default min functor. More...

struct	cub::ArgMin
	Arg min functor (keeps the value and offset of the first occurrence of the smallest item) More...

struct	cub::CastOp< B >
	Default cast functor. More...

class	cub::SwizzleScanOp< ScanOp >
	Binary operator wrapper for switching non-commutative scan arguments. More...

struct	cub::ReduceBySegmentOp< ReductionOpT >
	Reduce-by-segment functor. More...

struct	cub::ReduceByKeyOp< ReductionOpT >
	< Binary reduction operator to apply to values More...

struct	cub::If< IF, ThenType, ElseType >
	Type selection (`IF ? ThenType : ElseType`) More...

struct	cub::If< false, ThenType, ElseType >

struct	cub::Equals< A, B >
	Type equality test. More...

struct	cub::Equals< A, A >

struct	cub::Log2< N, CURRENT_VAL, COUNT >
	Statically determine log2(N), rounded up. More...

struct	cub::Log2< N, 0, COUNT >

struct	cub::PowerOfTwo< N >
	Statically determine if N is a power-of-two. More...

struct	cub::IsPointer< Tp >
	Pointer vs. iterator. More...

struct	cub::IsPointer< Tp * >

struct	cub::IsVolatile< Tp >
	Volatile modifier test. More...

struct	cub::IsVolatile< Tp volatile >

struct	cub::RemoveQualifiers< Tp, Up >
	Removes `const` and `volatile` qualifiers from type `Tp`. More...

struct	cub::RemoveQualifiers< Tp, volatile Up >

struct	cub::RemoveQualifiers< Tp, const Up >

struct	cub::RemoveQualifiers< Tp, const volatile Up >

struct	cub::NullType
	A simple "NULL" marker type. More...

struct	cub::Int2Type< A >
	Allows for the treatment of an integral constant as a type at compile-time (e.g., to achieve static call dispatch based on constant integral values) More...

struct	cub::AlignBytes< T >
	Structure alignment. More...

struct	cub::AlignBytes< volatile T >

struct	cub::AlignBytes< const T >

struct	cub::AlignBytes< const volatile T >

struct	cub::UnitWord< T >
	Unit-words of data movement. More...

struct	cub::UnitWord< float2 >

struct	cub::UnitWord< float4 >

struct	cub::UnitWord< char2 >

struct	cub::UnitWord< volatile T >

struct	cub::UnitWord< const T >

struct	cub::UnitWord< const volatile T >

struct	cub::CubVector< T, vec_elements >
	Exposes a member typedef `Type` that names the corresponding CUDA vector type if one exists. Otherwise `Type` refers to the CubVector structure itself, which will wrap the corresponding `x`, `y`, etc. vector fields. More...

struct	cub::CubVector< T, 1 >

struct	cub::CubVector< T, 2 >

struct	cub::CubVector< T, 3 >

struct	cub::CubVector< T, 4 >

struct	cub::Uninitialized< T >
	A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions. More...

struct	cub::KeyValuePair< _Key, _Value >
	A key identifier paired with a corresponding value. More...

struct	cub::ArrayWrapper< T, COUNT >
	A wrapper for passing simple static arrays as kernel parameters. More...

struct	cub::DoubleBuffer< T >
	Double-buffer storage wrapper for multi-pass stream transformations that require more than one storage array for streaming intermediate results back and forth. More...

struct	cub::EnableIf< Condition, T >
	Simple enable-if (similar to Boost) More...

struct	cub::EnableIf< false, T >

struct	cub::BinaryOpHasIdxParam< T, BinaryOp >
	Determine whether or not BinaryOp's functor is of the form `bool operator()(const T& a, const T&b)` or `bool operator()(const T& a, const T&b, unsigned int idx)` More...

struct	cub::BaseTraits< _CATEGORY, _PRIMITIVE, _NULL_TYPE, _UnsignedBits, T >
	Basic type traits. More...

struct	cub::BaseTraits< UNSIGNED_INTEGER, true, false, _UnsignedBits, T >

struct	cub::BaseTraits< SIGNED_INTEGER, true, false, _UnsignedBits, T >

struct	cub::FpLimits< _T >

struct	cub::FpLimits< float >

struct	cub::FpLimits< double >

struct	cub::BaseTraits< FLOATING_POINT, true, false, _UnsignedBits, T >

struct	cub::NumericTraits< T >
	Numeric type traits. More...

struct	cub::NumericTraits< NullType >

struct	cub::NumericTraits< char >

struct	cub::NumericTraits< signed char >

struct	cub::NumericTraits< short >

struct	cub::NumericTraits< int >

struct	cub::NumericTraits< long >

struct	cub::NumericTraits< long long >

struct	cub::NumericTraits< unsigned char >

struct	cub::NumericTraits< unsigned short >

struct	cub::NumericTraits< unsigned int >

struct	cub::NumericTraits< unsigned long >

struct	cub::NumericTraits< unsigned long long >

struct	cub::NumericTraits< float >

struct	cub::NumericTraits< double >

struct	cub::NumericTraits< bool >

struct	cub::Traits< T >
	Type traits. More...

Macros
#define	CUB_ALIGN(bytes) __attribute__((aligned(bytes)))
	Align struct.

#define	CUB_MAX(a, b) (((b) > (a)) ? (b) : (a))
	Select maximum(a, b)

#define	CUB_MIN(a, b) (((b) < (a)) ? (b) : (a))
	Select minimum(a, b)

#define	CUB_QUOTIENT_FLOOR(x, y) ((x) / (y))
	Quotient of x/y rounded down to nearest integer.

#define	CUB_QUOTIENT_CEILING(x, y) (((x) + (y) - 1) / (y))
	Quotient of x/y rounded up to nearest integer.

#define	CUB_ROUND_UP_NEAREST(x, y) ((((x) + (y) - 1) / (y)) * y)
	x rounded up to the nearest multiple of y

#define	CUB_ROUND_DOWN_NEAREST(x, y) (((x) / (y)) * y)
	x rounded down to the nearest multiple of y

#define	CUB_CAT_(a, b) a ## b

#define	CUB_CAT(a, b) CUB_CAT_(a, b)

#define	CUB_STATIC_ASSERT(cond, msg) typedef int CUB_CAT(cub_static_assert, __LINE__)[(cond) ? 1 : -1]
	Static assert.

#define	__CUB_ALIGN_BYTES(t, b)

#define	CUB_DEFINE_VECTOR_TYPE(base_type, short_type)

#define	CUB_DEFINE_DETECT_NESTED_TYPE(detector_name, nested_type_name)
	Defines a structure `detector_name` that is templated on type `T`. The `detector_name` struct exposes a constant member `VALUE` indicating whether or not parameter `T` exposes a nested type `nested_type_name`.

Enumerations
enum	{ cub::MAX_VEC_ELEMENTS = 4 }

enum	cub::Category { NOT_A_NUMBER , SIGNED_INTEGER , UNSIGNED_INTEGER , FLOATING_POINT }
	Basic type traits categories. More...

Sequential prefix scan over statically-sized array types
template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanExclusive (T inclusive, T exclusive, T input, T output, ScanOp scan_op, Int2Type< LENGTH >)

template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanExclusive (T input, T output, ScanOp scan_op, T prefix, bool apply_prefix=true)
	Perform a sequential exclusive prefix scan over `LENGTH` elements of the `input` array, seeded with the specified `prefix`. The aggregate is returned.

template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanExclusive (T(&input)[LENGTH], T(&output)[LENGTH], ScanOp scan_op, T prefix, bool apply_prefix=true)
	Perform a sequential exclusive prefix scan over the statically-sized `input` array, seeded with the specified `prefix`. The aggregate is returned.

template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanInclusive (T inclusive, T input, T output, ScanOp scan_op, Int2Type< LENGTH >)

template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanInclusive (T input, T output, ScanOp scan_op)
	Perform a sequential inclusive prefix scan over `LENGTH` elements of the `input` array. The aggregate is returned.

template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanInclusive (T(&input)[LENGTH], T(&output)[LENGTH], ScanOp scan_op)
	Perform a sequential inclusive prefix scan over the statically-sized `input` array. The aggregate is returned.

template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanInclusive (T input, T output, ScanOp scan_op, T prefix, bool apply_prefix=true)
	Perform a sequential inclusive prefix scan over `LENGTH` elements of the `input` array, seeded with the specified `prefix`. The aggregate is returned.

template<int LENGTH, typename T , typename ScanOp >
__device__ __forceinline__ T	cub::internal::ThreadScanInclusive (T(&input)[LENGTH], T(&output)[LENGTH], ScanOp scan_op, T prefix, bool apply_prefix=true)
	Perform a sequential inclusive prefix scan over the statically-sized `input` array, seeded with the specified `prefix`. The aggregate is returned.

Detailed Description

Macro Definition Documentation

◆ __CUB_ALIGN_BYTES

#define __CUB_ALIGN_BYTES	(	t,
		b
	)

Value:

template <> struct AlignBytes<t> \

{ enum { ALIGN_BYTES = b }; typedef __align__(b) t Type; };

Definition at line 312 of file util_type.cuh.

◆ CUB_ALIGN

#define CUB_ALIGN ( bytes ) __attribute__((aligned(bytes)))

Align struct.

Definition at line 55 of file util_macro.cuh.

◆ CUB_CAT

#define CUB_CAT	(	a,
		b
	)	CUB_CAT_(a, b)

Definition at line 93 of file util_macro.cuh.

◆ CUB_CAT_

#define CUB_CAT_	(	a,
		b
	)	a ## b

Definition at line 92 of file util_macro.cuh.

◆ CUB_DEFINE_DETECT_NESTED_TYPE

#define CUB_DEFINE_DETECT_NESTED_TYPE	(	detector_name,
		nested_type_name
	)

Value:

    template <typename T>                                               \
    struct detector_name                                                \
    {                                                                   \
        template <typename C>                                           \
        static char& test(typename C::nested_type_name*);               \
        template <typename>                                             \
        static int& test(...);                                          \
        enum                                                            \
        {                                                               \
            VALUE = sizeof(test<T>(0)) < sizeof(int)                    \
        };                                                              \
    };

Defines a structure detector_name that is templated on type T. The detector_name struct exposes a constant member VALUE indicating whether or not parameter T exposes a nested type nested_type_name.

Definition at line 835 of file util_type.cuh.

◆ CUB_DEFINE_VECTOR_TYPE

#define CUB_DEFINE_VECTOR_TYPE	(	base_type,
		short_type
	)

Macro for expanding partially-specialized built-in vector types

Definition at line 523 of file util_type.cuh.

◆ CUB_MAX

#define CUB_MAX	(	a,
		b
	)	(((b) > (a)) ? (b) : (a))

Select maximum(a, b)

Definition at line 61 of file util_macro.cuh.

◆ CUB_MIN

#define CUB_MIN	(	a,
		b
	)	(((b) < (a)) ? (b) : (a))

Select minimum(a, b)

Definition at line 66 of file util_macro.cuh.

◆ CUB_QUOTIENT_CEILING

#define CUB_QUOTIENT_CEILING	(	x,
		y
	)	(((x) + (y) - 1) / (y))

Quotient of x/y rounded up to nearest integer.

Definition at line 76 of file util_macro.cuh.

◆ CUB_QUOTIENT_FLOOR

#define CUB_QUOTIENT_FLOOR	(	x,
		y
	)	((x) / (y))

Quotient of x/y rounded down to nearest integer.

Definition at line 71 of file util_macro.cuh.

◆ CUB_ROUND_DOWN_NEAREST

#define CUB_ROUND_DOWN_NEAREST	(	x,
		y
	)	(((x) / (y)) * y)

x rounded down to the nearest multiple of y

Definition at line 86 of file util_macro.cuh.

◆ CUB_ROUND_UP_NEAREST

#define CUB_ROUND_UP_NEAREST	(	x,
		y
	)	((((x) + (y) - 1) / (y)) * y)

x rounded up to the nearest multiple of y

Definition at line 81 of file util_macro.cuh.

◆ CUB_STATIC_ASSERT

#define CUB_STATIC_ASSERT	(	cond,
		msg
	)	typedef int CUB_CAT(cub_static_assert, __LINE__)[(cond) ? 1 : -1]

Static assert.

Definition at line 97 of file util_macro.cuh.

Enumeration Type Documentation

◆ anonymous enum

anonymous enum

Enumerator
MAX_VEC_ELEMENTS	The maximum number of elements in CUDA vector types.

Definition at line 458 of file util_type.cuh.

◆ Category

enum cub::Category

Basic type traits categories.

Definition at line 928 of file util_type.cuh.

Function Documentation

◆ ThreadScanExclusive() [1/3]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanExclusive	(	T *	input,
		T *	output,
		ScanOp	scan_op,
		T	prefix,
		bool	apply_prefix = `true`
	)

Perform a sequential exclusive prefix scan over LENGTH elements of the input array, seeded with the specified prefix. The aggregate is returned.

Template Parameters

LENGTH	LengthT of `input` and `output` arrays
T	[inferred] The data type to be scanned.
ScanOp	[inferred] Binary scan operator type having member `T operator()(const T &a, const T &b)`

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator
[in]	prefix	Prefix to seed scan with
[in]	apply_prefix	Whether or not the calling thread should apply its prefix. If not, the first output element is undefined. (Handy for preventing thread-0 from applying a prefix.)

Definition at line 95 of file thread_scan.cuh.

◆ ThreadScanExclusive() [2/3]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanExclusive	(	T	inclusive,
		T	exclusive,
		T *	input,
		T *	output,
		ScanOp	scan_op,
		Int2Type< LENGTH >
	)

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator

Definition at line 63 of file thread_scan.cuh.

◆ ThreadScanExclusive() [3/3]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanExclusive	(	T(&)	input[LENGTH],
		T(&)	output[LENGTH],
		ScanOp	scan_op,
		T	prefix,
		bool	apply_prefix = `true`
	)

Perform a sequential exclusive prefix scan over the statically-sized input array, seeded with the specified prefix. The aggregate is returned.

Template Parameters

LENGTH	[inferred] LengthT of `input` and `output` arrays
T	[inferred] The data type to be scanned.
ScanOp	[inferred] Binary scan operator type having member `T operator()(const T &a, const T &b)`

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator
[in]	prefix	Prefix to seed scan with
[in]	apply_prefix	Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

Definition at line 125 of file thread_scan.cuh.

◆ ThreadScanInclusive() [1/5]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanInclusive	(	T *	input,
		T *	output,
		ScanOp	scan_op
	)

Perform a sequential inclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Template Parameters

LENGTH	LengthT of `input` and `output` arrays
T	[inferred] The data type to be scanned.
ScanOp	[inferred] Binary scan operator type having member `T operator()(const T &a, const T &b)`

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator

Definition at line 176 of file thread_scan.cuh.

◆ ThreadScanInclusive() [2/5]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanInclusive	(	T *	input,
		T *	output,
		ScanOp	scan_op,
		T	prefix,
		bool	apply_prefix = `true`
	)

Perform a sequential inclusive prefix scan over LENGTH elements of the input array, seeded with the specified prefix. The aggregate is returned.

Template Parameters

LENGTH	LengthT of `input` and `output` arrays
T	[inferred] The data type to be scanned.
ScanOp	[inferred] Binary scan operator type having member `T operator()(const T &a, const T &b)`

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator
[in]	prefix	Prefix to seed scan with
[in]	apply_prefix	Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

Definition at line 220 of file thread_scan.cuh.

◆ ThreadScanInclusive() [3/5]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanInclusive	(	T	inclusive,
		T *	input,
		T *	output,
		ScanOp	scan_op,
		Int2Type< LENGTH >
	)

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator

Definition at line 147 of file thread_scan.cuh.

◆ ThreadScanInclusive() [4/5]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanInclusive	(	T(&)	input[LENGTH],
		T(&)	output[LENGTH],
		ScanOp	scan_op
	)

Perform a sequential inclusive prefix scan over the statically-sized input array. The aggregate is returned.

Template Parameters

LENGTH	[inferred] LengthT of `input` and `output` arrays
T	[inferred] The data type to be scanned.
ScanOp	[inferred] Binary scan operator type having member `T operator()(const T &a, const T &b)`

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator

Definition at line 200 of file thread_scan.cuh.

◆ ThreadScanInclusive() [5/5]

template<int LENGTH, typename T , typename ScanOp >

__device__ __forceinline__ T cub::internal::ThreadScanInclusive	(	T(&)	input[LENGTH],
		T(&)	output[LENGTH],
		ScanOp	scan_op,
		T	prefix,
		bool	apply_prefix = `true`
	)

Perform a sequential inclusive prefix scan over the statically-sized input array, seeded with the specified prefix. The aggregate is returned.

Template Parameters

LENGTH	[inferred] LengthT of `input` and `output` arrays
T	[inferred] The data type to be scanned.
ScanOp	[inferred] Binary scan operator type having member `T operator()(const T &a, const T &b)`

Parameters

[in]	input	Input array
[out]	output	Output array (may be aliased to `input`)
[in]	scan_op	Binary scan operator
[in]	prefix	Prefix to seed scan with
[in]	apply_prefix	Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

Definition at line 250 of file thread_scan.cuh.

Data Structures

Macros

Enumerations

Sequential prefix scan over statically-sized array types

Detailed Description

Macro Definition Documentation

◆ __CUB_ALIGN_BYTES

◆ CUB_ALIGN

◆ CUB_CAT

◆ CUB_CAT_

◆ CUB_DEFINE_DETECT_NESTED_TYPE

◆ CUB_DEFINE_VECTOR_TYPE

◆ CUB_MAX

◆ CUB_MIN

◆ CUB_QUOTIENT_CEILING

◆ CUB_QUOTIENT_FLOOR

◆ CUB_ROUND_DOWN_NEAREST

◆ CUB_ROUND_UP_NEAREST

◆ CUB_STATIC_ASSERT

Enumeration Type Documentation

◆ anonymous enum

◆ Category

Function Documentation

◆ ThreadScanExclusive() [1/3]

◆ ThreadScanExclusive() [2/3]

◆ ThreadScanExclusive() [3/3]

◆ ThreadScanInclusive() [1/5]

◆ ThreadScanInclusive() [2/5]

◆ ThreadScanInclusive() [3/5]

◆ ThreadScanInclusive() [4/5]

◆ ThreadScanInclusive() [5/5]