cub::DeviceFind#

struct DeviceFind#

Tuning#

The FindIf algorithms that accept an environment can be tuned by passing a custom policy selector that returns a @ref FindIfPolicy, as shown in the example below:

struct FindPolicySelector
{
  __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::FindIfPolicy
  {
    return {.threads_per_block = 128,
            .items_per_thread  = cc > cuda::compute_capability{9, 0} ? 16 : 7,
            .vec_size          = 4,
            .load_modifier     = cub::LOAD_LDG};
  }
};

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

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

int expected = 5;

The LowerBoundSortedValues and UpperBoundSortedValues algorithms that accept an environment can be tuned by passing a custom policy selector that returns a @ref FindBoundSortedValuesPolicy, as shown in the example below:

struct FindBoundSortedValuesPolicySelector
{
  __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::FindBoundSortedValuesPolicy
  {
    return {.threads_per_block = cc >= cuda::compute_capability{8, 0} ? 512 : 256,
            .items_per_thread  = 7,
            .load_modifier     = cub::LOAD_DEFAULT};
  }
};

thrust::device_vector<int> d_range  = {0, 2, 4, 6, 8};
thrust::device_vector<int> d_values = {0, 3, 4, 7};
thrust::device_vector<int> d_output(4, thrust::no_init);

auto error = cub::DeviceFind::LowerBoundSortedValues(
  d_range.begin(),
  d_range.size(),
  d_values.begin(),
  d_values.size(),
  d_output.begin(),
  cuda::std::less{},
  cuda::execution::tune(FindBoundSortedValuesPolicySelector{}));
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceFind::LowerBoundSortedValues failed with status: " << error << '\n';
}

thrust::device_vector<int> expected = {0, 2, 2, 4};

Public Static Functions

template<typename InputIteratorT, typename OutputIteratorT, typename ScanOpT, typename NumItemsT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t FindIf(
void *d_temp_storage,
size_t &temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
NumItemsT num_items,
const EnvT &env = {}
)#

Finds the first element in the input sequence that satisfies the given predicate.

  • The search terminates at the first element where the predicate evaluates to true.

  • The index of the found element is written to d_out.

  • If no element satisfies the predicate, num_items is written to d_out.

  • The range [d_out, d_out + 1) shall not overlap [d_in, d_in + num_items) 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.

Added in version 3.3.0.

Snippet#

The code snippet below illustrates the finding of the first element that satisfies the predicate.

struct is_greater_than_t
{
  int threshold;
  __host__ __device__ bool operator()(int value) const
  {
    return value > threshold;
  }
};

constexpr int num_items         = 8;
thrust::device_vector<int> d_in = {0, 1, 2, 3, 4, 5, 6, 7};
thrust::device_vector<int> d_out(1, thrust::no_init);
is_greater_than_t predicate{4};

size_t temp_storage_bytes = 0;
cub::DeviceFind::FindIf(nullptr, temp_storage_bytes, d_in.begin(), d_out.begin(), predicate, num_items);

thrust::device_vector<char> temp_storage(temp_storage_bytes, thrust::no_init);

cub::DeviceFind::FindIf(
  thrust::raw_pointer_cast(temp_storage.data()),
  temp_storage_bytes,
  d_in.begin(),
  d_out.begin(),
  predicate,
  num_items);

int expected = 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 the result index (may be a simple pointer type)

  • ScanOpT – [inferred] Unary predicate functor type having member bool operator()(const T &a)

  • NumItemsT – [inferred] An integral type representing the number of input elements

  • 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] Random-access iterator to the input sequence of data items

  • d_out – [out] Random-access iterator to the output location for the index of the found element

  • scan_op – [in] Unary predicate functor for determining whether an element satisfies the search condition

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

  • env – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t LowerBound(
void *d_temp_storage,
size_t &temp_storage_bytes,
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
const EnvT &env = {}
)#

Overview#

