cub::DeviceSelect#

struct DeviceSelect#

DeviceSelect provides device-wide, parallel operations for compacting selected items from sequences of data items residing within device-accessible memory. It is similar to DevicePartition, except that non-selected items are discarded, whereas DevicePartition retains them.

Overview#

These operations apply a selection criterion to selectively copy items from a specified input sequence to a compact output sequence.

Usage Considerations#

  • Dynamic parallelism. DeviceSelect methods can be called within kernel code on devices in which CUDA dynamic parallelism is supported.

Performance#

The work-complexity of select-flagged, select-if, and select-unique as a function of input size is linear, resulting in performance throughput that plateaus with problem sizes large enough to saturate the GPU.

Tuning#

All non-ByKey algorithms in DeviceSelect that accept an environment can be tuned by passing a custom policy selector that returns a @ref SelectPolicy, as shown in the example below:

struct SelectPolicySelector
{
  __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::SelectPolicy
  {
    return {.threads_per_block = 128,
            .items_per_thread  = cc > cuda::compute_capability{9, 0} ? 16 : 10,
            .load_algorithm    = cub::BLOCK_LOAD_DIRECT,
            .load_modifier     = cub::LOAD_DEFAULT,
            .scan_algorithm    = cub::BLOCK_SCAN_WARP_SCANS,
            .lookback_delay    = {cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}};
  }
};

auto d_in           = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto d_out          = thrust::device_vector<int>(4, thrust::no_init);
auto d_num_selected = thrust::device_vector<int>(1, thrust::no_init);

const auto error = cub::DeviceSelect::If(
  d_in.begin(),
  d_out.begin(),
  d_num_selected.begin(),
  d_in.size(),
  [] __host__ __device__(int v) {
    return v < 5;
  },
  cuda::execution::tune(SelectPolicySelector{}));
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::If failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{1, 2, 3, 4};
int expected_num_selected = 4;

All *ByKey algorithms in DeviceSelect that accept an environment can be tuned by passing a custom policy selector that returns a @ref UniqueByKeyPolicy, as shown in the example below:

struct UniqueByKeyPolicySelector
{
  __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::UniqueByKeyPolicy
  {
    return {.threads_per_block = 256,
            .items_per_thread  = cc > cuda::compute_capability{9, 0} ? 12 : 10,
            .load_algorithm    = cub::BLOCK_LOAD_DIRECT,
            .load_modifier     = cub::LOAD_DEFAULT,
            .scan_algorithm    = cub::BLOCK_SCAN_WARP_SCANS,
            .lookback_delay    = cub::LookbackDelayPolicy{
                 .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}};
  }
};

auto keys_in          = thrust::device_vector<int>{0, 2, 2, 9, 5, 5, 5, 8};
auto values_in        = thrust::device_vector<int>{0, 1, 2, 3, 4, 5, 6, 7};
auto keys_out         = thrust::device_vector<int>(8, thrust::no_init);
auto values_out       = thrust::device_vector<int>(8, thrust::no_init);
auto num_selected_out = thrust::device_vector<int>(1, thrust::no_init);

const auto error = cub::DeviceSelect::UniqueByKey(
  keys_in.begin(),
  values_in.begin(),
  keys_out.begin(),
  values_out.begin(),
  num_selected_out.begin(),
  keys_in.size(),
  cuda::execution::tune(UniqueByKeyPolicySelector{}));
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::UniqueByKey failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_keys{0, 2, 9, 5, 8};
thrust::device_vector<int> expected_values{0, 1, 3, 4, 7};

Public Static Functions

template<typename InputIteratorT, typename FlagIterator, typename OutputIteratorT, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<!::cuda::std::is_integral_v<NumSelectedIteratorT>, int> = 0>
static inline cudaError_t Flagged(
void *d_temp_storage,
size_t &temp_storage_bytes,
InputIteratorT d_in,
FlagIterator d_flags,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Uses the d_flags sequence to selectively copy the corresponding items from d_in into d_out. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • The value type of d_flags must be castable to bool (e.g., bool, char, int, etc.).

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap [d_in, d_in + num_items),
    [d_flags, d_flags + num_items) nor d_num_selected_out in any way.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>  // or equivalently <cub/device/device_select.cuh>

// Declare, allocate, and initialize device-accessible pointers for input,
// flags, and output
int  num_items;              // e.g., 8
int  *d_in;                  // e.g., [1, 2, 3, 4, 5, 6, 7, 8]
char *d_flags;               // e.g., [1, 0, 0, 1, 0, 1, 1, 0]
int  *d_out;                 // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int  *d_num_selected_out;    // e.g., [ ]
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::Flagged(
  d_temp_storage, temp_storage_bytes,
  d_in, d_flags, d_out, d_num_selected_out, num_items);

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::Flagged(
  d_temp_storage, temp_storage_bytes,
  d_in, d_flags, d_out, d_num_selected_out, num_items);

// d_out                 <-- [1, 4, 6, 7]
// d_num_selected_out    <-- [4]

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_in – [in] Pointer to the input sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename FlagIterator, typename OutputIteratorT, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t Flagged(
InputIteratorT d_in,
FlagIterator d_flags,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Uses the d_flags sequence to selectively copy the corresponding items from d_in into d_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The value type of d_flags must be castable to bool (e.g., bool, char, int, etc.).

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap [d_in, d_in + num_items),
    [d_flags, d_flags + num_items) nor d_num_selected_out in any way.

Snippet#

The code snippet below illustrates the compaction of flagged items from an int device vector using environment-based API:

auto input        = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto flags        = thrust::device_vector<char>{1, 0, 0, 1, 0, 1, 1, 0};
auto output       = thrust::device_vector<int>(4);
auto num_selected = thrust::device_vector<int>(1);

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::Flagged(
  input.begin(),
  flags.begin(),
  output.begin(),
  num_selected.begin(),
  static_cast<::cuda::std::int64_t>(input.size()),
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::Flagged failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{1, 4, 6, 7};
thrust::device_vector<int> expected_num_selected{4};

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_in – [in] Pointer to the input sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename FlagIterator, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t Flagged(
IteratorT d_data,
FlagIterator d_flags,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Uses the d_flags sequence to selectively compact items in d_data. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The value type of d_flags must be castable to bool (e.g., bool, char, int, etc.).

  • Copies of the selected items are compacted in-place and maintain their original relative ordering.

  • The d_data may equal d_flags. The range [d_data, d_data + num_items) shall not overlap
    [d_flags, d_flags + num_items) in any other way.

Snippet#

The code snippet below illustrates the in-place compaction of items selected from an int device vector using environment-based API:

auto data         = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto flags        = thrust::device_vector<char>{1, 0, 0, 1, 0, 1, 1, 0};
auto num_selected = thrust::device_vector<int>(1);

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::Flagged(
  data.begin(), flags.begin(), num_selected.begin(), static_cast<::cuda::std::int64_t>(data.size()), stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::Flagged in-place failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{1, 4, 6, 7};
thrust::device_vector<int> expected_num_selected{4};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing selected items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_data – [inout] Pointer to the sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename OutputIteratorT, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_unary_predicate<SelectOp, InputIteratorT>, int> = 0>
static inline cudaError_t If(
InputIteratorT d_in,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor to selectively copy items from d_in into d_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector using environment-based API:

auto input        = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto output       = thrust::device_vector<int>(4);
auto num_selected = thrust::device_vector<int>(1);
less_than_t<int> le{5};

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::If(
  input.begin(),
  output.begin(),
  num_selected.begin(),
  static_cast<::cuda::std::int64_t>(input.size()),
  le,
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::If failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{1, 2, 3, 4};
thrust::device_vector<int> expected_num_selected{4};

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_in – [in] Pointer to the input sequence of data items

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_unary_predicate<SelectOp, IteratorT>, int> = 0>
static inline cudaError_t If(
IteratorT d_data,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor to selectively compact items in d_data. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • Copies of the selected items are compacted in d_data and maintain their original relative ordering.

Snippet#

The code snippet below illustrates the in-place compaction of items selected from an int device vector using environment-based API:

auto data         = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto num_selected = thrust::device_vector<int>(1);
less_than_t<int> le{5};

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::If(
  data.begin(), num_selected.begin(), static_cast<::cuda::std::int64_t>(data.size()), le, stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::If in-place failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{1, 2, 3, 4};
thrust::device_vector<int> expected_num_selected{4};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_data – [inout] Pointer to the sequence of data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename FlagIterator, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t Flagged(
void *d_temp_storage,
size_t &temp_storage_bytes,
IteratorT d_data,
FlagIterator d_flags,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Uses the d_flags sequence to selectively compact the items in d_data`. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • The value type of d_flags must be castable to bool (e.g., bool, char, int, etc.).

  • Copies of the selected items are compacted in-place and maintain their original relative ordering.

  • The d_data may equal d_flags. The range [d_data, d_data + num_items) shall not overlap
    [d_flags, d_flags + num_items) in any other way.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>  // or equivalently <cub/device/device_select.cuh>

// Declare, allocate, and initialize device-accessible pointers for input,
// flags, and output
int  num_items;              // e.g., 8
int  *d_data;                // e.g., [1, 2, 3, 4, 5, 6, 7, 8]
char *d_flags;               // e.g., [1, 0, 0, 1, 0, 1, 1, 0]
int  *d_num_selected_out;    // e.g., [ ]
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::Flagged(
  d_temp_storage, temp_storage_bytes,
  d_in, d_flags, d_num_selected_out, num_items);

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::Flagged(
  d_temp_storage, temp_storage_bytes,
  d_in, d_flags, d_num_selected_out, num_items);

// d_data                <-- [1, 4, 6, 7]
// d_num_selected_out    <-- [4]

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing selected items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_data – [inout] Pointer to the sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename OutputIteratorT, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_unary_predicate<SelectOp, InputIteratorT>, int> = 0>
static inline cudaError_t If(
void *d_temp_storage,
size_t &temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor to selectively copy items from d_in into d_out. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>   // or equivalently <cub/device/device_select.cuh>

// Functor type for selecting values less than some criteria
struct LessThan
{
    int compare;

    __host__ __device__ __forceinline__
    LessThan(int compare) : compare(compare) {}

    __host__ __device__ __forceinline__
    bool operator()(const int &a) const {
        return (a < compare);
    }
};

// Declare, allocate, and initialize device-accessible pointers
// for input and output
int      num_items;              // e.g., 8
int      *d_in;                  // e.g., [0, 2, 3, 9, 5, 2, 81, 8]
int      *d_out;                 // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int      *d_num_selected_out;    // e.g., [ ]
LessThan select_op(7);
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::If(
  d_temp_storage, temp_storage_bytes,
  d_in, d_out, d_num_selected_out, num_items, select_op);

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::If(
  d_temp_storage, temp_storage_bytes,
  d_in, d_out, d_num_selected_out, num_items, select_op);

// d_out                 <-- [0, 2, 3, 5, 2]
// d_num_selected_out    <-- [5]

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_in – [in] Pointer to the input sequence of data items

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_unary_predicate<SelectOp, IteratorT>, int> = 0>
static inline cudaError_t If(
void *d_temp_storage,
size_t &temp_storage_bytes,
IteratorT d_data,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor to selectively compact items in d_data. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • Copies of the selected items are compacted in d_data and maintain
    their original relative ordering.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>   // or equivalently <cub/device/device_select.cuh>

// Functor type for selecting values less than some criteria
struct LessThan
{
    int compare;

    __host__ __device__ __forceinline__
    LessThan(int compare) : compare(compare) {}

    __host__ __device__ __forceinline__
    bool operator()(const int &a) const {
        return (a < compare);
    }
};

// Declare, allocate, and initialize device-accessible pointers
// for input and output
int      num_items;              // e.g., 8
int      *d_data;                // e.g., [0, 2, 3, 9, 5, 2, 81, 8]
int      *d_num_selected_out;    // e.g., [ ]
LessThan select_op(7);
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::If(
  d_temp_storage, temp_storage_bytes,
  d_data, d_num_selected_out, num_items, select_op);

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::If(
  d_temp_storage, temp_storage_bytes,
  d_data, d_num_selected_out, num_items, select_op);

// d_data                <-- [0, 2, 3, 5, 2]
// d_num_selected_out    <-- [5]

Template Parameters:

  • IteratorT – [inferred] Random-access input iterator type for reading and writing items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_data – [inout] Pointer to the sequence of data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename FlagIterator, typename OutputIteratorT, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t FlaggedIf(
void *d_temp_storage,
size_t &temp_storage_bytes,
InputIteratorT d_in,
FlagIterator d_flags,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor applied to d_flags to selectively copy the corresponding items from d_in into d_out. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • The expression select_op(flag) must be convertible to bool, where the type of flag corresponds to the value type of FlagIterator.

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

struct is_even_t
{
  __host__ __device__ bool operator()(int flag) const
  {
    return !(flag % 2);
  }
};

constexpr int num_items         = 8;
c2h::device_vector<int> d_in    = {0, 1, 2, 3, 4, 5, 6, 7};
c2h::device_vector<int> d_flags = {8, 6, 7, 5, 3, 0, 9, 3};
c2h::device_vector<int> d_out(num_items);
c2h::device_vector<int> d_num_selected_out(num_items);
is_even_t is_even{};

// Determine temporary device storage requirements
size_t temp_storage_bytes = 0;
cub::DeviceSelect::FlaggedIf(
  nullptr,
  temp_storage_bytes,
  d_in.begin(),
  d_flags.begin(),
  d_out.begin(),
  d_num_selected_out.data(),
  num_items,
  is_even);

// Allocate temporary storage
c2h::device_vector<char> temp_storage(temp_storage_bytes);

// Run selection
cub::DeviceSelect::FlaggedIf(
  thrust::raw_pointer_cast(temp_storage.data()),
  temp_storage_bytes,
  d_in.begin(),
  d_flags.begin(),
  d_out.begin(),
  d_num_selected_out.data(),
  num_items,
  is_even);

c2h::device_vector<int> expected{0, 1, 5};

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_in – [in] Pointer to the input sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename FlagIterator, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t FlaggedIf(
void *d_temp_storage,
size_t &temp_storage_bytes,
IteratorT d_data,
FlagIterator d_flags,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor applied to d_flags to selectively compact the corresponding items in d_data. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • The expression select_op(flag) must be convertible to bool, where the type of flag corresponds to the value type of FlagIterator.

  • Copies of the selected items are compacted in-place and maintain their original relative ordering.

  • The d_data may equal d_flags. The range [d_data, d_data + num_items) shall not overlap
    [d_flags, d_flags + num_items) in any other way.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

struct is_even_t
{
  __host__ __device__ bool operator()(int flag) const
  {
    return !(flag % 2);
  }
};

constexpr int num_items         = 8;
c2h::device_vector<int> d_data  = {0, 1, 2, 3, 4, 5, 6, 7};
c2h::device_vector<int> d_flags = {8, 6, 7, 5, 3, 0, 9, 3};
c2h::device_vector<int> d_num_selected_out(num_items);
is_even_t is_even{};

// Determine temporary device storage requirements
size_t temp_storage_bytes = 0;
cub::DeviceSelect::FlaggedIf(
  nullptr, temp_storage_bytes, d_data.begin(), d_flags.begin(), d_num_selected_out.data(), num_items, is_even);

// Allocate temporary storage
c2h::device_vector<char> temp_storage(temp_storage_bytes);

// Run selection
cub::DeviceSelect::FlaggedIf(
  thrust::raw_pointer_cast(temp_storage.data()),
  temp_storage_bytes,
  d_data.begin(),
  d_flags.begin(),
  d_num_selected_out.data(),
  num_items,
  is_even);

c2h::device_vector<int> expected{0, 1, 5};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing selected items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_data – [inout] Pointer to the sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename FlagIterator, typename OutputIteratorT, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_unary_predicate<SelectOp, FlagIterator>, int> = 0>
static inline cudaError_t FlaggedIf(
InputIteratorT d_in,
FlagIterator d_flags,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor applied to d_flags to selectively copy the corresponding items from d_in into d_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The expression select_op(flag) must be convertible to bool, where the type of flag corresponds to the value type of FlagIterator.

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector using environment-based API:

auto input        = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto flags        = thrust::device_vector<int>{2, 1, 1, 4, 1, 6, 6, 1};
auto output       = thrust::device_vector<int>(4);
auto num_selected = thrust::device_vector<int>(1);
mod_n<int> select_op{2};

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::FlaggedIf(
  input.begin(),
  flags.begin(),
  output.begin(),
  num_selected.begin(),
  static_cast<::cuda::std::int64_t>(input.size()),
  select_op,
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::FlaggedIf failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{1, 4, 6, 7};
thrust::device_vector<int> expected_num_selected{4};

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_in – [in] Pointer to the input sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename FlagIterator, typename NumSelectedIteratorT, typename SelectOp, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_unary_predicate<SelectOp, FlagIterator>, int> = 0>
static inline cudaError_t FlaggedIf(
IteratorT d_data,
FlagIterator d_flags,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
SelectOp select_op,
const EnvT &env = {}
)#

Uses the select_op functor applied to d_flags to selectively compact items in d_data. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The expression select_op(flag) must be convertible to bool, where the type of flag corresponds to the value type of FlagIterator.

  • Copies of the selected items are compacted in-place and maintain their original relative ordering.

  • The d_data may equal d_flags. The range [d_data, d_data + num_items) shall not overlap
    [d_flags, d_flags + num_items) in any other way.

Snippet#

The code snippet below illustrates the in-place compaction of items selected from an int device vector using environment-based API:

auto data         = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto flags        = thrust::device_vector<int>{2, 1, 1, 4, 1, 6, 6, 1};
auto num_selected = thrust::device_vector<int>(1);
mod_n<int> select_op{2};

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::FlaggedIf(
  data.begin(),
  flags.begin(),
  num_selected.begin(),
  static_cast<::cuda::std::int64_t>(data.size()),
  select_op,
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::FlaggedIf in-place failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{1, 4, 6, 7};
thrust::device_vector<int> expected_num_selected{4};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing selected items (may be a simple pointer type)

  • FlagIterator – [inferred] Random-access input iterator type for reading selection flags (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • SelectOp – [inferred] Selection operator type having member bool operator()(const T &a)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_data – [inout] Pointer to the sequence of data items

  • d_flags – [in] Pointer to the input sequence of selection flags

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • select_op – [in] Unary selection operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename OutputIteratorT, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<!::cuda::std::indirect_binary_predicate<EnvT, InputIteratorT, InputIteratorT>, int> = 0>
static inline cudaError_t Unique(
InputIteratorT d_in,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Given an input sequence d_in having runs of consecutive equal-valued keys, only the first key from each run is selectively copied to d_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The == equality operator is used to determine whether keys are equivalent.

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector using environment-based API:

auto input        = thrust::device_vector<int>{0, 2, 2, 9, 5, 5, 5, 8};
auto output       = thrust::device_vector<int>(5);
auto num_selected = thrust::device_vector<int>(1);

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::Unique(
  input.begin(), output.begin(), num_selected.begin(), static_cast<::cuda::std::int64_t>(input.size()), stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::Unique failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{0, 2, 9, 5, 8};
thrust::device_vector<int> expected_num_selected{5};

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_in – [in] Pointer to the input sequence of data items

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename OutputIteratorT, typename NumSelectedIteratorT, typename EqualityOpT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_binary_predicate<EqualityOpT, InputIteratorT, InputIteratorT>, int> = 0>
static inline cudaError_t Unique(
InputIteratorT d_in,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
EqualityOpT equality_op,
const EnvT &env = {}
)#

Given an input sequence d_in having runs of consecutive equal-valued keys, only the first key from each run is selectively copied to d_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The user-provided equality operator, equality_op, is used to determine whether keys are equivalent.

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector using a custom equality operator:

// Unique modulo 3 — consecutive elements are "equal" if they have the same remainder mod 3
auto input        = thrust::device_vector<int>{0, 3, 6, 1, 4, 7, 2, 5};
auto output       = thrust::device_vector<int>(8);
auto num_selected = thrust::device_vector<int>(1);

eq_mod3_t<int> eq_mod3{};

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::Unique(
  input.begin(),
  output.begin(),
  num_selected.begin(),
  static_cast<::cuda::std::int64_t>(input.size()),
  eq_mod3,
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::Unique with custom equality_op failed with status: " << error << '\n';
}

// 0,3,6 all == mod 3 (consecutive), keep first (0); 1,4,7 all == mod 3 (consecutive), keep first (1);
// 2,5 == mod 3 (consecutive), keep first (2)
thrust::device_vector<int> expected_output{0, 1, 2};
thrust::device_vector<int> expected_num_selected{3};

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EqualityOpT – [inferred] Type of equality_op

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_in – [in] Pointer to the input sequence of data items

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • equality_op – [in] Binary equality operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<!::cuda::std::indirect_binary_predicate<EnvT, IteratorT, IteratorT>, int> = 0>
static inline cudaError_t Unique(
IteratorT d_data,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Given an input sequence d_data having runs of consecutive equal-valued keys, only the first key from each run is selectively compacted in-place. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The == equality operator is used to determine whether keys are equivalent.

  • Copies of the selected items are compacted in d_data and maintain their original relative ordering.

Snippet#

The code snippet below illustrates the in-place compaction of items selected from an int device vector using environment-based API:

auto data         = thrust::device_vector<int>{0, 2, 2, 9, 5, 5, 5, 8};
auto num_selected = thrust::device_vector<int>(1);

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::Unique(
  data.begin(), num_selected.begin(), static_cast<::cuda::std::int64_t>(data.size()), stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::Unique in-place failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_output{0, 2, 9, 5, 8};
thrust::device_vector<int> expected_num_selected{5};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_data – [inout] Pointer to the sequence of data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename NumSelectedIteratorT, typename EqualityOpT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_binary_predicate<EqualityOpT, IteratorT, IteratorT>, int> = 0>
static inline cudaError_t Unique(
IteratorT d_data,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
EqualityOpT equality_op,
const EnvT &env = {}
)#

Given an input sequence d_data having runs of consecutive equal-valued keys, only the first key from each run is selectively compacted in-place. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The user-provided equality operator, equality_op, is used to determine whether keys are equivalent.

  • Copies of the selected items are compacted in d_data and maintain their original relative ordering.

Snippet#

The code snippet below illustrates the in-place compaction of items selected from an int device vector using a custom equality operator:

// Unique modulo 3 — consecutive elements are "equal" if they have the same remainder mod 3
auto data         = thrust::device_vector<int>{0, 3, 6, 1, 4, 7, 2, 5};
auto num_selected = thrust::device_vector<int>(1);

eq_mod3_t<int> eq_mod3{};

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::Unique(
  data.begin(), num_selected.begin(), static_cast<::cuda::std::int64_t>(data.size()), eq_mod3, stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::Unique in-place with custom equality_op failed with status: " << error << '\n';
}

// 0,3,6 all == mod 3 (consecutive), keep first (0); 1,4,7 all == mod 3 (consecutive), keep first (1);
// 2,5 == mod 3 (consecutive), keep first (2)
thrust::device_vector<int> expected_output{0, 1, 2};
thrust::device_vector<int> expected_num_selected{3};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EqualityOpT – [inferred] Type of equality_op

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_data – [inout] Pointer to the sequence of data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • equality_op – [in] Binary equality operator

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename KeyInputIteratorT, typename ValueInputIteratorT, typename KeyOutputIteratorT, typename ValueOutputIteratorT, typename NumSelectedIteratorT, typename NumItemsT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::is_integral_v<NumItemsT> && !::cuda::std::indirect_binary_predicate<EnvT, KeyInputIteratorT, KeyInputIteratorT>, int> = 0>
static inline cudaError_t UniqueByKey(
KeyInputIteratorT d_keys_in,
ValueInputIteratorT d_values_in,
KeyOutputIteratorT d_keys_out,
ValueOutputIteratorT d_values_out,
NumSelectedIteratorT d_num_selected_out,
NumItemsT num_items,
EnvT env = {}
)#

Given an input sequence d_keys_in and d_values_in with runs of key-value pairs with consecutive equal-valued keys, only the first key and its value from each run is selectively copied to d_keys_out and d_values_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The user-provided equality operator, equality_op, is used to determine whether keys are equivalent.

  • Copies of the selected items are compacted into d_keys_out and d_values_out and maintain their original relative ordering.

  • In-place operations are not supported. There must be no overlap between any of the provided ranges.

The code snippet below illustrates the compaction of items selected from an int device vector using environment-based API:

auto keys_in      = thrust::device_vector<int>{0, 2, 2, 9, 5, 5, 5, 8};
auto values_in    = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto keys_out     = thrust::device_vector<int>(5);
auto values_out   = thrust::device_vector<int>(5);
auto num_selected = thrust::device_vector<int>(1);

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::UniqueByKey(
  keys_in.begin(),
  values_in.begin(),
  keys_out.begin(),
  values_out.begin(),
  num_selected.begin(),
  keys_in.size(),
  cuda::std::equal_to<>{},
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::UniqueByKey failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_keys{0, 2, 9, 5, 8};
thrust::device_vector<int> expected_values{1, 2, 4, 5, 8};
thrust::device_vector<int> expected_num_selected{5};
@tparam KeyInputIteratorT

[inferred] Random-access input iterator type for reading input keys (may be a simple pointer type)

@tparam ValueInputIteratorT

[inferred] Random-access input iterator type for reading input values (may be a simple pointer type)

@tparam KeyOutputIteratorT

[inferred] Random-access output iterator type for writing selected keys (may be a simple pointer type)

@tparam ValueOutputIteratorT

[inferred] Random-access output iterator type for writing selected values (may be a simple pointer type)

@tparam NumSelectedIteratorT

[inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

@tparam NumItemsT

[inferred] Type of num_items

@tparam EqualityOpT

[inferred] Type of equality_op

@tparam EnvT

[inferred] Environment type (e.g., cuda::std::execution::env<…>)

@param[in] d_keys_in

Pointer to the input sequence of keys

@param[in] d_values_in

Pointer to the input sequence of values

@param[out] d_keys_out

Pointer to the output sequence of selected keys

@param[out] d_values_out

Pointer to the output sequence of selected values

@param[out] d_num_selected_out

Pointer to the total number of items selected (i.e., length of d_keys_out or d_values_out)

@param[in] num_items

Total number of input items (i.e., length of d_keys_in or d_values_in)

@param[in] equality_op

Binary predicate to determine equality

@param[in] env

[optional] Execution environment. Default is cuda::std::execution::env{}. */

template <typename KeyInputIteratorT,

typename ValueInputIteratorT, typename KeyOutputIteratorT, typename ValueOutputIteratorT, typename NumSelectedIteratorT, typename NumItemsT, typename EqualityOpT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::is_integral_v<NumItemsT>&& ::cuda::std:

  indirect_binary_predicate<EqualityOpT, KeyInputIteratorT, KeyInputIteratorT>,
int> = 0>
[[nodiscard]] static cudaError_t UniqueByKey(

KeyInputIteratorT d_keys_in, ValueInputIteratorT d_values_in, KeyOutputIteratorT d_keys_out, ValueOutputIteratorT d_values_out, NumSelectedIteratorT d_num_selected_out, NumItemsT num_items, EqualityOpT equality_op, EnvT env = {})

{

_CCCL_NVTX_RANGE_SCOPE(“cub::DeviceSelect::UniqueByKey”);

using offset_t = detail::choose_offset_t<NumItemsT>;

using default_policy_selector =
detail::unique_by_key::policy_selector_from_types<detail::it_value_t<KeyInputIteratorT>,

detail::it_value_t<ValueInputIteratorT>>;

return detail::dispatch_with_env_and_tuning<default_policy_selector>(
env, [&](auto policy_selector, void storage, size_t& bytes, auto stream) {
return detail::unique_by_key::dispatch(

storage, bytes, d_keys_in, d_values_in, d_keys_out, d_values_out, d_num_selected_out, equality_op, static_cast<offset_t>(num_items), stream, policy_selector);

});

}

/ * verbatim embed:rst:leading-asterisk

Given an input sequence d_keys_in and d_values_in with runs of key-value pairs with consecutive equal-valued keys, only the first key and its value from each run is selectively copied to d_keys_out and d_values_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

This is an environment-based API that allows customization of:

  • Stream: Query via cuda::get_stream

  • Memory resource: Query via cuda::mr::get_memory_resource

  • The == equality operator is used to determine whether keys are equivalent.

  • Copies of the selected items are compacted into d_keys_out and d_values_out and maintain their original relative ordering.

  • In-place operations are not supported. There must be no overlap between any of the provided ranges.

The code snippet below illustrates the compaction of items selected from an int device vector using environment-based API and default key equality:

// Same setup/expectations as the explicit equality_op test above, but relying on default equality.
auto keys_in      = thrust::device_vector<int>{0, 2, 2, 9, 5, 5, 5, 8};
auto values_in    = thrust::device_vector<int>{1, 2, 3, 4, 5, 6, 7, 8};
auto keys_out     = thrust::device_vector<int>(5);
auto values_out   = thrust::device_vector<int>(5);
auto num_selected = thrust::device_vector<int>(1);

cuda::stream stream{cuda::devices[0]};
cuda::stream_ref stream_ref{stream};

auto error = cub::DeviceSelect::UniqueByKey(
  keys_in.begin(),
  values_in.begin(),
  keys_out.begin(),
  values_out.begin(),
  num_selected.begin(),
  keys_in.size(),
  stream_ref);

if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceSelect::UniqueByKey without equality_op failed with status: " << error << '\n';
}

thrust::device_vector<int> expected_keys{0, 2, 9, 5, 8};
thrust::device_vector<int> expected_values{1, 2, 4, 5, 8};
thrust::device_vector<int> expected_num_selected{5};

Template Parameters:

  • KeyInputIteratorT – [inferred] Random-access input iterator type for reading input keys (may be a simple pointer type)

  • ValueInputIteratorT – [inferred] Random-access input iterator type for reading input values (may be a simple pointer type)

  • KeyOutputIteratorT – [inferred] Random-access output iterator type for writing selected keys (may be a simple pointer type)

  • ValueOutputIteratorT – [inferred] Random-access output iterator type for writing selected values (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • NumItemsT – [inferred] Type of num_items

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_keys_in – [in] Pointer to the input sequence of keys

  • d_values_in – [in] Pointer to the input sequence of values

  • d_keys_out – [out] Pointer to the output sequence of selected keys

  • d_values_out – [out] Pointer to the output sequence of selected values

  • d_num_selected_out – [out] Pointer to the total number of items selected (i.e., length of d_keys_out or d_values_out)

  • num_items – [in] Total number of input items (i.e., length of d_keys_in or d_values_in)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename OutputIteratorT, typename NumSelectedIteratorT, typename EqualityOpT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_binary_predicate<EqualityOpT, InputIteratorT, InputIteratorT>, int> = 0>
static inline cudaError_t Unique(
void *d_temp_storage,
size_t &temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
EqualityOpT equality_op,
const EnvT &env = {}
)#

Given an input sequence d_in having runs of consecutive equal-valued keys, only the first key from each run is selectively copied to d_out. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

  • The user-provided equality operator, equality_op, is used to determine whether keys are equivalent

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>   // or equivalently <cub/device/device_select.cuh>

// Declare, allocate, and initialize device-accessible pointers
// for input and output
int  num_items;              // e.g., 8
int  *d_in;                  // e.g., [0, 2, 2, 9, 5, 5, 5, 8]
int  *d_out;                 // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int  *d_num_selected_out;    // e.g., [ ]
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::Unique(
  d_temp_storage, temp_storage_bytes,
  d_in, d_out, d_num_selected_out, num_items, cuda::std::equal_to<>{});

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::Unique(
  d_temp_storage, temp_storage_bytes,
  d_in, d_out, d_num_selected_out, num_items, cuda::std::equal_to<>{});

// d_out                 <-- [0, 2, 9, 5, 8]
// d_num_selected_out    <-- [5]

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EqualityOpT – [inferred] Type of equality_op

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_in – [in] Pointer to the input sequence of data items

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • equality_op – [in] Binary predicate to determine equality

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename InputIteratorT, typename OutputIteratorT, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<!::cuda::std::indirect_binary_predicate<EnvT, InputIteratorT, InputIteratorT>, int> = 0>
static inline cudaError_t Unique(
void *d_temp_storage,
size_t &temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Given an input sequence d_in having runs of consecutive equal-valued keys, only the first key from each run is selectively copied to d_out. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • The == equality operator is used to determine whether keys are equivalent

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • The range [d_out, d_out + *d_num_selected_out) shall not overlap
    [d_in, d_in + num_items) nor d_num_selected_out in any way.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>   // or equivalently <cub/device/device_select.cuh>

// Declare, allocate, and initialize device-accessible pointers
// for input and output
int  num_items;              // e.g., 8
int  *d_in;                  // e.g., [0, 2, 2, 9, 5, 5, 5, 8]
int  *d_out;                 // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int  *d_num_selected_out;    // e.g., [ ]
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::Unique(
  d_temp_storage, temp_storage_bytes,
  d_in, d_out, d_num_selected_out, num_items);

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::Unique(
  d_temp_storage, temp_storage_bytes,
  d_in, d_out, d_num_selected_out, num_items);

// d_out                 <-- [0, 2, 9, 5, 8]
// d_num_selected_out    <-- [5]

Template Parameters:

  • InputIteratorT – [inferred] Random-access input iterator type for reading input items (may be a simple pointer type)

  • OutputIteratorT – [inferred] Random-access output iterator type for writing selected items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_in – [in] Pointer to the input sequence of data items

  • d_out – [out] Pointer to the output sequence of selected data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected (i.e., length of d_out)

  • num_items – [in] Total number of input items (i.e., length of d_in)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename NumSelectedIteratorT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<!::cuda::std::indirect_binary_predicate<EnvT, IteratorT, IteratorT>, int> = 0>
static inline cudaError_t Unique(
void *d_temp_storage,
size_t &temp_storage_bytes,
IteratorT d_data,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
const EnvT &env = {}
)#

Given an input sequence d_data having runs of consecutive equal-valued keys, only the first key from each run is selectively compacted in-place. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

  • The == equality operator is used to determine whether keys are equivalent

  • Copies of the selected items are compacted in d_data and maintain their original relative ordering.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the in-place compaction of items selected from an int device vector.

constexpr int num_items                       = 8;
thrust::device_vector<int> d_data             = {0, 2, 2, 9, 5, 5, 5, 8};
thrust::device_vector<int> d_num_selected_out = {0};

// Determine temporary device storage requirements
size_t temp_storage_bytes = 0;
cub::DeviceSelect::Unique(nullptr, temp_storage_bytes, d_data.begin(), d_num_selected_out.begin(), num_items);

// Allocate temporary storage
thrust::device_vector<char> temp_storage(temp_storage_bytes, thrust::no_init);

// Run selection
cub::DeviceSelect::Unique(
  thrust::raw_pointer_cast(temp_storage.data()),
  temp_storage_bytes,
  d_data.begin(),
  d_num_selected_out.begin(),
  num_items);

// Resize input to new length
d_data.resize(d_num_selected_out[0]);

thrust::device_vector<int> expected{0, 2, 9, 5, 8};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_data – [inout] Pointer to the sequence of data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename IteratorT, typename NumSelectedIteratorT, typename EqualityOpT, typename EnvT = ::cuda::std::execution::env<>, ::cuda::std::enable_if_t<::cuda::std::indirect_binary_predicate<EqualityOpT, IteratorT, IteratorT>, int> = 0>
static inline cudaError_t Unique(
void *d_temp_storage,
size_t &temp_storage_bytes,
IteratorT d_data,
NumSelectedIteratorT d_num_selected_out,
::cuda::std::int64_t num_items,
EqualityOpT equality_op,
const EnvT &env = {}
)#

Given an input sequence d_data having runs of consecutive equal-valued keys, only the first key from each run is selectively compacted in-place. The total number of items selected is written to d_num_selected_out.

Added in version 3.4.0: First appears in CUDA Toolkit 13.4.

  • The user-provided equality operator, equality_op, is used to determine whether keys are equivalent.

  • Copies of the selected items are compacted in d_data and maintain their original relative ordering.

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the in-place compaction of items selected from an int device vector.

struct my_equality_op
{
  __host__ __device__ bool operator()(int lhs, int rhs) const
  {
    return lhs == rhs;
  }
};

struct my_equality_op
{
  __host__ __device__ bool operator()(int lhs, int rhs) const
  {
    return lhs == rhs;
  }
};

Template Parameters:

  • IteratorT – [inferred] Random-access iterator type for reading and writing items (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • EqualityOpT – [inferred] Type of equality_op

  • EnvT – [inferred] Environment type (e.g., cuda::std::execution::env<...>)

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_data – [inout] Pointer to the sequence of data items

  • d_num_selected_out – [out] Pointer to the output total number of items selected

  • num_items – [in] Total number of input items (i.e., length of d_data)

  • equality_op – [in] Binary predicate to determine equality

  • env – [in] [optional] Execution environment. Default is cuda::std::execution::env{}.

template<typename KeyInputIteratorT, typename ValueInputIteratorT, typename KeyOutputIteratorT, typename ValueOutputIteratorT, typename NumSelectedIteratorT, typename NumItemsT, typename EqualityOpT>
static inline ::cuda::std::enable_if_t<!::cuda::std::is_convertible_v<EqualityOpT, cudaStream_t>, cudaError_t> UniqueByKey(
void *d_temp_storage,
size_t &temp_storage_bytes,
KeyInputIteratorT d_keys_in,
ValueInputIteratorT d_values_in,
KeyOutputIteratorT d_keys_out,
ValueOutputIteratorT d_values_out,
NumSelectedIteratorT d_num_selected_out,
NumItemsT num_items,
EqualityOpT equality_op,
cudaStream_t stream = nullptr
)#

Given an input sequence d_keys_in and d_values_in with runs of key-value pairs with consecutive equal-valued keys, only the first key and its value from each run is selectively copied to d_keys_out and d_values_out. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • The user-provided equality operator, equality_op, is used to determine whether keys are equivalent

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • In-place operations are not supported. There must be no overlap between any of the provided ranges:

    • [d_keys_in,          d_keys_in    + num_items)

    • [d_keys_out,         d_keys_out   + *d_num_selected_out)

    • [d_values_in,        d_values_in  + num_items)

    • [d_values_out,       d_values_out + *d_num_selected_out)

    • [d_num_selected_out, d_num_selected_out + 1)

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>   // or equivalently <cub/device/device_select.cuh>

// Declare, allocate, and initialize device-accessible pointers
// for input and output
int  num_items;              // e.g., 8
int  *d_keys_in;             // e.g., [0, 2, 2, 9, 5, 5, 5, 8]
int  *d_values_in;           // e.g., [1, 2, 3, 4, 5, 6, 7, 8]
int  *d_keys_out;            // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int  *d_values_out;          // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int  *d_num_selected_out;    // e.g., [ ]
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::UniqueByKey(
  d_temp_storage, temp_storage_bytes,
  d_keys_in, d_values_in,
  d_keys_out, d_values_out, d_num_selected_out, num_items);

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::UniqueByKey(
  d_temp_storage, temp_storage_bytes,
  d_keys_in, d_values_in,
  d_keys_out, d_values_out, d_num_selected_out, num_items);

// d_keys_out            <-- [0, 2, 9, 5, 8]
// d_values_out          <-- [1, 2, 4, 5, 8]
// d_num_selected_out    <-- [5]

Template Parameters:

  • KeyInputIteratorT – [inferred] Random-access input iterator type for reading input keys (may be a simple pointer type)

  • ValueInputIteratorT – [inferred] Random-access input iterator type for reading input values (may be a simple pointer type)

  • KeyOutputIteratorT – [inferred] Random-access output iterator type for writing selected keys (may be a simple pointer type)

  • ValueOutputIteratorT – [inferred] Random-access output iterator type for writing selected values (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • NumItemsT – [inferred] Type of num_items

  • EqualityOpT – [inferred] Type of equality_op

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_keys_in – [in] Pointer to the input sequence of keys

  • d_values_in – [in] Pointer to the input sequence of values

  • d_keys_out – [out] Pointer to the output sequence of selected keys

  • d_values_out – [out] Pointer to the output sequence of selected values

  • d_num_selected_out – [out] Pointer to the total number of items selected (i.e., length of d_keys_out or d_values_out)

  • num_items – [in] Total number of input items (i.e., length of d_keys_in or d_values_in)

  • equality_op – [in] Binary predicate to determine equality

  • stream – [in]

    [optional] CUDA stream to launch kernels within. Default is stream0.

template<typename KeyInputIteratorT, typename ValueInputIteratorT, typename KeyOutputIteratorT, typename ValueOutputIteratorT, typename NumSelectedIteratorT, typename NumItemsT>
static inline cudaError_t UniqueByKey(
void *d_temp_storage,
size_t &temp_storage_bytes,
KeyInputIteratorT d_keys_in,
ValueInputIteratorT d_values_in,
KeyOutputIteratorT d_keys_out,
ValueOutputIteratorT d_values_out,
NumSelectedIteratorT d_num_selected_out,
NumItemsT num_items,
cudaStream_t stream = nullptr
)#

Given an input sequence d_keys_in and d_values_in with runs of key-value pairs with consecutive equal-valued keys, only the first key and its value from each run is selectively copied to d_keys_out and d_values_out. The total number of items selected is written to d_num_selected_out.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

  • The == equality operator is used to determine whether keys are equivalent

  • Copies of the selected items are compacted into d_out and maintain their original relative ordering.

  • In-place operations are not supported. There must be no overlap between any of the provided ranges:

    • [d_keys_in,          d_keys_in    + num_items)

    • [d_keys_out,         d_keys_out   + *d_num_selected_out)

    • [d_values_in,        d_values_in  + num_items)

    • [d_values_out,       d_values_out + *d_num_selected_out)

    • [d_num_selected_out, d_num_selected_out + 1)

  • Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See Two-Phase API (explicit temporary storage management) for usage guidance.

Snippet#

The code snippet below illustrates the compaction of items selected from an int device vector.

#include <cub/cub.cuh>   // or equivalently <cub/device/device_select.cuh>

// Declare, allocate, and initialize device-accessible pointers
// for input and output
int  num_items;              // e.g., 8
int  *d_keys_in;             // e.g., [0, 2, 2, 9, 5, 5, 5, 8]
int  *d_values_in;           // e.g., [1, 2, 3, 4, 5, 6, 7, 8]
int  *d_keys_out;            // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int  *d_values_out;          // e.g., [ ,  ,  ,  ,  ,  ,  ,  ]
int  *d_num_selected_out;    // e.g., [ ]
...

// Determine temporary device storage requirements
void     *d_temp_storage = nullptr;
size_t   temp_storage_bytes = 0;
cub::DeviceSelect::UniqueByKey(
  d_temp_storage, temp_storage_bytes,
  d_keys_in, d_values_in,
  d_keys_out, d_values_out, d_num_selected_out, num_items);

// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Run selection
cub::DeviceSelect::UniqueByKey(
  d_temp_storage, temp_storage_bytes,
  d_keys_in, d_values_in,
  d_keys_out, d_values_out, d_num_selected_out, num_items);

// d_keys_out            <-- [0, 2, 9, 5, 8]
// d_values_out          <-- [1, 2, 4, 5, 8]
// d_num_selected_out    <-- [5]

Template Parameters:

  • KeyInputIteratorT – [inferred] Random-access input iterator type for reading input keys (may be a simple pointer type)

  • ValueInputIteratorT – [inferred] Random-access input iterator type for reading input values (may be a simple pointer type)

  • KeyOutputIteratorT – [inferred] Random-access output iterator type for writing selected keys (may be a simple pointer type)

  • ValueOutputIteratorT – [inferred] Random-access output iterator type for writing selected values (may be a simple pointer type)

  • NumSelectedIteratorT – [inferred] Output iterator type for recording the number of items selected (may be a simple pointer type)

  • NumItemsT – [inferred] Type of num_items

Parameters:

  • d_temp_storage – [in] Temporary storage for this operation. If d_temp_storage is nullptr, the required size is written to temp_storage_bytes without dereferencing iterators or launching kernels. Otherwise, d_temp_storage must point to a device-accessible allocation of at least temp_storage_bytes bytes. No special alignment is required. See :ref:device-temp-storage for usage guidance.

  • temp_storage_bytes – [inout] Reference to size in bytes of d_temp_storage allocation

  • d_keys_in – [in] Pointer to the input sequence of keys

  • d_values_in – [in] Pointer to the input sequence of values

  • d_keys_out – [out] Pointer to the output sequence of selected keys

  • d_values_out – [out] Pointer to the output sequence of selected values

  • d_num_selected_out – [out] Pointer to the total number of items selected (i.e., length of d_keys_out or d_values_out)

  • num_items – [in] Total number of input items (i.e., length of d_keys_in or d_values_in)

  • stream – [in]

    [optional] CUDA stream to launch kernels within. Default is stream0.