For each value in [d_values, d_values + values_num_items), performs a binary search in the range [d_range, d_range + range_num_items), using comp as the comparator to find the iterator to the first element of said range which is not ordered before value.

  • The range [d_range, d_range + range_num_items) must be sorted consistently with comp.

Added in version 3.3.0.

Snippet#

The code snippet below illustrates the lower bound search.

thrust::device_vector<int> d_range  = {0, 2, 4, 6, 8};
thrust::device_vector<int> d_values = {1, 3, 5, 7};
thrust::device_vector<int> d_output(4);

size_t temp_storage_bytes = 0;
cub::DeviceFind::LowerBound(
  nullptr,
  temp_storage_bytes,
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{});

thrust::device_vector<char> temp_storage(temp_storage_bytes, thrust::no_init);

cub::DeviceFind::LowerBound(
  thrust::raw_pointer_cast(temp_storage.data()),
  temp_storage_bytes,
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{});

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

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT using CompareOpT as the predicate.

  • RangeNumItemsT – is an integral type representing the number of elements in the range to be searched.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT using CompareOpT as the predicate.

  • ValuesNumItemsT – is a model of integral type representing the number of elements in the range of values to be searched for.

  • OutputIteratorT – is a model of Random Access Iterator, whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering, which forms a Relation with the value types of RangeIteratorT and ValuesIteratorT.

  • 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_range – [in] Iterator to the beginning of the ordered range to be searched.

  • range_num_items – [in] Number of elements in the ordered range to be searched.

  • d_values – [in] Iterator to the beginning of the range of values to be searched for.

  • values_num_items – [in] Number of elements in the range of values to be searched for.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object which returns true if its first argument is ordered before the second in the Strict Weak Ordering of the range to be searched.

  • env – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t UpperBound(
void *d_temp_storage,
size_t &temp_storage_bytes,
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
const EnvT &env = {}
)#

Overview#

For each value in [d_values, d_values + values_num_items), performs a binary search in the range [d_range, d_range + range_num_items), using comp as the comparator to find the iterator to the first element of said range which is ordered after value.

  • The range [d_range, d_range + range_num_items) must be sorted consistently with comp.

Added in version 3.3.0.

Snippet#

The code snippet below illustrates the upper bound search.

thrust::device_vector<int> d_range  = {0, 2, 4, 6, 8};
thrust::device_vector<int> d_values = {1, 3, 5, 7};
thrust::device_vector<int> d_output(4);

size_t temp_storage_bytes = 0;
cub::DeviceFind::UpperBound(
  nullptr,
  temp_storage_bytes,
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{});

thrust::device_vector<char> temp_storage(temp_storage_bytes, thrust::no_init);

cub::DeviceFind::UpperBound(
  thrust::raw_pointer_cast(temp_storage.data()),
  temp_storage_bytes,
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{});

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

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT using CompareOpT as the predicate.

  • RangeNumItemsT – is an integral type representing the number of elements in the range to be searched.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT using CompareOpT as the predicate.

  • ValuesNumItemsT – is a model of integral type representing the number of elements in the range of values to be searched for.

  • OutputIteratorT – is a model of Random Access Iterator, whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering, which forms a Relation with the value types of RangeIteratorT and ValuesIteratorT.

  • 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_range – [in] Iterator to the beginning of the ordered range to be searched.

  • range_num_items – [in] Number of elements in the ordered range to be searched.

  • d_values – [in] Iterator to the beginning of the range of values to be searched for.

  • values_num_items – [in] Number of elements in the range of values to be searched for.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object which returns true if its first argument is ordered before the second in the Strict Weak Ordering of the range to be searched.

  • env – [in]

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

template<typename InputIteratorT, typename OutputIteratorT, typename ScanOpT, typename NumItemsT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t FindIf(
InputIteratorT d_in,
OutputIteratorT d_out,
ScanOpT scan_op,
NumItemsT num_items,
const EnvT &env = {}
)#

Finds the first element in the input sequence that satisfies the given predicate.

  • The search terminates at the first element where the predicate evaluates to true.

  • The index of the found element is written to d_out.

  • If no element satisfies the predicate, num_items is written to d_out.

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

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

*/

/ * Snippet

The code snippet below illustrates the finding of the first element that satisfies the predicate.

struct is_greater_than_t
{
  int threshold;
  __host__ __device__ bool operator()(int value) const
  {
    return value > threshold;
  }
};

constexpr int num_items         = 8;
thrust::device_vector<int> d_in = {0, 1, 2, 3, 4, 5, 6, 7};
thrust::device_vector<int> d_out(1);
is_greater_than_t predicate{4};

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

auto error = cub::DeviceFind::FindIf(d_in.begin(), d_out.begin(), predicate, num_items, stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceFind::FindIf failed with status: " << error << '\n';
}

int expected = 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 the result index (may be a simple pointer type)

  • ScanOpT – [inferred] Unary predicate functor type having member bool operator()(const T &a)

  • NumItemsT – [inferred] An integral type representing the number of input elements

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

Parameters:

  • d_in – [in] Random-access iterator to the input sequence of data items

  • d_out – [out] Random-access iterator to the output location for the index of the found element

  • scan_op – [in] Unary predicate functor for determining whether an element satisfies the search condition

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

  • env – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t LowerBound(
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
const EnvT &env = {}
)#

For each value in [d_values, d_values + values_num_items), performs a binary search in the range [d_range, d_range + range_num_items), using comp as the comparator to find the iterator to the first element of said range which is not ordered before value.

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 range [d_range, d_range + range_num_items) must be sorted consistently with comp.

Snippet#

The code snippet below illustrates the lower bound search.

thrust::device_vector<int> d_range  = {0, 2, 4, 6, 8};
thrust::device_vector<int> d_values = {1, 3, 5, 7};
thrust::device_vector<int> d_output(4);

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

auto error = cub::DeviceFind::LowerBound(
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{},
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceFind::LowerBound failed with status: " << error << '\n';
}

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

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT using CompareOpT as the predicate.

  • RangeNumItemsT – is an integral type representing the number of elements in the range to be searched.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT using CompareOpT as the predicate.

  • ValuesNumItemsT – is a model of integral type representing the number of elements in the range of values to be searched for.

  • OutputIteratorT – is a model of Random Access Iterator, whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering, which forms a Relation with the value types of RangeIteratorT and ValuesIteratorT.

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

Parameters:

  • d_range – [in] Iterator to the beginning of the ordered range to be searched.

  • range_num_items – [in] Number of elements in the ordered range to be searched.

  • d_values – [in] Iterator to the beginning of the range of values to be searched for.

  • values_num_items – [in] Number of elements in the range of values to be searched for.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object which returns true if its first argument is ordered before the second in the Strict Weak Ordering of the range to be searched.

  • env – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t UpperBound(
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
const EnvT &env = {}
)#

For each value in [d_values, d_values + values_num_items), performs a binary search in the range [d_range, d_range + range_num_items), using comp as the comparator to find the iterator to the first element of said range which is ordered after value.

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 range [d_range, d_range + range_num_items) must be sorted consistently with comp.

Snippet#

The code snippet below illustrates the upper bound search.

thrust::device_vector<int> d_range  = {0, 2, 4, 6, 8};
thrust::device_vector<int> d_values = {1, 3, 5, 7};
thrust::device_vector<int> d_output(4);

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

auto error = cub::DeviceFind::UpperBound(
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{},
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceFind::UpperBound failed with status: " << error << '\n';
}

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

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT using CompareOpT as the predicate.

  • RangeNumItemsT – is an integral type representing the number of elements in the range to be searched.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT using CompareOpT as the predicate.

  • ValuesNumItemsT – is a model of integral type representing the number of elements in the range of values to be searched for.

  • OutputIteratorT – is a model of Random Access Iterator, whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering, which forms a Relation with the value types of RangeIteratorT and ValuesIteratorT.

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

Parameters:

  • d_range – [in] Iterator to the beginning of the ordered range to be searched.

  • range_num_items – [in] Number of elements in the ordered range to be searched.

  • d_values – [in] Iterator to the beginning of the range of values to be searched for.

  • values_num_items – [in] Number of elements in the range of values to be searched for.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object which returns true if its first argument is ordered before the second in the Strict Weak Ordering of the range to be searched.

  • env – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT>
static inline cudaError_t LowerBoundSortedValues(
void *d_temp_storage,
size_t &temp_storage_bytes,
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
cudaStream_t stream = nullptr
)#

Overview#

Accelerated variant of LowerBound() that exploits the additional precondition that [d_values, d_values + values_num_items) is also sorted consistently with comp.

For each value in [d_values, d_values + values_num_items), performs a search in [d_range, d_range + range_num_items) to find the iterator to the first element that is not ordered before value.

Because both sequences are sorted, the algorithm uses the Merge-Path algorithm (Oded et al., IPDPS 2012) to partition the combined traversal across thread blocks, achieving O(N+M) total device work rather than the O(M log N) of independent binary searches.

  • Both [d_range, d_range + range_num_items) and [d_values, d_values + values_num_items) must be sorted consistently with comp.

  • 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.

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

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT via CompareOpT.

  • RangeNumItemsT – is an integral type representing the number of elements in the range.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT via CompareOpT.

  • ValuesNumItemsT – is an integral type representing the number of values to search for.

  • OutputIteratorT – is a model of Random Access Iterator whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering over the value types of both iterator types.

Parameters:

  • d_temp_storage – [in] Device-accessible allocation of temporary storage. When nullptr, the required allocation size is written to temp_storage_bytes and no work is done.

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

  • d_range – [in] Iterator to the beginning of the ordered haystack range.

  • range_num_items – [in] Number of elements in the haystack range.

  • d_values – [in] Iterator to the beginning of the sorted range of needles.

  • values_num_items – [in] Number of needle elements.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object (Strict Weak Ordering).

  • stream – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t LowerBoundSortedValues(
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
EnvT env = {}
)#

Accelerated variant of LowerBound() that exploits the additional precondition that [d_values, d_values + values_num_items) is also sorted consistently with comp.

Added in version 3.5.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

  • Both [d_range, d_range + range_num_items) and [d_values, d_values + values_num_items) must be sorted consistently with comp.

Snippet#

The code snippet below illustrates the lower bound search.

thrust::device_vector<int> d_range  = {0, 2, 4, 6, 8};
thrust::device_vector<int> d_values = {0, 3, 4, 7};
thrust::device_vector<int> d_output(4);

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

auto error = cub::DeviceFind::LowerBoundSortedValues(
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{},
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceFind::LowerBoundSortedValues failed with status: " << error << '\n';
}

thrust::device_vector<int> expected = {0, 2, 2, 4};

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT using CompareOpT as the predicate.

  • RangeNumItemsT – is an integral type representing the number of elements in the range to be searched.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT using CompareOpT as the predicate.

  • ValuesNumItemsT – is a model of integral type representing the number of elements in the range of values to be searched for.

  • OutputIteratorT – is a model of Random Access Iterator, whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering, which forms a Relation with the value types of RangeIteratorT and ValuesIteratorT.

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

Parameters:

  • d_range – [in] Iterator to the beginning of the ordered haystack range.

  • range_num_items – [in] Number of elements in the haystack range.

  • d_values – [in] Iterator to the beginning of the sorted range of needles.

  • values_num_items – [in] Number of needle elements.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object (Strict Weak Ordering).

  • env – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT>
static inline cudaError_t UpperBoundSortedValues(
void *d_temp_storage,
size_t &temp_storage_bytes,
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
cudaStream_t stream = nullptr
)#

Overview#

Accelerated variant of UpperBound() that exploits the additional precondition that [d_values, d_values + values_num_items) is also sorted consistently with comp.

For each value in [d_values, d_values + values_num_items), performs a search in [d_range, d_range + range_num_items) to find the iterator to the first element that is ordered after value.

Because both sequences are sorted, the algorithm uses the Merge-Path algorithm (Oded et al., IPDPS 2012) to partition the combined traversal across thread blocks, achieving O(N+M) total device work rather than the O(M log N) of independent binary searches.

  • Both [d_range, d_range + range_num_items) and [d_values, d_values + values_num_items) must be sorted consistently with comp.

  • 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.

Added in version 3.5.0.

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT via CompareOpT.

  • RangeNumItemsT – is an integral type representing the number of elements in the range.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT via CompareOpT.

  • ValuesNumItemsT – is an integral type representing the number of values to search for.

  • OutputIteratorT – is a model of Random Access Iterator whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering over the value types of both iterator types.

Parameters:

  • d_temp_storage – [in] Device-accessible allocation of temporary storage. When nullptr, the required allocation size is written to temp_storage_bytes and no work is done.

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

  • d_range – [in] Iterator to the beginning of the ordered haystack range.

  • range_num_items – [in] Number of elements in the haystack range.

  • d_values – [in] Iterator to the beginning of the sorted range of needles.

  • values_num_items – [in] Number of needle elements.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object (Strict Weak Ordering).

  • stream – [in]

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

template<typename RangeIteratorT, typename RangeNumItemsT, typename ValuesIteratorT, typename ValuesNumItemsT, typename OutputIteratorT, typename CompareOpT, typename EnvT = ::cuda::std::execution::env<>>
static inline cudaError_t UpperBoundSortedValues(
RangeIteratorT d_range,
RangeNumItemsT range_num_items,
ValuesIteratorT d_values,
ValuesNumItemsT values_num_items,
OutputIteratorT d_output,
CompareOpT comp,
EnvT env = {}
)#

Accelerated variant of UpperBound() that exploits the additional precondition that [d_values, d_values + values_num_items) is also sorted consistently with comp.

Added in version 3.5.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

  • Both [d_range, d_range + range_num_items) and [d_values, d_values + values_num_items) must be sorted consistently with comp.

Snippet#

The code snippet below illustrates the upper bound search.

thrust::device_vector<int> d_range  = {0, 2, 4, 6, 8};
thrust::device_vector<int> d_values = {0, 3, 4, 7};
thrust::device_vector<int> d_output(4);

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

auto error = cub::DeviceFind::UpperBoundSortedValues(
  d_range.begin(),
  static_cast<int>(d_range.size()),
  d_values.begin(),
  static_cast<int>(d_values.size()),
  d_output.begin(),
  cuda::std::less{},
  stream_ref);
if (error != cudaSuccess)
{
  std::cerr << "cub::DeviceFind::UpperBoundSortedValues failed with status: " << error << '\n';
}

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

Template Parameters:

  • RangeIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of ValuesIteratorT using CompareOpT as the predicate.

  • RangeNumItemsT – is an integral type representing the number of elements in the range to be searched.

  • ValuesIteratorT – is a model of Random Access Iterator, whose value type forms a Relation with the value type of RangeIteratorT using CompareOpT as the predicate.

  • ValuesNumItemsT – is a model of integral type representing the number of elements in the range of values to be searched for.

  • OutputIteratorT – is a model of Random Access Iterator, whose value type is assignable from RangeIteratorT’s difference type.

  • CompareOpT – is a model of Strict Weak Ordering, which forms a Relation with the value types of RangeIteratorT and ValuesIteratorT.

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

Parameters:

  • d_range – [in] Iterator to the beginning of the ordered haystack range.

  • range_num_items – [in] Number of elements in the haystack range.

  • d_values – [in] Iterator to the beginning of the sorted range of needles.

  • values_num_items – [in] Number of needle elements.

  • d_output – [out] Iterator to the beginning of the output range.

  • comp – [in] Comparison function object (Strict Weak Ordering).

  • env – [in]

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