API Reference Guide#

This chapter describes the rocThrust C++ API.

Memory Management#

group memory_management

All Thrust functionalities related to memory allocation and deallocation.

Functions

template<typename T>
inline void device_delete(thrust::device_ptr<T> ptr, const size_t n = 1)#

device_delete deletes a device_ptr allocated with device_new.

See also

device_ptr

See also

device_new

Parameters
  • ptr – The device_ptr to delete, assumed to have been allocated with device_new.

  • n – The number of objects to destroy at ptr. Defaults to 1 similar to device_new.

inline void device_free(thrust::device_ptr<void> ptr)#

device_free deallocates memory allocated by the function device_malloc.

The following code snippet demonstrates how to use device_free to deallocate memory allocated by device_malloc.

#include <thrust/device_malloc.h>
#include <thrust/device_free.h>
...
// allocate some integers with device_malloc
const int N = 100;
thrust::device_ptr<int> int_array = thrust::device_malloc<int>(N);

// manipulate integers
...

// deallocate with device_free
thrust::device_free(int_array);

See also

device_ptr

See also

device_malloc

Parameters

ptr – A device_ptr pointing to memory to be deallocated.

template<typename T>
inline thrust::device_ptr<T> device_malloc(const std::size_t n)#

This version of device_malloc allocates sequential device storage for new objects of the given type.

The following code snippet demonstrates how to use device_malloc to allocate a range of device memory.

#include <thrust/device_malloc.h>
#include <thrust/device_free.h>
...
// allocate some integers with device_malloc
const int N = 100;
thrust::device_ptr<int> int_array = thrust::device_malloc<int>(N);

// manipulate integers
...

// deallocate with device_free
thrust::device_free(int_array);

See also

device_ptr

See also

device_free

Parameters

n – The number of objects of type T to allocate sequentially in device memory.

Returns

A device_ptr to the newly allocated memory.

inline thrust::device_ptr<void> device_malloc(const std::size_t n)#

This version of device_malloc allocates untyped sequential device storage.

The following code snippet demonstrates how to use device_malloc to allocate a range of device memory.

#include <thrust/device_malloc.h>
#include <thrust/device_free.h>
...
// allocate some memory with device_malloc
const int N = 100;
thrust::device_ptr<void> void_ptr = thrust::device_malloc(N);

// manipulate memory
...

// deallocate with device_free
thrust::device_free(void_ptr);

See also

device_ptr

See also

device_free

Parameters

n – The number of bytes to allocate sequentially in device memory.

Returns

A device_ptr to the newly allocated memory.

template<typename T>
device_ptr<T> device_new(device_ptr<void> p, const size_t n = 1)#

device_new implements the placement new operator for types resident in device memory. device_new calls T’s null constructor on a array of objects in device memory. No memory is allocated by this function.

See also

device_ptr

Parameters
  • p – A device_ptr to a region of device memory into which to construct one or many Ts.

  • n – The number of objects to construct at p.

Returns

p, casted to T’s type.

template<typename T>
device_ptr<T> device_new(device_ptr<void> p, const T &exemplar, const size_t n = 1)#

device_new implements the placement new operator for types resident in device memory. device_new calls T’s copy constructor on a array of objects in device memory. No memory is allocated by this function.

See also

device_ptr

See also

fill

Parameters
  • p – A device_ptr to a region of device memory into which to construct one or many Ts.

  • exemplar – The value from which to copy.

  • n – The number of objects to construct at p.

Returns

p, casted to T’s type.

template<typename T>
device_ptr<T> device_new(const size_t n = 1)#

device_new implements the new operator for types resident in device memory. It allocates device memory large enough to hold n new objects of type T.

Parameters

n – The number of objects to allocate. Defaults to 1.

Returns

A device_ptr to the newly allocated region of device memory.

template<typename T> __host__ __device__ device_ptr< T > device_pointer_cast (T *ptr)

Create a device_ptr from a raw pointer.

Template Parameters

T – Any type.

Parameters

ptr – A raw pointer to a T in device memory.

Pre

ptr points to a location in device memory.

Returns

A device_ptr<T> pointing to ptr.

template<typename T> __host__ __device__ device_ptr< T > device_pointer_cast (device_ptr< T > const &dptr)

Create a device_ptr from another device_ptr.

Template Parameters

T – Any type.

Parameters

dptr – A device_ptr to a T.

template<typename T> __host__ __device__ void swap (device_reference< T > &x, device_reference< T > &y)

swaps the value of one device_reference with another. x The first device_reference of interest. y The second device_reference of interest.

template<typename DerivedPolicy> __host__ __device__ pointer< void, DerivedPolicy > malloc (const thrust::detail::execution_policy_base< DerivedPolicy > &system, std::size_t n)

This version of malloc allocates untyped uninitialized storage associated with a given system.

The following code snippet demonstrates how to use malloc to allocate a range of memory associated with Thrust’s device system.

#include <thrust/memory.h>
...
// allocate some memory with thrust::malloc
const int N = 100;
thrust::device_system_tag device_sys;
thrust::pointer<void,thrust::device_space_tag> void_ptr = thrust::malloc(device_sys, N);

// manipulate memory
...

// deallocate void_ptr with thrust::free
thrust::free(device_sys, void_ptr);

See also

free

See also

device_malloc

Parameters
  • system – The Thrust system with which to associate the storage.

  • n – The number of bytes of storage to allocate.

Template Parameters

DerivedPolicy – The name of the derived execution policy.

Returns

If allocation succeeds, a pointer to the allocated storage; a null pointer otherwise. The pointer must be deallocated with thrust::free.

Pre

DerivedPolicy must be publically derived from thrust::execution_policy<DerivedPolicy>.

template<typename T, typename DerivedPolicy> __host__ __device__ pointer< T, DerivedPolicy > malloc (const thrust::detail::execution_policy_base< DerivedPolicy > &system, std::size_t n)

This version of malloc allocates typed uninitialized storage associated with a given system.

The following code snippet demonstrates how to use malloc to allocate a range of memory to accomodate integers associated with Thrust’s device system.

#include <thrust/memory.h>
...
// allocate storage for 100 ints with thrust::malloc
const int N = 100;
thrust::device_system_tag device_sys;
thrust::pointer<int,thrust::device_system_tag> ptr = thrust::malloc<int>(device_sys, N);

// manipulate memory
...

// deallocate ptr with thrust::free
thrust::free(device_sys, ptr);

See also

free

See also

device_malloc

Parameters
  • system – The Thrust system with which to associate the storage.

  • n – The number of elements of type T which the storage should accomodate.

Template Parameters

DerivedPolicy – The name of the derived execution policy.

Returns

If allocation succeeds, a pointer to an allocation large enough to accomodate n elements of type T; a null pointer otherwise. The pointer must be deallocated with thrust::free.

Pre

DerivedPolicy must be publically derived from thrust::execution_policy<DerivedPolicy>.

template<typename T, typename DerivedPolicy> __host__ __device__ thrust::pair< thrust::pointer< T, DerivedPolicy >, typename thrust::pointer< T, DerivedPolicy >::difference_type > get_temporary_buffer (const thrust::detail::execution_policy_base< DerivedPolicy > &system, typename thrust::pointer< T, DerivedPolicy >::difference_type n)

get_temporary_buffer returns a pointer to storage associated with a given Thrust system sufficient to store up to n objects of type T. If not enough storage is available to accomodate n objects, an implementation may return a smaller buffer. The number of objects the returned buffer can accomodate is also returned.

Thrust uses get_temporary_buffer internally when allocating temporary storage required by algorithm implementations.

The storage allocated with get_temporary_buffer must be returned to the system with return_temporary_buffer.

The following code snippet demonstrates how to use get_temporary_buffer to allocate a range of memory to accomodate integers associated with Thrust’s device system.

#include <thrust/memory.h>
...
// allocate storage for 100 ints with thrust::get_temporary_buffer
const int N = 100;

typedef thrust::pair<
  thrust::pointer<int,thrust::device_system_tag>,
  std::ptrdiff_t
> ptr_and_size_t;

thrust::device_system_tag device_sys;
ptr_and_size_t ptr_and_size = thrust::get_temporary_buffer<int>(device_sys, N);

// manipulate up to 100 ints
for(int i = 0; i < ptr_and_size.second; ++i)
{
  *ptr_and_size.first = i;
}

// deallocate storage with thrust::return_temporary_buffer
thrust::return_temporary_buffer(device_sys, ptr_and_size.first);

See also

malloc

Parameters
  • system – The Thrust system with which to associate the storage.

  • n – The requested number of objects of type T the storage should accomodate.

Template Parameters

DerivedPolicy – The name of the derived execution policy.

Returns

A pair p such that p.first is a pointer to the allocated storage and p.second is the number of contiguous objects of type T that the storage can accomodate. If no storage can be allocated, p.first if no storage can be obtained. The storage must be returned to the system using return_temporary_buffer.

Pre

DerivedPolicy must be publically derived from thrust::execution_policy<DerivedPolicy>.

template<typename DerivedPolicy, typename Pointer> __host__ __device__ void free (const thrust::detail::execution_policy_base< DerivedPolicy > &system, Pointer ptr)

free deallocates the storage previously allocated by thrust::malloc.

The following code snippet demonstrates how to use free to deallocate a range of memory previously allocated with thrust::malloc.

#include <thrust/memory.h>
...
// allocate storage for 100 ints with thrust::malloc
const int N = 100;
thrust::device_system_tag device_sys;
thrust::pointer<int,thrust::device_system_tag> ptr = thrust::malloc<int>(device_sys, N);

// mainpulate memory
...

// deallocate ptr with thrust::free
thrust::free(device_sys, ptr);
Parameters
  • system – The Thrust system with which the storage is associated.

  • ptr – A pointer previously returned by thrust::malloc. If ptr is null, free does nothing.

Template Parameters

DerivedPolicy – The name of the derived execution policy.

Pre

ptr shall have been returned by a previous call to thrust::malloc(system, n) or thrust::malloc<T>(system, n) for some type T.

template<typename DerivedPolicy, typename Pointer> __host__ __device__ void return_temporary_buffer (const thrust::detail::execution_policy_base< DerivedPolicy > &system, Pointer p, std::ptrdiff_t n)

return_temporary_buffer deallocates storage associated with a given Thrust system previously allocated by get_temporary_buffer.

Thrust uses return_temporary_buffer internally when deallocating temporary storage required by algorithm implementations.

The following code snippet demonstrates how to use return_temporary_buffer to deallocate a range of memory previously allocated by get_temporary_buffer.

#include <thrust/memory.h>
...
// allocate storage for 100 ints with thrust::get_temporary_buffer
const int N = 100;

typedef thrust::pair<
  thrust::pointer<int,thrust::device_system_tag>,
  std::ptrdiff_t
> ptr_and_size_t;

thrust::device_system_tag device_sys;
ptr_and_size_t ptr_and_size = thrust::get_temporary_buffer<int>(device_sys, N);

// manipulate up to 100 ints
for(int i = 0; i < ptr_and_size.second; ++i)
{
  *ptr_and_size.first = i;
}

// deallocate storage with thrust::return_temporary_buffer
thrust::return_temporary_buffer(device_sys, ptr_and_size.first);

See also

free

Parameters
  • system – The Thrust system with which the storage is associated.

  • p – A pointer previously returned by thrust::get_temporary_buffer. If ptr is null, return_temporary_buffer does nothing.

  • n

Template Parameters

DerivedPolicy – The name of the derived execution policy.

Pre

p shall have been previously allocated by thrust::get_temporary_buffer.

template<typename Pointer> __host__ __device__ thrust::detail::pointer_traits< Pointer >::raw_pointer raw_pointer_cast (Pointer ptr)

raw_pointer_cast creates a “raw” pointer from a pointer-like type, simply returning the wrapped pointer, should it exist.

Parameters

ptr – The pointer of interest.

Returns

ptr.get(), if the expression is well formed; ptr, otherwise.

template<typename T> __host__ __device__ detail::raw_reference< T >::type raw_reference_cast (T &ref)

raw_reference_cast creates a “raw” reference from a wrapped reference type, simply returning the underlying reference, should it exist.

If the argument is not a reference wrapper, the result is a reference to the argument.

See also

raw_pointer_cast

Note

There are two versions of raw_reference_cast. One for const references, and one for non-const.

Parameters

ref – The reference of interest.

Returns

*raw_pointer_cast(&ref).

template<typename T> __host__ __device__ detail::raw_reference< constT >::type raw_reference_cast (const T &ref)

raw_reference_cast creates a “raw” reference from a wrapped reference type, simply returning the underlying reference, should it exist.

If the argument is not a reference wrapper, the result is a reference to the argument.

See also

raw_pointer_cast

Note

There are two versions of raw_reference_cast. One for const references, and one for non-const.

Parameters

ref – The reference of interest.

Returns

*raw_reference_cast(&ref).

template<typename T>
class device_reference : public thrust::reference<T, thrust::device_ptr<T>, thrust::device_reference<T>>#
#include <device_reference.h>

device_reference acts as a reference-like object to an object stored in device memory. device_reference is not intended to be used directly; rather, this type is the result of deferencing a device_ptr. Similarly, taking the address of a device_reference yields a device_ptr.

device_reference may often be used from host code in place of operations defined on its associated value_type. For example, when device_reference refers to an arithmetic type, arithmetic operations on it are legal:

#include <thrust/device_vector.h>

int main(void)
{
  thrust::device_vector<int> vec(1, 13);

  thrust::device_reference<int> ref_to_thirteen = vec[0];

  int x = ref_to_thirteen + 1;

  // x is 14

  return 0;
}

Similarly, we can print the value of ref_to_thirteen in the above code by using an iostream:

#include <thrust/device_vector.h>
#include <iostream>

int main(void)
{
  thrust::device_vector<int> vec(1, 13);

  thrust::device_reference<int> ref_to_thirteen = vec[0];

  std::cout << ref_to_thirteen << std::endl;

  // 13 is printed

  return 0;
}

Of course, we needn’t explicitly create a device_reference in the previous example, because one is returned by device_vector's bracket operator. A more natural way to print the value of a device_vector element might be:

#include <thrust/device_vector.h>
#include <iostream>

int main(void)
{
  thrust::device_vector<int> vec(1, 13);

  std::cout << vec[0] << std::endl;

  // 13 is printed

  return 0;
}

These kinds of operations should be used sparingly in performance-critical code, because they imply a potentially expensive copy between host and device space.

Some operations which are possible with regular objects are impossible with their corresponding device_reference objects due to the requirements of the C++ language. For example, because the member access operator cannot be overloaded, member variables and functions of a referent object cannot be directly accessed through its device_reference.

The following code, which generates a compiler error, illustrates:

#include <thrust/device_vector.h>

struct foo
{
  int x;
};

int main(void)
{
  thrust::device_vector<foo> foo_vec(1);

  thrust::device_reference<foo> foo_ref = foo_vec[0];

  foo_ref.x = 13; // ERROR: x cannot be accessed through foo_ref

  return 0;
}

Instead, a host space copy must be created to access foo's x member:

#include <thrust/device_vector.h>

struct foo
{
  int x;
};

int main(void)
{
  thrust::device_vector<foo> foo_vec(1);

  // create a local host-side foo object
  foo host_foo;
  host_foo.x = 13;

  thrust::device_reference<foo> foo_ref = foo_vec[0];

  foo_ref = host_foo;

  // foo_ref's x member is 13

  return 0;
}

Another common case where a device_reference cannot directly be used in place of its referent object occurs when passing them as parameters to functions like printf which have varargs parameters. Because varargs parameters must be Plain Old Data, a device_reference to a POD type requires a cast when passed to printf:

#include <stdio.h>
#include <thrust/device_vector.h>

int main(void)
{
  thrust::device_vector<int> vec(1,13);

  // vec[0] must be cast to int when passing to printf
  printf("%d\n", (int) vec[0]);

  return 0;
}

See also

device_ptr

See also

device_vector

template<typename T>
class device_ptr : public thrust::pointer<T, thrust::device_system_tag, thrust::device_reference<T>, thrust::device_ptr<T>>#
#include <device_ptr.h>

forward declaration to WAR circular #includes

device_ptr is a pointer-like object which points to an object that resides in memory associated with the device system.

device_ptr has pointer semantics: it may be dereferenced safely from anywhere, including the host, and may be manipulated with pointer arithmetic.

device_ptr can be created with device_new, device_malloc, device_malloc_allocator, device_allocator, or device_pointer_cast, or by explicitly calling its constructor with a raw pointer.

The raw pointer contained in a device_ptr may be obtained via get member function or the raw_pointer_cast free function.

Algorithms operating on device_ptr types will automatically be dispatched to the device system.

See also

device_new

See also

device_malloc

See also

device_allocator

See also

raw_pointer_cast

Note

device_ptr is not a smart pointer; it is the programmer’s responsibility to deallocate memory pointed to by device_ptr.

Functions

template<typename T, typename MR> __host__ __device__ bool operator== (const allocator< T, MR > &lhs, const allocator< T, MR > &rhs) noexcept

Compares the allocators for equality by comparing the underlying memory resources.

template<typename T, typename MR> __host__ __device__ bool operator!= (const allocator< T, MR > &lhs, const allocator< T, MR > &rhs) noexcept

Compares the allocators for inequality by comparing the underlying memory resources.

template<typename Upstream>
class device_ptr_memory_resource : public thrust::mr::memory_resource<device_ptr<void>>#
#include <device_allocator.h>

Memory resource adaptor that turns any memory resource that returns a fancy with the same tag as device_ptr, and adapts it to a resource that returns a device_ptr.

template<typename T>
class device_allocator : public thrust::mr::stateless_resource_allocator<T, device_ptr_memory_resource<device_memory_resource>>#
#include <device_allocator.h>

An allocator which creates new elements in memory accessible by devices.

template<typename T>
class device_malloc_allocator#
#include <device_malloc_allocator.h>

device_malloc_allocator is a device memory allocator that employs the device_malloc function for allocation.

device_malloc_allocator is deprecated in favor of thrust::mr memory resource-based allocators.

See also

device_malloc

See also

device_ptr

See also

device_allocator

template<typename T>
class device_new_allocator#
#include <device_new_allocator.h>

device_new_allocator is a device memory allocator that employs the device_new function for allocation.

See also

device_new

See also

device_ptr

template<typename T, class MR>
class allocator : private mr::validator<MR>#
#include <allocator.h>

An mr::allocator is a template that fulfills the C++ requirements for Allocators, allowing to use the NPA-based memory resources where an Allocator is required. Unlike memory resources, but like other allocators, mr::allocator is typed and bound to allocate object of a specific type, however it can be freely rebound to other types.

Template Parameters
  • T – the type that will be allocated by this allocator.

  • MR – the upstream memory resource to use for memory allocation. Must derive from thrust::mr::memory_resource and must be final (in C++11 and beyond).

template<typename T, typename Pointer>
class polymorphic_allocator : public mr::allocator<T, polymorphic_adaptor_resource<Pointer>>#
#include <allocator.h>

An allocator whose behaviour depends on a memory resource that we can dynamically configure at runtime.

Template Parameters
  • T – - the type that will be allocated by this allocator

  • Pointer – - the pointer type that will be used to create the memory resource

template<typename T, typename Upstream>
class stateless_resource_allocator : public thrust::mr::allocator<T, Upstream>#
#include <allocator.h>

A helper allocator class that uses global instances of a given upstream memory resource. Requires the memory resource to be default constructible.

Template Parameters
  • T – the type that will be allocated by this allocator.

  • Upstream – the upstream memory resource to use for memory allocation. Must derive from thrust::mr::memory_resource and must be final (in C++11 and beyond).

Typedefs

template<typename T>
using universal_ptr = thrust::system::__THRUST_DEVICE_SYSTEM_NAMESPACE::universal_pointer<T>#

universal_ptr stores a pointer to an object allocated in memory accessible to both hosts and devices.

Algorithms dispatched with this type of pointer will be dispatched to either host or device, depending on which backend you are using. Explicit policies (thrust::device, etc) can be used to specify where an algorithm should be run.

universal_ptr has pointer semantics: it may be dereferenced safely from both hosts and devices and may be manipulated with pointer arithmetic.

universal_ptr can be created with universal_allocator or by explicitly calling its constructor with a raw pointer.

The raw pointer encapsulated by a universal_ptr may be obtained by either its get method or the raw_pointer_cast free function.

See also

host_ptr For the documentation of the complete interface which is shared by universal_ptr.

See also

raw_pointer_cast

Note

universal_ptr is not a smart pointer; it is the programmer’s responsibility to deallocate memory pointed to by universal_ptr.

Functions

template<typename Pointer> __host__ __device__ bool operator== (const memory_resource< Pointer > &lhs, const memory_resource< Pointer > &rhs) noexcept

Compares the memory resources for equality, first by identity, then by is_equal.

template<typename Pointer> __host__ __device__ bool operator!= (const memory_resource< Pointer > &lhs, const memory_resource< Pointer > &rhs) noexcept

Compares the memory resources for inequality, first by identity, then by is_equal.

template<typename MR> __host__ MR * get_global_resource ()

Returns a global instance of MR, created as a function local static variable.

Template Parameters

MR – type of a memory resource to get an instance from. Must be DefaultConstructible.

Returns

a pointer to a global instance of MR.

template<typename Upstream, typename Bookkeeper>
class disjoint_unsynchronized_pool_resource : public mr::memory_resource<Upstream::pointer>, private mr::validator2<Upstream, Bookkeeper>#
#include <disjoint_pool.h>

A memory resource adaptor allowing for pooling and caching allocations from Upstream, using Bookkeeper for management of that cached and pooled memory, allowing to cache portions of memory inaccessible from the host.

On a typical memory resource, calls to allocate and deallocate actually allocate and deallocate memory. Pooling memory resources only allocate and deallocate memory from an external resource (the upstream memory resource) when there’s no suitable memory currently cached; otherwise, they use memory they have acquired beforehand, to make memory allocation faster and more efficient.

The disjoint version of the pool resources uses a separate upstream memory resource, Bookkeeper, to allocate memory necessary to manage the cached memory. There may be many reasons to do that; the canonical one is that Upstream allocates memory that is inaccessible to the code of the pool resource, which means that it cannot embed the necessary information in memory obtained from Upstream; for instance, Upstream can be a CUDA non-managed memory resource, or a CUDA managed memory resource whose memory we would prefer to not migrate back and forth between host and device when executing bookkeeping code.

This is not the only case where it makes sense to use a disjoint pool resource, though. In a multi-core environment it may be beneficial to avoid stealing cache lines from other cores by writing over bookkeeping information embedded in an allocated block of memory. In such a case, one can imagine wanting to use a disjoint pool where both the upstream and the bookkeeper are of the same type, to allocate memory consistently, but separately for those two purposes.

Template Parameters
  • Upstream – the type of memory resources that will be used for allocating memory blocks to be handed off to the user

  • Bookkeeper – the type of memory resources that will be used for allocating bookkeeping memory

template<typename Pointer = void*>
class memory_resource#
#include <memory_resource.h>

memory_resource is the base class for all other memory resources.

Template Parameters

Pointer – the pointer type that is allocated and deallocated by the memory resource derived from this base class. If this is void *, this class derives from std::pmr::memory_resource.

Subclassed by mr::new_delete_resource, mr::polymorphic_adaptor_resource< Pointer >

class new_delete_resource : public mr::memory_resource<>#
#include <new.h>

A memory resource that uses global operators new and delete to allocate and deallocate memory. Uses alignment-enabled overloads when available, otherwise uses regular overloads and implements alignment requirements by itself.

template<typename Upstream>
class unsynchronized_pool_resource : public mr::memory_resource<Upstream::pointer>, private mr::validator<Upstream>#
#include <pool.h>

A memory resource adaptor allowing for pooling and caching allocations from Upstream, using memory allocated from it for both blocks then allocated to the user and for internal bookkeeping of the cached memory.

On a typical memory resource, calls to allocate and deallocate actually allocate and deallocate memory. Pooling memory resources only allocate and deallocate memory from an external resource (the upstream memory resource) when there’s no suitable memory currently cached; otherwise, they use memory they have acquired beforehand, to make memory allocation faster and more efficient.

The non-disjoint version of the pool resource uses a single upstream memory resource. Every allocation is larger than strictly necessary to fulfill the end-user’s request, because it needs to account for the memory overhead of tracking the memory blocks and chunks inside those same memory regions. Nevertheless, this version should be more memory-efficient than the disjoint_unsynchronized_pool_resource, because it doesn’t need to allocate additional blocks of memory from a separate resource, which in turn would necessitate the bookkeeping overhead in the upstream resource.

This version requires that memory allocated from Upstream is accessible from device. It supports smart references, meaning that the non-managed CUDA resource, returning a device-tagged pointer, will work, but will be much less efficient than the disjoint version, which wouldn’t need to touch device memory at all, and therefore wouldn’t need to transfer it back and forth between the host and the device whenever an allocation or a deallocation happens.

Template Parameters

Upstream – the type of memory resources that will be used for allocating memory blocks

struct pool_options#
#include <pool_options.h>

A type used for configuring pooling resource adaptors, to fine-tune their behavior and parameters.

Iterators#

group iterators

Functions

template<typename InputIterator, typename Distance> __host__ __device__ void advance (InputIterator &i, Distance n)

advance(i, n) increments the iterator i by the distance n. If n > 0 it is equivalent to executing ++i n times, and if n < 0 it is equivalent to executing i n times. If n == 0, the call has no effect.

The following code snippet demonstrates how to use advance to increment an iterator a given number of times.

#include <thrust/advance.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> vec(13);
thrust::device_vector<int>::iterator iter = vec.begin();

thrust::advance(iter, 7);

// iter - vec.begin() == 7

Parameters
  • i – The iterator to be advanced.

  • n – The distance by which to advance the iterator.

Template Parameters
  • InputIterator – is a model of Input Iterator.

  • Distance – is an integral type that is convertible to InputIterator's distance type.

Pre

n shall be negative only for bidirectional and random access iterators.

template<typename InputIterator> inline __host__ __device__ thrust::iterator_traits< InputIterator >::difference_type distance (InputIterator first, InputIterator last)

distance finds the distance between first and last, i.e. the number of times that first must be incremented until it is equal to last.

The following code snippet demonstrates how to use distance to compute the distance to one iterator from another.

#include <thrust/distance.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> vec(13);
thrust::device_vector<int>::iterator iter1 = vec.begin();
thrust::device_vector<int>::iterator iter2 = iter1 + 7;

int d = thrust::distance(iter1, iter2);

// d is 7

Parameters
  • first – The beginning of an input range of interest.

  • last – The end of an input range of interest.

Template Parameters

InputIterator – is a model of Input Iterator.

Returns

The distance between the beginning and end of the input range.

Pre

If InputIterator meets the requirements of random access iterator, last shall be reachable from first or first shall be reachable from last; otherwise, last shall be reachable from first.

Functions

template<typename ValueT, typename IndexT> inline __host__ __device__ constant_iterator< ValueT, IndexT > make_constant_iterator (ValueT x, IndexT i=int())

This version of make_constant_iterator creates a constant_iterator from values given for both value and index. The type of constant_iterator may be inferred by the compiler from the types of its parameters.

Parameters
  • x – The value of the returned constant_iterator's constant value.

  • i – The index of the returned constant_iterator within a sequence. The type of this parameter defaults to int. In the default case, the value of this parameter is 0.

Returns

A new constant_iterator with constant value & index as given by x & i.

template<typename V> inline __host__ __device__ constant_iterator< V > make_constant_iterator (V x)

This version of make_constant_iterator creates a constant_iterator using only a parameter for the desired constant value. The value of the returned constant_iterator's index is set to 0.

Parameters

x – The value of the returned constant_iterator's constant value.

Returns

A new constant_iterator with constant value equal to x and index equal to 0.

template<typename Incrementable> inline __host__ __device__ counting_iterator< Incrementable > make_counting_iterator (Incrementable x)

make_counting_iterator creates a counting_iterator using an initial value for its Incrementable counter.

Parameters

x – The initial value of the new counting_iterator's counter.

Returns

A new counting_iterator whose counter has been initialized to x.

inline __host__ __device__ discard_iterator make_discard_iterator (discard_iterator<>::difference_type i=discard_iterator<>::difference_type(0))

make_discard_iterator creates a discard_iterator from an optional index parameter.

Parameters

i – The index of the returned discard_iterator within a range. In the default case, the value of this parameter is 0.

Returns

A new discard_iterator with index as given by i.

template<typename ElementIterator, typename IndexIterator> __host__ __device__ permutation_iterator< ElementIterator, IndexIterator > make_permutation_iterator (ElementIterator e, IndexIterator i)

make_permutation_iterator creates a permutation_iterator from an ElementIterator pointing to a range of elements to “permute” and an IndexIterator pointing to a range of indices defining an indexing scheme on the values.

Parameters
  • e – An ElementIterator pointing to a range of values.

  • i – An IndexIterator pointing to an indexing scheme to use on e.

Returns

A new permutation_iterator which permutes the range e by i.

template<typename BidirectionalIterator> __host__ __device__ reverse_iterator< BidirectionalIterator > make_reverse_iterator (BidirectionalIterator x)

make_reverse_iterator creates a reverse_iterator from a BidirectionalIterator pointing to a range of elements to reverse.

Parameters

x – A BidirectionalIterator pointing to a range to reverse.

Returns

A new reverse_iterator which reverses the range x.

template<typename InputFunction, typename OutputFunction, typename Iterator> transform_input_output_iterator< InputFunction, OutputFunction, Iterator > __host__ __device__ make_transform_input_output_iterator (Iterator io, InputFunction input_function, OutputFunction output_function)

make_transform_input_output_iterator creates a transform_input_output_iterator from an Iterator a InputFunction and a OutputFunction

Parameters
  • io – An Iterator pointing to where the input to InputFunction will be read from and the result of OutputFunction will be written to

  • input_function – An InputFunction to be executed on values read from the iterator

  • output_function – An OutputFunction to be executed on values written to the iterator

template<class AdaptableUnaryFunction, class Iterator> inline __host__ __device__ transform_iterator< AdaptableUnaryFunction, Iterator > make_transform_iterator (Iterator it, AdaptableUnaryFunction fun)

make_transform_iterator creates a transform_iterator from an Iterator and AdaptableUnaryFunction.

Parameters
  • it – The Iterator pointing to the input range of the newly created transform_iterator.

  • fun – The AdaptableUnaryFunction used to transform the range pointed to by it in the newly created transform_iterator.

Returns

A new transform_iterator which transforms the range at it by fun.

template<typename UnaryFunction, typename OutputIterator> transform_output_iterator< UnaryFunction, OutputIterator > __host__ __device__ make_transform_output_iterator (OutputIterator out, UnaryFunction fun)

make_transform_output_iterator creates a transform_output_iterator from an OutputIterator and UnaryFunction.

Parameters
template<typename... Iterators> inline __host__ __device__ zip_iterator< thrust::tuple< Iterators... > > make_zip_iterator (thrust::tuple< Iterators... > t)

make_zip_iterator creates a zip_iterator from a tuple of iterators.

See also

zip_iterator

Parameters

t – The tuple of iterators to copy.

Returns

A newly created zip_iterator which zips the iterators encapsulated in t.

template<typename... Iterators> inline __host__ __device__ zip_iterator< thrust::tuple< Iterators... > > make_zip_iterator (Iterators... its)

make_zip_iterator creates a zip_iterator from iterators.

See also

zip_iterator

Parameters

its – The iterators to copy.

Returns

A newly created zip_iterator which zips the iterators.

template<typename Value, typename Incrementable = use_default, typename System = use_default>
class constant_iterator : public detail::constant_iterator_base::type#
#include <constant_iterator.h>

constant_iterator is an iterator which represents a pointer into a range of constant values. This iterator is useful for creating a range filled with the same value without explicitly storing it in memory. Using constant_iterator saves both memory capacity and bandwidth.

The following code snippet demonstrates how to create a constant_iterator whose value_type is int and whose value is 10.

#include <thrust/iterator/constant_iterator.h>

thrust::constant_iterator<int> iter(10);

*iter;    // returns 10
iter[0];  // returns 10
iter[1];  // returns 10
iter[13]; // returns 10

// and so on...

This next example demonstrates how to use a constant_iterator with the thrust::transform function to increment all elements of a sequence by the same value. We will create a temporary constant_iterator with the function make_constant_iterator function in order to avoid explicitly specifying its type:

#include <thrust/iterator/constant_iterator.h>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>

int main()
{
  thrust::device_vector<int> data(4);
  data[0] = 3;
  data[1] = 7;
  data[2] = 2;
  data[3] = 5;

  // add 10 to all values in data
  thrust::transform(data.begin(), data.end(),
                    thrust::make_constant_iterator(10),
                    data.begin(),
                    thrust::plus<int>());

  // data is now [13, 17, 12, 15]

  return 0;
}

template<typename Incrementable, typename System = use_default, typename Traversal = use_default, typename Difference = use_default>
class counting_iterator : public detail::counting_iterator_base::type#
#include <counting_iterator.h>

counting_iterator is an iterator which represents a pointer into a range of sequentially changing values. This iterator is useful for creating a range filled with a sequence without explicitly storing it in memory. Using counting_iterator saves memory capacity and bandwidth.

The following code snippet demonstrates how to create a counting_iterator whose value_type is int and which sequentially increments by 1.

#include <thrust/iterator/counting_iterator.h>
...
// create iterators
thrust::counting_iterator<int> first(10);
thrust::counting_iterator<int> last = first + 3;

first[0]   // returns 10
first[1]   // returns 11
first[100] // returns 110

// sum of [first, last)
thrust::reduce(first, last);   // returns 33 (i.e. 10 + 11 + 12)

// initialize vector to [0,1,2,..]
thrust::counting_iterator<int> iter(0);
thrust::device_vector<int> vec(500);
thrust::copy(iter, iter + vec.size(), vec.begin());

This next example demonstrates how to use a counting_iterator with the thrust::copy_if function to compute the indices of the non-zero elements of a device_vector. In this example, we use the make_counting_iterator function to avoid specifying the type of the counting_iterator.

#include <thrust/iterator/counting_iterator.h>
#include <thrust/copy.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>

int main()
{
 // this example computes indices for all the nonzero values in a sequence

 // sequence of zero and nonzero values
 thrust::device_vector<int> stencil(8);
 stencil[0] = 0;
 stencil[1] = 1;
 stencil[2] = 1;
 stencil[3] = 0;
 stencil[4] = 0;
 stencil[5] = 1;
 stencil[6] = 0;
 stencil[7] = 1;

 // storage for the nonzero indices
 thrust::device_vector<int> indices(8);

 // compute indices of nonzero elements
 typedef thrust::device_vector<int>::iterator IndexIterator;

 // use make_counting_iterator to define the sequence [0, 8)
 IndexIterator indices_end = thrust::copy_if(thrust::make_counting_iterator(0),
                                             thrust::make_counting_iterator(8),
                                             stencil.begin(),
                                             indices.begin(),
                                             thrust::identity<int>());
 // indices now contains [1,2,5,7]

 return 0;
}

template<typename System = use_default>
class discard_iterator : public detail::discard_iterator_base::type#
#include <discard_iterator.h>

discard_iterator is an iterator which represents a special kind of pointer that ignores values written to it upon dereference. This iterator is useful for ignoring the output of certain algorithms without wasting memory capacity or bandwidth. discard_iterator may also be used to count the size of an algorithm’s output which may not be known a priori.

The following code snippet demonstrates how to use discard_iterator to ignore ignore one of the output ranges of reduce_by_key

#include <thrust/iterator/discard_iterator.h>
#include <thrust/reduce.h>
#include <thrust/device_vector.h>

int main()
{
  thrust::device_vector<int> keys(7), values(7);

  keys[0] = 1;
  keys[1] = 3;
  keys[2] = 3;
  keys[3] = 3;
  keys[4] = 2;
  keys[5] = 2;
  keys[6] = 1;

  values[0] = 9;
  values[1] = 8;
  values[2] = 7;
  values[3] = 6;
  values[4] = 5;
  values[5] = 4;
  values[6] = 3;

  thrust::device_vector<int> result(4);

  // we are only interested in the reduced values
  // use discard_iterator to ignore the output keys
  thrust::reduce_by_key(keys.begin(), keys.end(),
                        values.begin(),
                        thrust::make_discard_iterator(),
                        result.begin());

  // result is now [9, 21, 9, 3]

  return 0;
}

template<typename Derived, typename Base, typename Value = use_default, typename System = use_default, typename Traversal = use_default, typename Reference = use_default, typename Difference = use_default>
class iterator_adaptor : public detail::iterator_adaptor_base::type#
#include <iterator_adaptor.h>

iterator_adaptor is an iterator which adapts an existing type of iterator to create a new type of iterator. Most of Thrust’s fancy iterators are defined via inheritance from iterator_adaptor. While composition of these existing Thrust iterators is often sufficient for expressing the desired functionality, it is occasionally more straightforward to derive from iterator_adaptor directly.

To see how to use iterator_adaptor to create a novel iterator type, let’s examine how to use it to define repeat_iterator, a fancy iterator which repeats elements from another range a given number of time:

#include <thrust/iterator/iterator_adaptor.h>

// derive repeat_iterator from iterator_adaptor
template<typename Iterator>
  class repeat_iterator
    : public thrust::iterator_adaptor<
        repeat_iterator<Iterator>, // the first template parameter is the name of the iterator we're creating
        Iterator                   // the second template parameter is the name of the iterator we're adapting
                                   // we can use the default for the additional template parameters
      >
{
  public:
    // shorthand for the name of the iterator_adaptor we're deriving from
    typedef thrust::iterator_adaptor<
      repeat_iterator<Iterator>,
      Iterator
    > super_t;

    __host__ __device__
    repeat_iterator(const Iterator &x, int n) : super_t(x), begin(x), n(n) {}

    // befriend thrust::iterator_core_access to allow it access to the private interface below
    friend class thrust::iterator_core_access;

  private:
    // repeat each element of the adapted range n times
    unsigned int n;

    // used to keep track of where we began
    const Iterator begin;

    // it is private because only thrust::iterator_core_access needs access to it
    __host__ __device__
    typename super_t::reference dereference() const
    {
      return *(begin + (this->base() - begin) / n);
    }
};

Except for the first two, iterator_adaptor's template parameters are optional. When omitted, or when the user specifies thrust::use_default in its place, iterator_adaptor will use a default type inferred from Base.

iterator_adaptor's functionality is derived from and generally equivalent to boost::iterator_adaptor. The exception is Thrust’s addition of the template parameter System, which is necessary to allow Thrust to dispatch an algorithm to one of several parallel backend systems.

iterator_adaptor is a powerful tool for creating custom iterators directly. However, the large set of iterator semantics which must be satisfied for algorithm compatibility can make iterator_adaptor difficult to use correctly. Unless you require the full expressivity of iterator_adaptor, consider building a custom iterator through composition of existing higher-level fancy iterators instead.

Interested users may refer to boost::iterator_adaptor’s documentation for further usage examples.

template<typename Derived, typename Value, typename System, typename Traversal, typename Reference, typename Difference = std::ptrdiff_t>
class iterator_facade#
#include <iterator_facade.h>

iterator_facade is a template which allows the programmer to define a novel iterator with a standards-conforming interface which Thrust can use to reason about algorithm acceleration opportunities.

Because most of a standard iterator’s interface is defined in terms of a small set of core primitives, iterator_facade defines the non-primitive portion mechanically. In principle a novel iterator could explicitly provide the entire interface in an ad hoc fashion but doing so might be tedious and prone to subtle errors.

Often iterator_facade is too primitive a tool to use for defining novel iterators. In these cases, iterator_adaptor or a specific fancy iterator should be used instead.

iterator_facade's functionality is derived from and generally equivalent to boost::iterator_facade. The exception is Thrust’s addition of the template parameter System, which is necessary to allow Thrust to dispatch an algorithm to one of several parallel backend systems. An additional exception is Thrust’s omission of the operator-> member function.

Interested users may refer to boost::iterator_facade’s documentation for usage examples.

Note

iterator_facade's arithmetic operator free functions exist with the usual meanings but are omitted here for brevity.

class iterator_core_access#
#include <iterator_facade.h>

iterator_core_access is the class which user iterator types derived from thrust::iterator_adaptor or thrust::iterator_facade must befriend to allow it to access their private interface.

template<typename ElementIterator, typename IndexIterator>
class permutation_iterator : public thrust::detail::permutation_iterator_base::type#
#include <permutation_iterator.h>

permutation_iterator is an iterator which represents a pointer into a reordered view of a given range. permutation_iterator is an imprecise name; the reordered view need not be a strict permutation. This iterator is useful for fusing a scatter or gather operation with other algorithms.

This iterator takes two arguments:

  • an iterator to the range V on which the “permutation” will be applied

  • the reindexing scheme that defines how the elements of V will be permuted.

Note that permutation_iterator is not limited to strict permutations of the given range V. The distance between begin and end of the reindexing iterators is allowed to be smaller compared to the size of the range V, in which case the permutation_iterator only provides a “permutation” of a subrange of V. The indices neither need to be unique. In this same context, it must be noted that the past-the-end permutation_iterator is completely defined by means of the past-the-end iterator to the indices.

The following code snippet demonstrates how to create a permutation_iterator which represents a reordering of the contents of a device_vector.

#include <thrust/iterator/permutation_iterator.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<float> values(8);
values[0] = 10.0f;
values[1] = 20.0f;
values[2] = 30.0f;
values[3] = 40.0f;
values[4] = 50.0f;
values[5] = 60.0f;
values[6] = 70.0f;
values[7] = 80.0f;

thrust::device_vector<int> indices(4);
indices[0] = 2;
indices[1] = 6;
indices[2] = 1;
indices[3] = 3;

typedef thrust::device_vector<float>::iterator ElementIterator;
typedef thrust::device_vector<int>::iterator   IndexIterator;

thrust::permutation_iterator<ElementIterator,IndexIterator> iter(values.begin(), indices.begin());

*iter;   // returns 30.0f;
iter[0]; // returns 30.0f;
iter[1]; // returns 70.0f;
iter[2]; // returns 20.0f;
iter[3]; // returns 40.0f;

// iter[4] is an out-of-bounds error

*iter   = -1.0f; // sets values[2] to -1.0f;
iter[0] = -1.0f; // sets values[2] to -1.0f;
iter[1] = -1.0f; // sets values[6] to -1.0f;
iter[2] = -1.0f; // sets values[1] to -1.0f;
iter[3] = -1.0f; // sets values[3] to -1.0f;

// values is now {10, -1, -1, -1, 50, 60, -1, 80}

template<typename BidirectionalIterator>
class reverse_iterator : public detail::reverse_iterator_base::type#
#include <reverse_iterator.h>

reverse_iterator is an iterator which represents a pointer into a reversed view of a given range. In this way, reverse_iterator allows backwards iteration through a bidirectional input range.

It is important to note that although reverse_iterator is constructed from a given iterator, it points to the element preceding it. In this way, the past-the-end reverse_iterator of a given range points to the element preceding the first element of the input range. By the same token, the first reverse_iterator of a given range is constructed from a past-the-end iterator of the original range yet points to the last element of the input.

The following code snippet demonstrates how to create a reverse_iterator which represents a reversed view of the contents of a device_vector.

#include <thrust/iterator/reverse_iterator.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<float> v(4);
v[0] = 0.0f;
v[1] = 1.0f;
v[2] = 2.0f;
v[3] = 3.0f;

typedef thrust::device_vector<float>::iterator Iterator;

// note that we point the iterator to the *end* of the device_vector
thrust::reverse_iterator<Iterator> iter(values.end());

*iter;   // returns 3.0f;
iter[0]; // returns 3.0f;
iter[1]; // returns 2.0f;
iter[2]; // returns 1.0f;
iter[3]; // returns 0.0f;

// iter[4] is an out-of-bounds error

Since reversing a range is a common operation, containers like device_vector have nested typedefs for declaration shorthand and methods for constructing reverse_iterators. The following code snippet is equivalent to the previous:

#include <thrust/device_vector.h>
...
thrust::device_vector<float> v(4);
v[0] = 0.0f;
v[1] = 1.0f;
v[2] = 2.0f;
v[3] = 3.0f;

// we use the nested type reverse_iterator to refer to a reversed view of
// a device_vector and the method rbegin() to create a reverse_iterator pointing
// to the beginning of the reversed device_vector
thrust::device_iterator<float>::reverse_iterator iter = values.rbegin();

*iter;   // returns 3.0f;
iter[0]; // returns 3.0f;
iter[1]; // returns 2.0f;
iter[2]; // returns 1.0f;
iter[3]; // returns 0.0f;

// iter[4] is an out-of-bounds error

// similarly, rend() points to the end of the reversed sequence:
assert(values.rend() == (iter + 4));

Finally, the following code snippet demonstrates how to use reverse_iterator to perform a reversed prefix sum operation on the contents of a device_vector:

#include <thrust/device_vector.h>
#include <thrust/scan.h>
...
thrust::device_vector<int> v(5);
v[0] = 0;
v[1] = 1;
v[2] = 2;
v[3] = 3;
v[4] = 4;

thrust::device_vector<int> result(5);

// exclusive scan v into result in reverse
thrust::exclusive_scan(v.rbegin(), v.rend(), result.begin());

// result is now {0, 4, 7, 9, 10}

template<typename InputFunction, typename OutputFunction, typename Iterator>
class transform_input_output_iterator : public detail::transform_input_output_iterator_base::type#
#include <transform_input_output_iterator.h>

transform_input_output_iterator is a special kind of iterator which applies transform functions when reading from or writing to dereferenced values. This iterator is useful for algorithms that operate on a type that needs to be serialized/deserialized from values in another iterator, avoiding the need to materialize intermediate results in memory. This also enables the transform functions to be fused with the operations that read and write to the transform_input_output_iterator.

The following code snippet demonstrates how to create a transform_input_output_iterator which performs different transformations when reading from and writing to the iterator.

#include <thrust/iterator/transform_input_output_iterator.h>
#include <thrust/device_vector.h>

 int main()
 {
   const size_t size = 4;
   thrust::device_vector<float> v(size);

   // Write 1.0f, 2.0f, 3.0f, 4.0f to vector
   thrust::sequence(v.begin(), v.end(), 1);

   // Iterator that returns negated values and writes squared values
   auto iter = thrust::make_transform_input_output_iterator(v.begin(),
       thrust::negate<float>{}, thrust::square<float>{});

   // Iterator negates values when reading
   std::cout << iter[0] << " ";  // -1.0f;
   std::cout << iter[1] << " ";  // -2.0f;
   std::cout << iter[2] << " ";  // -3.0f;
   std::cout << iter[3] << "\n"; // -4.0f;

   // Write 1.0f, 2.0f, 3.0f, 4.0f to iterator
   thrust::sequence(iter, iter + size, 1);

   // Values were squared before writing to vector
   std::cout << v[0] << " ";  // 1.0f;
   std::cout << v[1] << " ";  // 4.0f;
   std::cout << v[2] << " ";  // 9.0f;
   std::cout << v[3] << "\n"; // 16.0f;

 }

template<class AdaptableUnaryFunction, class Iterator, class Reference = use_default, class Value = use_default>
class transform_iterator : public detail::transform_iterator_base::type#
#include <transform_iterator.h>

transform_iterator is an iterator which represents a pointer into a range of values after transformation by a function. This iterator is useful for creating a range filled with the result of applying an operation to another range without either explicitly storing it in memory, or explicitly executing the transformation. Using transform_iterator facilitates kernel fusion by deferring the execution of a transformation until the value is needed while saving both memory capacity and bandwidth.

The following code snippet demonstrates how to create a transform_iterator which represents the result of sqrtf applied to the contents of a device_vector.

#include <thrust/iterator/transform_iterator.h>
#include <thrust/device_vector.h>

// note: functor inherits from unary_function
struct square_root : public thrust::unary_function<float,float>
{
  __host__ __device__
  float operator()(float x) const
  {
    return sqrtf(x);
  }
};

int main()
{
  thrust::device_vector<float> v(4);
  v[0] = 1.0f;
  v[1] = 4.0f;
  v[2] = 9.0f;
  v[3] = 16.0f;

  typedef thrust::device_vector<float>::iterator FloatIterator;

  thrust::transform_iterator<square_root, FloatIterator> iter(v.begin(), square_root());

  *iter;   // returns 1.0f
  iter[0]; // returns 1.0f;
  iter[1]; // returns 2.0f;
  iter[2]; // returns 3.0f;
  iter[3]; // returns 4.0f;

  // iter[4] is an out-of-bounds error
}

This next example demonstrates how to use a transform_iterator with the thrust::reduce function to compute the sum of squares of a sequence. We will create temporary transform_iterators with the make_transform_iterator function in order to avoid explicitly specifying their type:

#include <thrust/iterator/transform_iterator.h>
#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <iostream>

// note: functor inherits from unary_function
struct square : public thrust::unary_function<float,float>
{
  __host__ __device__
  float operator()(float x) const
  {
    return x * x;
  }
};

int main()
{
  // initialize a device array
  thrust::device_vector<float> v(4);
  v[0] = 1.0f;
  v[1] = 2.0f;
  v[2] = 3.0f;
  v[3] = 4.0f;

  float sum_of_squares =
   thrust::reduce(thrust::make_transform_iterator(v.begin(), square()),
                  thrust::make_transform_iterator(v.end(),   square()));

  std::cout << "sum of squares: " << sum_of_squares << std::endl;
  return 0;
}

Note that in the previous two examples the transform functor (namely square_root and square) inherits from thrust::unary_function. Inheriting from thrust::unary_function ensures that a functor is a valid AdaptableUnaryFunction and provides all the necessary typedef declarations. The transform_iterator can also be applied to a UnaryFunction that does not inherit from thrust::unary_function using an optional template argument. The following example illustrates how to use the third template argument to specify the result_type of the function.

#include <thrust/iterator/transform_iterator.h>
#include <thrust/device_vector.h>

// note: functor *does not* inherit from unary_function
struct square_root
{
  __host__ __device__
  float operator()(float x) const
  {
    return sqrtf(x);
  }
};

int main()
{
  thrust::device_vector<float> v(4);
  v[0] = 1.0f;
  v[1] = 4.0f;
  v[2] = 9.0f;
  v[3] = 16.0f;

  typedef thrust::device_vector<float>::iterator FloatIterator;

  // note: float result_type is specified explicitly
  thrust::transform_iterator<square_root, FloatIterator, float> iter(v.begin(), square_root());

  *iter;   // returns 1.0f
  iter[0]; // returns 1.0f;
  iter[1]; // returns 2.0f;
  iter[2]; // returns 3.0f;
  iter[3]; // returns 4.0f;

  // iter[4] is an out-of-bounds error
}

template<typename UnaryFunction, typename OutputIterator>
class transform_output_iterator : public detail::transform_output_iterator_base::type#
#include <transform_output_iterator.h>

transform_output_iterator is a special kind of output iterator which transforms a value written upon dereference. This iterator is useful for transforming an output from algorithms without explicitly storing the intermediate result in the memory and applying subsequent transformation, thereby avoiding wasting memory capacity and bandwidth. Using transform_iterator facilitates kernel fusion by deferring execution of transformation until the value is written while saving both memory capacity and bandwidth.

The following code snippet demonstrated how to create a transform_output_iterator which applies sqrtf to the assigning value.

#include <thrust/iterator/transform_output_iterator.h>
#include <thrust/device_vector.h>

// note: functor inherits form unary function
 // note: functor inherits from unary_function
 struct square_root : public thrust::unary_function<float,float>
 {
   __host__ __device__
   float operator()(float x) const
   {
     return sqrtf(x);
   }
 };

 int main()
 {
   thrust::device_vector<float> v(4);

   typedef thrust::device_vector<float>::iterator FloatIterator;
   thrust::transform_output_iterator<square_root, FloatIterator> iter(v.begin(), square_root());

   iter[0] =  1.0f;    // stores sqrtf( 1.0f)
   iter[1] =  4.0f;    // stores sqrtf( 4.0f)
   iter[2] =  9.0f;    // stores sqrtf( 9.0f)
   iter[3] = 16.0f;    // stores sqrtf(16.0f)
   // iter[4] is an out-of-bounds error

   v[0]; // returns 1.0f;
   v[1]; // returns 2.0f;
   v[2]; // returns 3.0f;
   v[3]; // returns 4.0f;

 }

template<typename IteratorTuple>
class zip_iterator : public detail::zip_iterator_base::type#
#include <zip_iterator.h>

zip_iterator is an iterator which represents a pointer into a range of tuples whose elements are themselves taken from a tuple of input iterators. This iterator is useful for creating a virtual array of structures while achieving the same performance and bandwidth as the structure of arrays idiom. zip_iterator also facilitates kernel fusion by providing a convenient means of amortizing the execution of the same operation over multiple ranges.

The following code snippet demonstrates how to create a zip_iterator which represents the result of “zipping” multiple ranges together.

#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> int_v(3);
int_v[0] = 0; int_v[1] = 1; int_v[2] = 2;

thrust::device_vector<float> float_v(3);
float_v[0] = 0.0f; float_v[1] = 1.0f; float_v[2] = 2.0f;

thrust::device_vector<char> char_v(3);
char_v[0] = 'a'; char_v[1] = 'b'; char_v[2] = 'c';

// typedef these iterators for shorthand
typedef thrust::device_vector<int>::iterator   IntIterator;
typedef thrust::device_vector<float>::iterator FloatIterator;
typedef thrust::device_vector<char>::iterator  CharIterator;

// typedef a tuple of these iterators
typedef thrust::tuple<IntIterator, FloatIterator, CharIterator> IteratorTuple;

// typedef the zip_iterator of this tuple
typedef thrust::zip_iterator<IteratorTuple> ZipIterator;

// finally, create the zip_iterator
ZipIterator iter(thrust::make_tuple(int_v.begin(), float_v.begin(), char_v.begin()));

*iter;   // returns (0, 0.0f, 'a')
iter[0]; // returns (0, 0.0f, 'a')
iter[1]; // returns (1, 1.0f, 'b')
iter[2]; // returns (2, 2.0f, 'c')

thrust::get<0>(iter[2]); // returns 2
thrust::get<1>(iter[0]); // returns 0.0f
thrust::get<2>(iter[1]); // returns 'b'

// iter[3] is an out-of-bounds error

Defining the type of a zip_iterator can be complex. The next code example demonstrates how to use the make_zip_iterator function with the make_tuple function to avoid explicitly specifying the type of the zip_iterator. This example shows how to use zip_iterator to copy multiple ranges with a single call to thrust::copy.

#include <thrust/zip_iterator.h>
#include <thrust/tuple.h>
#include <thrust/device_vector.h>

int main()
{
  thrust::device_vector<int> int_in(3), int_out(3);
  int_in[0] = 0;
  int_in[1] = 1;
  int_in[2] = 2;

  thrust::device_vector<float> float_in(3), float_out(3);
  float_in[0] =  0.0f;
  float_in[1] = 10.0f;
  float_in[2] = 20.0f;

  thrust::copy(thrust::make_zip_iterator(thrust::make_tuple(int_in.begin(), float_in.begin())),
               thrust::make_zip_iterator(thrust::make_tuple(int_in.end(),   float_in.end())),
               thrust::make_zip_iterator(thrust::make_tuple(int_out.begin(),float_out.begin())));

  // int_out is now [0, 1, 2]
  // float_out is now [0.0f, 10.0f, 20.0f]

  return 0;
}

See also

make_tuple

See also

tuple

See also

get

Typedefs

typedef std::input_iterator_tag input_host_iterator_tag#

input_host_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Input Host Iterator concept within the C++ type system.

typedef std::output_iterator_tag output_host_iterator_tag#

output_host_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Output Host Iterator concept within the C++ type system.

typedef std::forward_iterator_tag forward_host_iterator_tag#

forward_host_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Forward Host Iterator concept within the C++ type system.

typedef std::bidirectional_iterator_tag bidirectional_host_iterator_tag#

bidirectional_host_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Forward Host Iterator concept within the C++ type system.

typedef std::random_access_iterator_tag random_access_host_iterator_tag#

random_access_host_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Forward Host Iterator concept within the C++ type system.

struct input_device_iterator_tag : public thrust::detail::iterator_category_with_system_and_traversal<std::input_iterator_tag, thrust::device_system_tag, thrust::single_pass_traversal_tag>#
#include <iterator_categories.h>

input_device_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Input Device Iterator concept within the C++ type system.

struct output_device_iterator_tag : public thrust::detail::iterator_category_with_system_and_traversal<std::output_iterator_tag, thrust::device_system_tag, thrust::single_pass_traversal_tag>#
#include <iterator_categories.h>

output_device_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Output Device Iterator concept within the C++ type system.

struct forward_device_iterator_tag : public thrust::detail::iterator_category_with_system_and_traversal<std::forward_iterator_tag, thrust::device_system_tag, thrust::forward_traversal_tag>#
#include <iterator_categories.h>

forward_device_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Forward Device Iterator concept within the C++ type system.

struct bidirectional_device_iterator_tag : public thrust::detail::iterator_category_with_system_and_traversal<std::bidirectional_iterator_tag, thrust::device_system_tag, thrust::bidirectional_traversal_tag>#
#include <iterator_categories.h>

bidirectional_device_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Bidirectional Device Iterator concept within the C++ type system.

struct random_access_device_iterator_tag : public thrust::detail::iterator_category_with_system_and_traversal<std::random_access_iterator_tag, thrust::device_system_tag, thrust::random_access_traversal_tag>#
#include <iterator_categories.h>

random_access_device_iterator_tag is an empty class: it has no member functions, member variables, or nested types. It is used solely as a “tag”: a representation of the Random Access Device Iterator concept within the C++ type system.

Algorithms#

group algorithms

Functions

template<typename DerivedPolicy, typename InputIterator, typename T> __host__ __device__ InputIterator find (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator first, InputIterator last, const T &value)

find returns the first iterator i in the range [first, last) such that *i == value or last if no such iterator exists.

The algorithm’s execution is parallelized as determined by exec.

#include <thrust/find.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(4);

input[0] = 0;
input[1] = 5;
input[2] = 3;
input[3] = 7;

thrust::device_vector<int>::iterator iter;

iter = thrust::find(thrust::device, input.begin(), input.end(), 3); // returns input.first() + 2
iter = thrust::find(thrust::device, input.begin(), input.end(), 5); // returns input.first() + 1
iter = thrust::find(thrust::device, input.begin(), input.end(), 9); // returns input.end()

See also

find_if

See also

mismatch

Parameters
  • exec – The execution policy to use for parallelization.

  • first – Beginning of the sequence to search.

  • last – End of the sequence to search.

  • value – The value to find.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – is a model of Input Iterator and InputIterator's value_type is equality comparable to type T.

  • T – is a model of EqualityComparable.

Returns

The first iterator i such that *i == value or last.

template<typename InputIterator, typename T>
InputIterator find(InputIterator first, InputIterator last, const T &value)#

find returns the first iterator i in the range [first, last) such that *i == value or last if no such iterator exists.

#include <thrust/find.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(4);

input[0] = 0;
input[1] = 5;
input[2] = 3;
input[3] = 7;

thrust::device_vector<int>::iterator iter;

iter = thrust::find(input.begin(), input.end(), 3); // returns input.first() + 2
iter = thrust::find(input.begin(), input.end(), 5); // returns input.first() + 1
iter = thrust::find(input.begin(), input.end(), 9); // returns input.end()

See also

find_if

See also

mismatch

Parameters
  • first – Beginning of the sequence to search.

  • last – End of the sequence to search.

  • value – The value to find.

Template Parameters
Returns

The first iterator i such that *i == value or last.

template<typename DerivedPolicy, typename InputIterator, typename Predicate> __host__ __device__ InputIterator find_if (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator first, InputIterator last, Predicate pred)

find_if returns the first iterator i in the range [first, last) such that pred(*i) is true or last if no such iterator exists.

The algorithm’s execution is parallelized as determined by exec.

#include <thrust/find.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...

struct greater_than_four
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 4;
  }
};

struct greater_than_ten
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 10;
  }
};

...
thrust::device_vector<int> input(4);

input[0] = 0;
input[1] = 5;
input[2] = 3;
input[3] = 7;

thrust::device_vector<int>::iterator iter;

iter = thrust::find_if(thrust::device, input.begin(), input.end(), greater_than_four()); // returns input.first() + 1

iter = thrust::find_if(thrust::device, input.begin(), input.end(), greater_than_ten());  // returns input.end()

See also

find

See also

find_if_not

See also

mismatch

Parameters
  • exec – The execution policy to use for parallelization.

  • first – Beginning of the sequence to search.

  • last – End of the sequence to search.

  • pred – A predicate used to test range elements.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – is a model of Input Iterator.

  • Predicate – is a model of Predicate.

Returns

The first iterator i such that pred(*i) is true, or last.

template<typename InputIterator, typename Predicate>
InputIterator find_if(InputIterator first, InputIterator last, Predicate pred)#

find_if returns the first iterator i in the range [first, last) such that pred(*i) is true or last if no such iterator exists.

#include <thrust/find.h>
#include <thrust/device_vector.h>

struct greater_than_four
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 4;
  }
};

struct greater_than_ten
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 10;
  }
};

...
thrust::device_vector<int> input(4);

input[0] = 0;
input[1] = 5;
input[2] = 3;
input[3] = 7;

thrust::device_vector<int>::iterator iter;

iter = thrust::find_if(input.begin(), input.end(), greater_than_four()); // returns input.first() + 1

iter = thrust::find_if(input.begin(), input.end(), greater_than_ten());  // returns input.end()

See also

find

See also

find_if_not

See also

mismatch

Parameters
  • first – Beginning of the sequence to search.

  • last – End of the sequence to search.

  • pred – A predicate used to test range elements.

Template Parameters
Returns

The first iterator i such that pred(*i) is true, or last.

template<typename DerivedPolicy, typename InputIterator, typename Predicate> __host__ __device__ InputIterator find_if_not (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator first, InputIterator last, Predicate pred)

find_if_not returns the first iterator i in the range [first, last) such that pred(*i) is false or last if no such iterator exists.

The algorithm’s execution is parallelized as determined by exec.

#include <thrust/find.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...

struct greater_than_four
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 4;
  }
};

struct greater_than_ten
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 10;
  }
};

...
thrust::device_vector<int> input(4);

input[0] = 0;
input[1] = 5;
input[2] = 3;
input[3] = 7;

thrust::device_vector<int>::iterator iter;

iter = thrust::find_if_not(thrust::device, input.begin(), input.end(), greater_than_four()); // returns input.first()

iter = thrust::find_if_not(thrust::device, input.begin(), input.end(), greater_than_ten());  // returns input.first()

See also

find

See also

find_if

See also

mismatch

Parameters
  • exec – The execution policy to use for parallelization.

  • first – Beginning of the sequence to search.

  • last – End of the sequence to search.

  • pred – A predicate used to test range elements.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – is a model of Input Iterator.

  • Predicate – is a model of Predicate.

Returns

The first iterator i such that pred(*i) is false, or last.

template<typename InputIterator, typename Predicate>
InputIterator find_if_not(InputIterator first, InputIterator last, Predicate pred)#

find_if_not returns the first iterator i in the range [first, last) such that pred(*i) is false or last if no such iterator exists.

#include <thrust/find.h>
#include <thrust/device_vector.h>

struct greater_than_four
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 4;
  }
};

struct greater_than_ten
{
  __host__ __device__
  bool operator()(int x)
  {
    return x > 10;
  }
};

...
thrust::device_vector<int> input(4);

input[0] = 0;
input[1] = 5;
input[2] = 3;
input[3] = 7;

thrust::device_vector<int>::iterator iter;

iter = thrust::find_if_not(input.begin(), input.end(), greater_than_four()); // returns input.first()

iter = thrust::find_if_not(input.begin(), input.end(), greater_than_ten());  // returns input.first()

See also

find

See also

find_if

See also

mismatch

Parameters
  • first – Beginning of the sequence to search.

  • last – End of the sequence to search.

  • pred – A predicate used to test range elements.

Template Parameters
Returns

The first iterator i such that pred(*i) is false, or last.

template<typename DerivedPolicy, typename InputIterator1, typename InputIterator2> __host__ __device__ thrust::pair< InputIterator1, InputIterator2 > mismatch (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator1 first1, InputIterator1 last1, InputIterator2 first2)

mismatch finds the first position where the two ranges [first1, last1) and [first2, first2 + (last1 - first1)) differ. The two versions of mismatch use different tests for whether elements differ.

This version of mismatch finds the first iterator i in [first1, last1) such that *i == *(first2 + (i - first1)) is false. The return value is a pair whose first element is i and whose second element is *(first2 + (i - first1)). If no such iterator i exists, the return value is a pair whose first element is last1 and whose second element is *(first2 + (last1 - first1)).

The algorithm’s execution is parallelized as determined by exec.

#include <thrust/mismatch.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> vec1(4);
thrust::device_vector<int> vec2(4);

vec1[0] = 0;  vec2[0] = 0; 
vec1[1] = 5;  vec2[1] = 5;
vec1[2] = 3;  vec2[2] = 8;
vec1[3] = 7;  vec2[3] = 7;

typedef thrust::device_vector<int>::iterator Iterator;
thrust::pair<Iterator,Iterator> result;

result = thrust::mismatch(thrust::device, vec1.begin(), vec1.end(), vec2.begin());

// result.first  is vec1.begin() + 2
// result.second is vec2.begin() + 2

See also

find

See also

find_if

Parameters
  • exec – The execution policy to use for parallelization.

  • first1 – The beginning of the first sequence.

  • last1 – The end of the first sequence.

  • first2 – The beginning of the second sequence.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator1 – is a model of Input Iterator and InputIterator1's value_type is equality comparable to InputIterator2's value_type.

  • InputIterator2 – is a model of Input Iterator.

Returns

The first position where the sequences differ.

template<typename InputIterator1, typename InputIterator2>
thrust::pair<InputIterator1, InputIterator2> mismatch(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2)#

mismatch finds the first position where the two ranges [first1, last1) and [first2, first2 + (last1 - first1)) differ. The two versions of mismatch use different tests for whether elements differ.

This version of mismatch finds the first iterator i in [first1, last1) such that *i == *(first2 + (i - first1)) is false. The return value is a pair whose first element is i and whose second element is *(first2 + (i - first1)). If no such iterator i exists, the return value is a pair whose first element is last1 and whose second element is *(first2 + (last1 - first1)).

#include <thrust/mismatch.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> vec1(4);
thrust::device_vector<int> vec2(4);

vec1[0] = 0;  vec2[0] = 0; 
vec1[1] = 5;  vec2[1] = 5;
vec1[2] = 3;  vec2[2] = 8;
vec1[3] = 7;  vec2[3] = 7;

typedef thrust::device_vector<int>::iterator Iterator;
thrust::pair<Iterator,Iterator> result;

result = thrust::mismatch(vec1.begin(), vec1.end(), vec2.begin());

// result.first  is vec1.begin() + 2
// result.second is vec2.begin() + 2

See also

find

See also

find_if

Parameters
  • first1 – The beginning of the first sequence.

  • last1 – The end of the first sequence.

  • first2 – The beginning of the second sequence.

Template Parameters
  • InputIterator1 – is a model of Input Iterator and InputIterator1's value_type is equality comparable to InputIterator2's value_type.

  • InputIterator2 – is a model of Input Iterator.

Returns

The first position where the sequences differ.

template<typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename BinaryPredicate> __host__ __device__ thrust::pair< InputIterator1, InputIterator2 > mismatch (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, BinaryPredicate pred)

mismatch finds the first position where the two ranges [first1, last1) and [first2, first2 + (last1 - first1)) differ. The two versions of mismatch use different tests for whether elements differ.

This version of mismatch finds the first iterator i in [first1, last1) such that pred(*i, *(first2 + (i - first1)) is false. The return value is a pair whose first element is i and whose second element is *(first2 + (i - first1)). If no such iterator i exists, the return value is a pair whose first element is last1 and whose second element is *(first2 + (last1 - first1)).

The algorithm’s execution is parallelized as determined by exec.

#include <thrust/mismatch.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> vec1(4);
thrust::device_vector<int> vec2(4);

vec1[0] = 0;  vec2[0] = 0; 
vec1[1] = 5;  vec2[1] = 5;
vec1[2] = 3;  vec2[2] = 8;
vec1[3] = 7;  vec2[3] = 7;

typedef thrust::device_vector<int>::iterator Iterator;
thrust::pair<Iterator,Iterator> result;

result = thrust::mismatch(thrust::device, vec1.begin(), vec1.end(), vec2.begin(), thrust::equal_to<int>());

// result.first  is vec1.begin() + 2
// result.second is vec2.begin() + 2

See also

find

See also

find_if

Parameters
  • exec – The execution policy to use for parallelization.

  • first1 – The beginning of the first sequence.

  • last1 – The end of the first sequence.

  • first2 – The beginning of the second sequence.

  • pred – The binary predicate to compare elements.

Template Parameters
Returns

The first position where the sequences differ.

template<typename InputIterator1, typename InputIterator2, typename BinaryPredicate>
thrust::pair<InputIterator1, InputIterator2> mismatch(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, BinaryPredicate pred)#

mismatch finds the first position where the two ranges [first1, last1) and [first2, first2 + (last1 - first1)) differ. The two versions of mismatch use different tests for whether elements differ.

This version of mismatch finds the first iterator i in [first1, last1) such that pred(*i, *(first2 + (i - first1)) is false. The return value is a pair whose first element is i and whose second element is *(first2 + (i - first1)). If no such iterator i exists, the return value is a pair whose first element is last1 and whose second element is *(first2 + (last1 - first1)).

#include <thrust/mismatch.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> vec1(4);
thrust::device_vector<int> vec2(4);

vec1[0] = 0;  vec2[0] = 0; 
vec1[1] = 5;  vec2[1] = 5;
vec1[2] = 3;  vec2[2] = 8;
vec1[3] = 7;  vec2[3] = 7;

typedef thrust::device_vector<int>::iterator Iterator;
thrust::pair<Iterator,Iterator> result;

result = thrust::mismatch(vec1.begin(), vec1.end(), vec2.begin(), thrust::equal_to<int>());

// result.first  is vec1.begin() + 2
// result.second is vec2.begin() + 2

See also

find

See also

find_if

Parameters
  • first1 – The beginning of the first sequence.

  • last1 – The end of the first sequence.

  • first2 – The beginning of the second sequence.

  • pred – The binary predicate to compare elements.

Template Parameters
Returns

The first position where the sequences differ.

template<typename DerivedPolicy, typename ForwardIterator, typename Predicate> __host__ __device__ ForwardIterator partition_point (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, Predicate pred)

partition_point returns an iterator pointing to the end of the true partition of a partitioned range. partition_point requires the input range [first,last) to be a partition; that is, all elements which satisfy pred shall appear before those that do not.

The algorithm’s execution is parallelized as determined by exec.

#include <thrust/partition.h>
#include <thrust/execution_policy.h>

struct is_even
{
  __host__ __device__
  bool operator()(const int &x)
  {
    return (x % 2) == 0;
  }
};

...

int A[] = {2, 4, 6, 8, 10, 1, 3, 5, 7, 9};
int * B = thrust::partition_point(thrust::host, A, A + 10, is_even());
// B - A is 5
// [A, B) contains only even values

See also

partition

See also

find_if_not

Note

Though similar, partition_point is not redundant with find_if_not. partition_point's precondition provides an opportunity for a faster implemention.

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the range to consider.

  • last – The end of the range to consider.

  • pred – A function object which decides to which partition each element of the range [first, last) belongs.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator, and ForwardIterator's value_type is convertible to Predicate's argument_type.

  • Predicate – is a model of Predicate.

Returns

An iterator mid such that all_of(first, mid, pred) and none_of(mid, last, pred) are both true.

Pre

The range [first, last) shall be partitioned by pred.

template<typename ForwardIterator, typename Predicate>
ForwardIterator partition_point(ForwardIterator first, ForwardIterator last, Predicate pred)#

partition_point returns an iterator pointing to the end of the true partition of a partitioned range. partition_point requires the input range [first,last) to be a partition; that is, all elements which satisfy pred shall appear before those that do not.

#include <thrust/partition.h>

struct is_even
{
  __host__ __device__
  bool operator()(const int &x)
  {
    return (x % 2) == 0;
  }
};

...

int A[] = {2, 4, 6, 8, 10, 1, 3, 5, 7, 9};
int * B = thrust::partition_point(A, A + 10, is_even());
// B - A is 5
// [A, B) contains only even values

See also

partition

See also

find_if_not

Note

Though similar, partition_point is not redundant with find_if_not. partition_point's precondition provides an opportunity for a faster implemention.

Parameters
  • first – The beginning of the range to consider.

  • last – The end of the range to consider.

  • pred – A function object which decides to which partition each element of the range [first, last) belongs.

Template Parameters
  • ForwardIterator – is a model of Forward Iterator, and ForwardIterator's value_type is convertible to Predicate's argument_type.

  • Predicate – is a model of Predicate.

Returns

An iterator mid such that all_of(first, mid, pred) and none_of(mid, last, pred) are both true.

Pre

The range [first, last) shall be partitioned by pred.

Functions

template<typename DerivedPolicy, typename ForwardIterator, typename LessThanComparable> __host__ __device__ ForwardIterator lower_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const LessThanComparable &value)

lower_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the first position where value could be inserted without violating the ordering. This version of lower_bound uses operator< for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), *j < value.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use lower_bound to search for values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::lower_bound(thrust::device, input.begin(), input.end(), 0); // returns input.begin()
thrust::lower_bound(thrust::device, input.begin(), input.end(), 1); // returns input.begin() + 1
thrust::lower_bound(thrust::device, input.begin(), input.end(), 2); // returns input.begin() + 1
thrust::lower_bound(thrust::device, input.begin(), input.end(), 3); // returns input.begin() + 2
thrust::lower_bound(thrust::device, input.begin(), input.end(), 8); // returns input.begin() + 4
thrust::lower_bound(thrust::device, input.begin(), input.end(), 9); // returns input.end()

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • LessThanComparable – is a model of LessThanComparable.

Returns

The furthermost iterator i, such that *i < value.

template<class ForwardIterator, class LessThanComparable>
ForwardIterator lower_bound(ForwardIterator first, ForwardIterator last, const LessThanComparable &value)#

lower_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the first position where value could be inserted without violating the ordering. This version of lower_bound uses operator< for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), *j < value.

The following code snippet demonstrates how to use lower_bound to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::lower_bound(input.begin(), input.end(), 0); // returns input.begin()
thrust::lower_bound(input.begin(), input.end(), 1); // returns input.begin() + 1
thrust::lower_bound(input.begin(), input.end(), 2); // returns input.begin() + 1
thrust::lower_bound(input.begin(), input.end(), 3); // returns input.begin() + 2
thrust::lower_bound(input.begin(), input.end(), 8); // returns input.begin() + 4
thrust::lower_bound(input.begin(), input.end(), 9); // returns input.end()

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
Returns

The furthermost iterator i, such that *i < value.

template<typename DerivedPolicy, typename ForwardIterator, typename T, typename StrictWeakOrdering> __host__ __device__ ForwardIterator lower_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)

lower_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the first position where value could be inserted without violating the ordering. This version of lower_bound uses function object comp for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), comp(*j, value) is true.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use lower_bound to search for values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::lower_bound(input.begin(), input.end(), 0, thrust::less<int>()); // returns input.begin()
thrust::lower_bound(input.begin(), input.end(), 1, thrust::less<int>()); // returns input.begin() + 1
thrust::lower_bound(input.begin(), input.end(), 2, thrust::less<int>()); // returns input.begin() + 1
thrust::lower_bound(input.begin(), input.end(), 3, thrust::less<int>()); // returns input.begin() + 2
thrust::lower_bound(input.begin(), input.end(), 8, thrust::less<int>()); // returns input.begin() + 4
thrust::lower_bound(input.begin(), input.end(), 9, thrust::less<int>()); // returns input.end()

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • T – is comparable to ForwardIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Returns

The furthermost iterator i, such that comp(*i, value) is true.

template<class ForwardIterator, class T, class StrictWeakOrdering>
ForwardIterator lower_bound(ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)#

lower_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the first position where value could be inserted without violating the ordering. This version of lower_bound uses function object comp for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), comp(*j, value) is true.

The following code snippet demonstrates how to use lower_bound to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::lower_bound(input.begin(), input.end(), 0, thrust::less<int>()); // returns input.begin()
thrust::lower_bound(input.begin(), input.end(), 1, thrust::less<int>()); // returns input.begin() + 1
thrust::lower_bound(input.begin(), input.end(), 2, thrust::less<int>()); // returns input.begin() + 1
thrust::lower_bound(input.begin(), input.end(), 3, thrust::less<int>()); // returns input.begin() + 2
thrust::lower_bound(input.begin(), input.end(), 8, thrust::less<int>()); // returns input.begin() + 4
thrust::lower_bound(input.begin(), input.end(), 9, thrust::less<int>()); // returns input.end()

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
Returns

The furthermost iterator i, such that comp(*i, value) is true.

template<typename DerivedPolicy, typename ForwardIterator, typename LessThanComparable> __host__ __device__ ForwardIterator upper_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const LessThanComparable &value)

upper_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the last position where value could be inserted without violating the ordering. This version of upper_bound uses operator< for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), value < *j is false.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use upper_bound to search for values in a ordered range using the thrust::device execution policy for parallelism:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::upper_bound(thrust::device, input.begin(), input.end(), 0); // returns input.begin() + 1
thrust::upper_bound(thrust::device, input.begin(), input.end(), 1); // returns input.begin() + 1
thrust::upper_bound(thrust::device, input.begin(), input.end(), 2); // returns input.begin() + 2
thrust::upper_bound(thrust::device, input.begin(), input.end(), 3); // returns input.begin() + 2
thrust::upper_bound(thrust::device, input.begin(), input.end(), 8); // returns input.end()
thrust::upper_bound(thrust::device, input.begin(), input.end(), 9); // returns input.end()

See also

lower_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • LessThanComparable – is a model of LessThanComparable.

Returns

The furthermost iterator i, such that value < *i is false.

template<class ForwardIterator, class LessThanComparable>
ForwardIterator upper_bound(ForwardIterator first, ForwardIterator last, const LessThanComparable &value)#

upper_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the last position where value could be inserted without violating the ordering. This version of upper_bound uses operator< for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), value < *j is false.

The following code snippet demonstrates how to use upper_bound to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::upper_bound(input.begin(), input.end(), 0); // returns input.begin() + 1
thrust::upper_bound(input.begin(), input.end(), 1); // returns input.begin() + 1
thrust::upper_bound(input.begin(), input.end(), 2); // returns input.begin() + 2
thrust::upper_bound(input.begin(), input.end(), 3); // returns input.begin() + 2
thrust::upper_bound(input.begin(), input.end(), 8); // returns input.end()
thrust::upper_bound(input.begin(), input.end(), 9); // returns input.end()

See also

lower_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
Returns

The furthermost iterator i, such that value < *i is false.

template<typename DerivedPolicy, typename ForwardIterator, typename T, typename StrictWeakOrdering> __host__ __device__ ForwardIterator upper_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)

upper_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the last position where value could be inserted without violating the ordering. This version of upper_bound uses function object comp for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), comp(value, *j) is false.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use upper_bound to search for values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::upper_bound(thrust::device, input.begin(), input.end(), 0, thrust::less<int>()); // returns input.begin() + 1
thrust::upper_bound(thrust::device, input.begin(), input.end(), 1, thrust::less<int>()); // returns input.begin() + 1
thrust::upper_bound(thrust::device, input.begin(), input.end(), 2, thrust::less<int>()); // returns input.begin() + 2
thrust::upper_bound(thrust::device, input.begin(), input.end(), 3, thrust::less<int>()); // returns input.begin() + 2
thrust::upper_bound(thrust::device, input.begin(), input.end(), 8, thrust::less<int>()); // returns input.end()
thrust::upper_bound(thrust::device, input.begin(), input.end(), 9, thrust::less<int>()); // returns input.end()

See also

lower_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • T – is comparable to ForwardIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Returns

The furthermost iterator i, such that comp(value, *i) is false.

template<class ForwardIterator, class T, class StrictWeakOrdering>
ForwardIterator upper_bound(ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)#

upper_bound is a version of binary search: it attempts to find the element value in an ordered range [first, last). Specifically, it returns the last position where value could be inserted without violating the ordering. This version of upper_bound uses function object comp for comparison and returns the furthermost iterator i in [first, last) such that, for every iterator j in [first, i), comp(value, *j) is false.

The following code snippet demonstrates how to use upper_bound to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::upper_bound(input.begin(), input.end(), 0, thrust::less<int>()); // returns input.begin() + 1
thrust::upper_bound(input.begin(), input.end(), 1, thrust::less<int>()); // returns input.begin() + 1
thrust::upper_bound(input.begin(), input.end(), 2, thrust::less<int>()); // returns input.begin() + 2
thrust::upper_bound(input.begin(), input.end(), 3, thrust::less<int>()); // returns input.begin() + 2
thrust::upper_bound(input.begin(), input.end(), 8, thrust::less<int>()); // returns input.end()
thrust::upper_bound(input.begin(), input.end(), 9, thrust::less<int>()); // returns input.end()

See also

lower_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
Returns

The furthermost iterator i, such that comp(value, *i) is false.

template<typename DerivedPolicy, typename ForwardIterator, typename LessThanComparable> __host__ __device__ bool binary_search (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const LessThanComparable &value)

binary_search is a version of binary search: it attempts to find the element value in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists. Specifically, this version returns true if and only if there exists an iterator i in [first, last) such that *i < value and value < *i are both false.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use binary_search to search for values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::binary_search(thrust::device, input.begin(), input.end(), 0); // returns true
thrust::binary_search(thrust::device, input.begin(), input.end(), 1); // returns false
thrust::binary_search(thrust::device, input.begin(), input.end(), 2); // returns true
thrust::binary_search(thrust::device, input.begin(), input.end(), 3); // returns false
thrust::binary_search(thrust::device, input.begin(), input.end(), 8); // returns true
thrust::binary_search(thrust::device, input.begin(), input.end(), 9); // returns false

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • LessThanComparable – is a model of LessThanComparable.

Returns

true if an equivalent element exists in [first, last), otherwise false.

template<class ForwardIterator, class LessThanComparable>
bool binary_search(ForwardIterator first, ForwardIterator last, const LessThanComparable &value)#

binary_search is a version of binary search: it attempts to find the element value in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists. Specifically, this version returns true if and only if there exists an iterator i in [first, last) such that *i < value and value < *i are both false.

The following code snippet demonstrates how to use binary_search to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::binary_search(input.begin(), input.end(), 0); // returns true
thrust::binary_search(input.begin(), input.end(), 1); // returns false
thrust::binary_search(input.begin(), input.end(), 2); // returns true
thrust::binary_search(input.begin(), input.end(), 3); // returns false
thrust::binary_search(input.begin(), input.end(), 8); // returns true
thrust::binary_search(input.begin(), input.end(), 9); // returns false

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
Returns

true if an equivalent element exists in [first, last), otherwise false.

template<typename DerivedPolicy, typename ForwardIterator, typename T, typename StrictWeakOrdering> __host__ __device__ bool binary_search (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)

binary_search is a version of binary search: it attempts to find the element value in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists. Specifically, this version returns true if and only if there exists an iterator i in [first, last) such that comp(*i, value) and comp(value, *i) are both false.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use binary_search to search for values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::binary_search(thrust::device, input.begin(), input.end(), 0, thrust::less<int>()); // returns true
thrust::binary_search(thrust::device, input.begin(), input.end(), 1, thrust::less<int>()); // returns false
thrust::binary_search(thrust::device, input.begin(), input.end(), 2, thrust::less<int>()); // returns true
thrust::binary_search(thrust::device, input.begin(), input.end(), 3, thrust::less<int>()); // returns false
thrust::binary_search(thrust::device, input.begin(), input.end(), 8, thrust::less<int>()); // returns true
thrust::binary_search(thrust::device, input.begin(), input.end(), 9, thrust::less<int>()); // returns false

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • T – is comparable to ForwardIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Returns

true if an equivalent element exists in [first, last), otherwise false.

template<class ForwardIterator, class T, class StrictWeakOrdering>
bool binary_search(ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)#

binary_search is a version of binary search: it attempts to find the element value in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists. Specifically, this version returns true if and only if there exists an iterator i in [first, last) such that comp(*i, value) and comp(value, *i) are both false.

The following code snippet demonstrates how to use binary_search to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::binary_search(input.begin(), input.end(), 0, thrust::less<int>()); // returns true
thrust::binary_search(input.begin(), input.end(), 1, thrust::less<int>()); // returns false
thrust::binary_search(input.begin(), input.end(), 2, thrust::less<int>()); // returns true
thrust::binary_search(input.begin(), input.end(), 3, thrust::less<int>()); // returns false
thrust::binary_search(input.begin(), input.end(), 8, thrust::less<int>()); // returns true
thrust::binary_search(input.begin(), input.end(), 9, thrust::less<int>()); // returns false

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
Returns

true if an equivalent element exists in [first, last), otherwise false.

template<typename DerivedPolicy, typename ForwardIterator, typename LessThanComparable> __host__ __device__ thrust::pair< ForwardIterator, ForwardIterator > equal_range (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const LessThanComparable &value)

equal_range is a version of binary search: it attempts to find the element value in an ordered range [first, last). The value returned by equal_range is essentially a combination of the values returned by lower_bound and upper_bound: it returns a pair of iterators i and j such that i is the first position where value could be inserted without violating the ordering and j is the last position where value could be inserted without violating the ordering. It follows that every element in the range [i, j) is equivalent to value, and that [i, j) is the largest subrange of [first, last) that has this property.

This version of equal_range returns a pair of iterators [i, j), where i is the furthermost iterator in [first, last) such that, for every iterator k in [first, i), *k < value. j is the furthermost iterator in [first, last) such that, for every iterator k in [first, j), value < *k is false. For every iterator k in [i, j), neither value < *k nor *k < value is true.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use equal_range to search for values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::equal_range(thrust::device, input.begin(), input.end(), 0); // returns [input.begin(), input.begin() + 1)
thrust::equal_range(thrust::device, input.begin(), input.end(), 1); // returns [input.begin() + 1, input.begin() + 1)
thrust::equal_range(thrust::device, input.begin(), input.end(), 2); // returns [input.begin() + 1, input.begin() + 2)
thrust::equal_range(thrust::device, input.begin(), input.end(), 3); // returns [input.begin() + 2, input.begin() + 2)
thrust::equal_range(thrust::device, input.begin(), input.end(), 8); // returns [input.begin() + 4, input.end)
thrust::equal_range(thrust::device, input.begin(), input.end(), 9); // returns [input.end(), input.end)

See also

lower_bound

See also

upper_bound

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • LessThanComparable – is a model of LessThanComparable.

Returns

A pair of iterators [i, j) that define the range of equivalent elements.

template<class ForwardIterator, class LessThanComparable>
thrust::pair<ForwardIterator, ForwardIterator> equal_range(ForwardIterator first, ForwardIterator last, const LessThanComparable &value)#

equal_range is a version of binary search: it attempts to find the element value in an ordered range [first, last). The value returned by equal_range is essentially a combination of the values returned by lower_bound and upper_bound: it returns a pair of iterators i and j such that i is the first position where value could be inserted without violating the ordering and j is the last position where value could be inserted without violating the ordering. It follows that every element in the range [i, j) is equivalent to value, and that [i, j) is the largest subrange of [first, last) that has this property.

This version of equal_range returns a pair of iterators [i, j), where i is the furthermost iterator in [first, last) such that, for every iterator k in [first, i), *k < value. j is the furthermost iterator in [first, last) such that, for every iterator k in [first, j), value < *k is false. For every iterator k in [i, j), neither value < *k nor *k < value is true.

The following code snippet demonstrates how to use equal_range to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::equal_range(input.begin(), input.end(), 0); // returns [input.begin(), input.begin() + 1)
thrust::equal_range(input.begin(), input.end(), 1); // returns [input.begin() + 1, input.begin() + 1)
thrust::equal_range(input.begin(), input.end(), 2); // returns [input.begin() + 1, input.begin() + 2)
thrust::equal_range(input.begin(), input.end(), 3); // returns [input.begin() + 2, input.begin() + 2)
thrust::equal_range(input.begin(), input.end(), 8); // returns [input.begin() + 4, input.end)
thrust::equal_range(input.begin(), input.end(), 9); // returns [input.end(), input.end)

See also

lower_bound

See also

upper_bound

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

Template Parameters
Returns

A pair of iterators [i, j) that define the range of equivalent elements.

template<typename DerivedPolicy, typename ForwardIterator, typename T, typename StrictWeakOrdering> __host__ __device__ thrust::pair< ForwardIterator, ForwardIterator > equal_range (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)

equal_range is a version of binary search: it attempts to find the element value in an ordered range [first, last). The value returned by equal_range is essentially a combination of the values returned by lower_bound and upper_bound: it returns a pair of iterators i and j such that i is the first position where value could be inserted without violating the ordering and j is the last position where value could be inserted without violating the ordering. It follows that every element in the range [i, j) is equivalent to value, and that [i, j) is the largest subrange of [first, last) that has this property.

This version of equal_range returns a pair of iterators [i, j). i is the furthermost iterator in [first, last) such that, for every iterator k in [first, i), comp(*k, value) is true. j is the furthermost iterator in [first, last) such that, for every iterator k in [first, last), comp(value, *k) is false. For every iterator k in [i, j), neither comp(value, *k) nor comp(*k, value) is true.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use equal_range to search for values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::equal_range(thrust::device, input.begin(), input.end(), 0, thrust::less<int>()); // returns [input.begin(), input.begin() + 1)
thrust::equal_range(thrust::device, input.begin(), input.end(), 1, thrust::less<int>()); // returns [input.begin() + 1, input.begin() + 1)
thrust::equal_range(thrust::device, input.begin(), input.end(), 2, thrust::less<int>()); // returns [input.begin() + 1, input.begin() + 2)
thrust::equal_range(thrust::device, input.begin(), input.end(), 3, thrust::less<int>()); // returns [input.begin() + 2, input.begin() + 2)
thrust::equal_range(thrust::device, input.begin(), input.end(), 8, thrust::less<int>()); // returns [input.begin() + 4, input.end)
thrust::equal_range(thrust::device, input.begin(), input.end(), 9, thrust::less<int>()); // returns [input.end(), input.end)

See also

lower_bound

See also

upper_bound

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • T – is comparable to ForwardIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Returns

A pair of iterators [i, j) that define the range of equivalent elements.

template<class ForwardIterator, class T, class StrictWeakOrdering>
thrust::pair<ForwardIterator, ForwardIterator> equal_range(ForwardIterator first, ForwardIterator last, const T &value, StrictWeakOrdering comp)#

equal_range is a version of binary search: it attempts to find the element value in an ordered range [first, last). The value returned by equal_range is essentially a combination of the values returned by lower_bound and upper_bound: it returns a pair of iterators i and j such that i is the first position where value could be inserted without violating the ordering and j is the last position where value could be inserted without violating the ordering. It follows that every element in the range [i, j) is equivalent to value, and that [i, j) is the largest subrange of [first, last) that has this property.

This version of equal_range returns a pair of iterators [i, j). i is the furthermost iterator in [first, last) such that, for every iterator k in [first, i), comp(*k, value) is true. j is the furthermost iterator in [first, last) such that, for every iterator k in [first, last), comp(value, *k) is false. For every iterator k in [i, j), neither comp(value, *k) nor comp(*k, value) is true.

The following code snippet demonstrates how to use equal_range to search for values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::equal_range(input.begin(), input.end(), 0, thrust::less<int>()); // returns [input.begin(), input.begin() + 1)
thrust::equal_range(input.begin(), input.end(), 1, thrust::less<int>()); // returns [input.begin() + 1, input.begin() + 1)
thrust::equal_range(input.begin(), input.end(), 2, thrust::less<int>()); // returns [input.begin() + 1, input.begin() + 2)
thrust::equal_range(input.begin(), input.end(), 3, thrust::less<int>()); // returns [input.begin() + 2, input.begin() + 2)
thrust::equal_range(input.begin(), input.end(), 8, thrust::less<int>()); // returns [input.begin() + 4, input.end)
thrust::equal_range(input.begin(), input.end(), 9, thrust::less<int>()); // returns [input.end(), input.end)

See also

lower_bound

See also

upper_bound

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • value – The value to be searched.

  • comp – The comparison operator.

Template Parameters
Returns

A pair of iterators [i, j) that define the range of equivalent elements.

Functions

template<typename DerivedPolicy, typename ForwardIterator, typename InputIterator, typename OutputIterator> __host__ __device__ OutputIterator lower_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result)

lower_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of first position where value could be inserted without violating the ordering.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use lower_bound to search for multiple values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::lower_bound(thrust::device,
                    input.begin(), input.end(),
                    values.begin(), values.end(),
                    output.begin());

// output is now [0, 1, 1, 2, 4, 5]

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • InputIterator – is a model of Input Iterator. and InputIterator's value_type is LessThanComparable.

  • OutputIterator – is a model of Output Iterator. and ForwardIterator's difference_type is convertible to OutputIterator's value_type.

Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<class ForwardIterator, class InputIterator, class OutputIterator>
OutputIterator lower_bound(ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result)#

lower_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of first position where value could be inserted without violating the ordering.

The following code snippet demonstrates how to use lower_bound to search for multiple values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::lower_bound(input.begin(), input.end(),
                    values.begin(), values.end(),
                    output.begin());

// output is now [0, 1, 1, 2, 4, 5]

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

Template Parameters
Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<typename DerivedPolicy, typename ForwardIterator, typename InputIterator, typename OutputIterator, typename StrictWeakOrdering> __host__ __device__ OutputIterator lower_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result, StrictWeakOrdering comp)

lower_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of first position where value could be inserted without violating the ordering. This version of lower_bound uses function object comp for comparison.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use lower_bound to search for multiple values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::lower_bound(input.begin(), input.end(),
                    values.begin(), values.end(), 
                    output.begin(),
                    thrust::less<int>());

// output is now [0, 1, 1, 2, 4, 5]

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

  • comp – The comparison operator.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • InputIterator – is a model of Input Iterator. and InputIterator's value_type is comparable to ForwardIterator's value_type.

  • OutputIterator – is a model of Output Iterator. and ForwardIterator's difference_type is convertible to OutputIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<class ForwardIterator, class InputIterator, class OutputIterator, class StrictWeakOrdering>
OutputIterator lower_bound(ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result, StrictWeakOrdering comp)#

lower_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of first position where value could be inserted without violating the ordering. This version of lower_bound uses function object comp for comparison.

The following code snippet demonstrates how to use lower_bound to search for multiple values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::lower_bound(input.begin(), input.end(),
                    values.begin(), values.end(), 
                    output.begin(),
                    thrust::less<int>());

// output is now [0, 1, 1, 2, 4, 5]

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

  • comp – The comparison operator.

Template Parameters
  • ForwardIterator – is a model of Forward Iterator.

  • InputIterator – is a model of Input Iterator. and InputIterator's value_type is comparable to ForwardIterator's value_type.

  • OutputIterator – is a model of Output Iterator. and ForwardIterator's difference_type is convertible to OutputIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<typename DerivedPolicy, typename ForwardIterator, typename InputIterator, typename OutputIterator> __host__ __device__ OutputIterator upper_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result)

upper_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of last position where value could be inserted without violating the ordering.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use upper_bound to search for multiple values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::upper_bound(thrust::device,
                    input.begin(), input.end(),
                    values.begin(), values.end(),
                    output.begin());

// output is now [1, 1, 2, 2, 5, 5]

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • InputIterator – is a model of Input Iterator. and InputIterator's value_type is LessThanComparable.

  • OutputIterator – is a model of Output Iterator. and ForwardIterator's difference_type is convertible to OutputIterator's value_type.

Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<class ForwardIterator, class InputIterator, class OutputIterator>
OutputIterator upper_bound(ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result)#

upper_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of last position where value could be inserted without violating the ordering.

The following code snippet demonstrates how to use upper_bound to search for multiple values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::upper_bound(input.begin(), input.end(),
                    values.begin(), values.end(),
                    output.begin());

// output is now [1, 1, 2, 2, 5, 5]

See also

upper_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

Template Parameters
Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<typename DerivedPolicy, typename ForwardIterator, typename InputIterator, typename OutputIterator, typename StrictWeakOrdering> __host__ __device__ OutputIterator upper_bound (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result, StrictWeakOrdering comp)

upper_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of first position where value could be inserted without violating the ordering. This version of upper_bound uses function object comp for comparison.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use upper_bound to search for multiple values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::upper_bound(thrust::device,
                    input.begin(), input.end(),
                    values.begin(), values.end(), 
                    output.begin(),
                    thrust::less<int>());

// output is now [1, 1, 2, 2, 5, 5]

See also

lower_bound

See also

equal_range

See also

binary_search

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

  • comp – The comparison operator.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • InputIterator – is a model of Input Iterator. and InputIterator's value_type is comparable to ForwardIterator's value_type.

  • OutputIterator – is a model of Output Iterator. and ForwardIterator's difference_type is convertible to OutputIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<class ForwardIterator, class InputIterator, class OutputIterator, class StrictWeakOrdering>
OutputIterator upper_bound(ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result, StrictWeakOrdering comp)#

upper_bound is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). Specifically, it returns the index of first position where value could be inserted without violating the ordering. This version of upper_bound uses function object comp for comparison.

The following code snippet demonstrates how to use upper_bound to search for multiple values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<unsigned int> output(6);

thrust::upper_bound(input.begin(), input.end(),
                    values.begin(), values.end(), 
                    output.begin(),
                    thrust::less<int>());

// output is now [1, 1, 2, 2, 5, 5]

See also

lower_bound

See also

equal_range

See also

binary_search

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

  • comp – The comparison operator.

Template Parameters
  • ForwardIterator – is a model of Forward Iterator.

  • InputIterator – is a model of Input Iterator. and InputIterator's value_type is comparable to ForwardIterator's value_type.

  • OutputIterator – is a model of Output Iterator. and ForwardIterator's difference_type is convertible to OutputIterator's value_type.

  • StrictWeakOrdering – is a model of Strict Weak Ordering.

Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<typename DerivedPolicy, typename ForwardIterator, typename InputIterator, typename OutputIterator> __host__ __device__ OutputIterator binary_search (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result)

binary_search is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use binary_search to search for multiple values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<bool> output(6);

thrust::binary_search(thrust::device,
                      input.begin(), input.end(),
                      values.begin(), values.end(),
                      output.begin());

// output is now [true, false, true, false, true, false]

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator – is a model of Forward Iterator.

  • InputIterator – is a model of Input Iterator. and InputIterator's value_type is LessThanComparable.

  • OutputIterator – is a model of Output Iterator. and bool is convertible to OutputIterator's value_type.

Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<class ForwardIterator, class InputIterator, class OutputIterator>
OutputIterator binary_search(ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result)#

binary_search is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists.

The following code snippet demonstrates how to use binary_search to search for multiple values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<bool> output(6);

thrust::binary_search(input.begin(), input.end(),
                      values.begin(), values.end(),
                      output.begin());

// output is now [true, false, true, false, true, false]

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

Template Parameters
Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<typename DerivedPolicy, typename ForwardIterator, typename InputIterator, typename OutputIterator, typename StrictWeakOrdering> __host__ __device__ OutputIterator binary_search (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result, StrictWeakOrdering comp)

binary_search is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists. This version of binary_search uses function object comp for comparison.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use binary_search to search for multiple values in a ordered range using the thrust::device execution policy for parallelization:

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<bool> output(6);

thrust::binary_search(thrust::device,
                      input.begin(), input.end(),
                      values.begin(), values.end(),
                      output.begin(),
                      thrust::less<T>());

// output is now [true, false, true, false, true, false]

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

  • comp – The comparison operator.

Template Parameters
Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

template<class ForwardIterator, class InputIterator, class OutputIterator, class StrictWeakOrdering>
OutputIterator binary_search(ForwardIterator first, ForwardIterator last, InputIterator values_first, InputIterator values_last, OutputIterator result, StrictWeakOrdering comp)#

binary_search is a vectorized version of binary search: for each iterator v in [values_first, values_last) it attempts to find the value *v in an ordered range [first, last). It returns true if an element that is equivalent to value is present in [first, last) and false if no such element exists. This version of binary_search uses function object comp for comparison.

The following code snippet demonstrates how to use binary_search to search for multiple values in a ordered range.

#include <thrust/binary_search.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
...
thrust::device_vector<int> input(5);

input[0] = 0;
input[1] = 2;
input[2] = 5;
input[3] = 7;
input[4] = 8;

thrust::device_vector<int> values(6);
values[0] = 0; 
values[1] = 1;
values[2] = 2;
values[3] = 3;
values[4] = 8;
values[5] = 9;

thrust::device_vector<bool> output(6);

thrust::binary_search(input.begin(), input.end(),
                      values.begin(), values.end(),
                      output.begin(),
                      thrust::less<T>());

// output is now [true, false, true, false, true, false]

See also

lower_bound

See also

upper_bound

See also

equal_range

Parameters
  • first – The beginning of the ordered sequence.

  • last – The end of the ordered sequence.

  • values_first – The beginning of the search values sequence.

  • values_last – The end of the search values sequence.

  • result – The beginning of the output sequence.

  • comp – The comparison operator.

Template Parameters
Pre

The ranges [first,last) and [result, result + (last - first)) shall not overlap.

Functions

template<typename DerivedPolicy, typename InputIterator, typename OutputIterator> __host__ __device__ OutputIterator copy (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator first, InputIterator last, OutputIterator result)

copy copies elements from the range [first, last) to the range [result, result + (last - first)). That is, it performs the assignments *result = *first, *(result + 1) = *(first + 1), and so on. Generally, for every integer n from 0 to last - first, copy performs the assignment *(result + n) = *(first + n). Unlike std::copy, copy offers no guarantee on order of operation. As a result, calling copy with overlapping source and destination ranges has undefined behavior.

The return value is result + (last - first).

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use copy to copy from one range to another using the thrust::device parallelization policy:

#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...

thrust::device_vector<int> vec0(100);
thrust::device_vector<int> vec1(100);
...

thrust::copy(thrust::device, vec0.begin(), vec0.end(), vec1.begin());

// vec1 is now a copy of vec0
Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the sequence to copy.

  • last – The end of the sequence to copy.

  • result – The destination sequence.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – must be a model of Input Iterator and InputIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

Returns

The end of the destination sequence.

Pre

result may be equal to first, but result shall not be in the range [first, last) otherwise.

template<typename DerivedPolicy, typename InputIterator, typename Size, typename OutputIterator> __host__ __device__ OutputIterator copy_n (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator first, Size n, OutputIterator result)

copy_n copies elements from the range [first, first + n) to the range [result, result + n). That is, it performs the assignments *result = *first, *(result + 1) = *(first + 1), and so on. Generally, for every integer i from 0 to n, copy performs the assignment *(result + i) = *(first + i). Unlike std::copy_n, copy_n offers no guarantee on order of operation. As a result, calling copy_n with overlapping source and destination ranges has undefined behavior.

The return value is result + n.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use copy to copy from one range to another using the thrust::device parallelization policy:

#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
size_t n = 100;
thrust::device_vector<int> vec0(n);
thrust::device_vector<int> vec1(n);
...
thrust::copy_n(thrust::device, vec0.begin(), n, vec1.begin());

// vec1 is now a copy of vec0

See also

thrust::copy

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The beginning of the range to copy.

  • n – The number of elements to copy.

  • result – The beginning destination range.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – must be a model of Input Iterator and InputIterator's value_type must be convertible to OutputIterator's value_type.

  • Size – is an integral type.

  • OutputIterator – must be a model of Output Iterator.

Returns

The end of the destination range.

Pre

result may be equal to first, but result shall not be in the range [first, first + n) otherwise.

template<typename InputIterator, typename OutputIterator>
OutputIterator copy(InputIterator first, InputIterator last, OutputIterator result)#

copy copies elements from the range [first, last) to the range [result, result + (last - first)). That is, it performs the assignments *result = *first, *(result + 1) = *(first + 1), and so on. Generally, for every integer n from 0 to last - first, copy performs the assignment *(result + n) = *(first + n). Unlike std::copy, copy offers no guarantee on order of operation. As a result, calling copy with overlapping source and destination ranges has undefined behavior.

The return value is result + (last - first).

The following code snippet demonstrates how to use copy to copy from one range to another.

#include <thrust/copy.h>
#include <thrust/device_vector.h>
...

thrust::device_vector<int> vec0(100);
thrust::device_vector<int> vec1(100);
...

thrust::copy(vec0.begin(), vec0.end(),
             vec1.begin());

// vec1 is now a copy of vec0
Parameters
  • first – The beginning of the sequence to copy.

  • last – The end of the sequence to copy.

  • result – The destination sequence.

Template Parameters
  • InputIterator – must be a model of Input Iterator and InputIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

Returns

The end of the destination sequence.

Pre

result may be equal to first, but result shall not be in the range [first, last) otherwise.

template<typename InputIterator, typename Size, typename OutputIterator>
OutputIterator copy_n(InputIterator first, Size n, OutputIterator result)#

copy_n copies elements from the range [first, first + n) to the range [result, result + n). That is, it performs the assignments *result = *first, *(result + 1) = *(first + 1), and so on. Generally, for every integer i from 0 to n, copy performs the assignment *(result + i) = *(first + i). Unlike std::copy_n, copy_n offers no guarantee on order of operation. As a result, calling copy_n with overlapping source and destination ranges has undefined behavior.

The return value is result + n.

The following code snippet demonstrates how to use copy to copy from one range to another.

#include <thrust/copy.h>
#include <thrust/device_vector.h>
...
size_t n = 100;
thrust::device_vector<int> vec0(n);
thrust::device_vector<int> vec1(n);
...
thrust::copy_n(vec0.begin(), n, vec1.begin());

// vec1 is now a copy of vec0

See also

thrust::copy

Parameters
  • first – The beginning of the range to copy.

  • n – The number of elements to copy.

  • result – The beginning destination range.

Template Parameters
  • InputIterator – must be a model of Input Iterator and InputIterator's value_type must be convertible to OutputIterator's value_type.

  • Size – is an integral type.

  • OutputIterator – must be a model of Output Iterator.

Returns

The end of the destination range.

Pre

result may be equal to first, but result shall not be in the range [first, first + n) otherwise.

template<typename DerivedPolicy, typename ForwardIterator1, typename ForwardIterator2> __host__ __device__ ForwardIterator2 swap_ranges (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, ForwardIterator1 first1, ForwardIterator1 last1, ForwardIterator2 first2)

swap_ranges swaps each of the elements in the range [first1, last1) with the corresponding element in the range [first2, first2 + (last1 - first1)). That is, for each integer n such that 0 <= n < (last1 - first1), it swaps *(first1 + n) and *(first2 + n). The return value is first2 + (last1 - first1).

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use swap_ranges to swap the contents of two thrust::device_vectors using the thrust::device execution policy for parallelization:

#include <thrust/swap.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
thrust::device_vector<int> v1(2), v2(2);
v1[0] = 1;
v1[1] = 2;
v2[0] = 3;
v2[1] = 4;

thrust::swap_ranges(thrust::device, v1.begin(), v1.end(), v2.begin());

// v1[0] == 3, v1[1] == 4, v2[0] == 1, v2[1] == 2

See also

swap

Parameters
  • exec – The execution policy to use for parallelization.

  • first1 – The beginning of the first sequence to swap.

  • last1 – One position past the last element of the first sequence to swap.

  • first2 – The beginning of the second sequence to swap.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • ForwardIterator1 – is a model of Forward Iterator, and ForwardIterator1's value_type must be convertible to ForwardIterator2's value_type.

  • ForwardIterator2 – is a model of Forward Iterator, and ForwardIterator2's value_type must be convertible to ForwardIterator1's value_type.

Returns

An iterator pointing to one position past the last element of the second sequence to swap.

Pre

first1 may equal first2, but the range [first1, last1) shall not overlap the range [first2, first2 + (last1 - first1)) otherwise.

template<typename ForwardIterator1, typename ForwardIterator2>
ForwardIterator2 swap_ranges(ForwardIterator1 first1, ForwardIterator1 last1, ForwardIterator2 first2)#

swap_ranges swaps each of the elements in the range [first1, last1) with the corresponding element in the range [first2, first2 + (last1 - first1)). That is, for each integer n such that 0 <= n < (last1 - first1), it swaps *(first1 + n) and *(first2 + n). The return value is first2 + (last1 - first1).

The following code snippet demonstrates how to use swap_ranges to swap the contents of two thrust::device_vectors.

#include <thrust/swap.h>
#include <thrust/device_vector.h>
...
thrust::device_vector<int> v1(2), v2(2);
v1[0] = 1;
v1[1] = 2;
v2[0] = 3;
v2[1] = 4;

thrust::swap_ranges(v1.begin(), v1.end(), v2.begin());

// v1[0] == 3, v1[1] == 4, v2[0] == 1, v2[1] == 2

See also

swap

Parameters
  • first1 – The beginning of the first sequence to swap.

  • last1 – One position past the last element of the first sequence to swap.

  • first2 – The beginning of the second sequence to swap.

Template Parameters
  • ForwardIterator1 – is a model of Forward Iterator, and ForwardIterator1's value_type must be convertible to ForwardIterator2's value_type.

  • ForwardIterator2 – is a model of Forward Iterator, and ForwardIterator2's value_type must be convertible to ForwardIterator1's value_type.

Returns

An iterator pointing to one position past the last element of the second sequence to swap.

Pre

first1 may equal first2, but the range [first1, last1) shall not overlap the range [first2, first2 + (last1 - first1)) otherwise.

template<typename DerivedPolicy, typename InputIterator, typename ForwardIterator> __host__ __device__ ForwardIterator uninitialized_copy (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator first, InputIterator last, ForwardIterator result)

In thrust, the function thrust::device_new allocates memory for an object and then creates an object at that location by calling a constructor. Occasionally, however, it is useful to separate those two operations. If each iterator in the range [result, result + (last - first)) points to uninitialized memory, then uninitialized_copy creates a copy of [first, last) in that range. That is, for each iterator i in the input, uninitialized_copy creates a copy of *i in the location pointed to by the corresponding iterator in the output range by ForwardIterator's value_type's copy constructor with *i as its argument.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use uninitialized_copy to initialize a range of uninitialized memory using the thrust::device execution policy for parallelization:

#include <thrust/uninitialized_copy.h>
#include <thrust/device_malloc.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>

struct Int
{
  __host__ __device__
  Int(int x) : val(x) {}
  int val;
};  
...
const int N = 137;

Int val(46);
thrust::device_vector<Int> input(N, val);
thrust::device_ptr<Int> array = thrust::device_malloc<Int>(N);
thrust::uninitialized_copy(thrust::device, input.begin(), input.end(), array);

// Int x = array[i];
// x.val == 46 for all 0 <= i < N

See also

copy

See also

device_new

See also

device_malloc

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The first element of the input range to copy from.

  • last – The last element of the input range to copy from.

  • result – The first element of the output range to copy to.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – is a model of Input Iterator.

  • ForwardIterator – is a model of Forward Iterator, ForwardIterator is mutable, and ForwardIterator's value_type has a constructor that takes a single argument whose type is InputIterator's value_type.

Returns

An iterator pointing to the last element of the output range.

Pre

first may equal result, but the range [first, last) and the range [result, result + (last - first)) shall not overlap otherwise.

template<typename InputIterator, typename ForwardIterator>
ForwardIterator uninitialized_copy(InputIterator first, InputIterator last, ForwardIterator result)#

In thrust, the function thrust::device_new allocates memory for an object and then creates an object at that location by calling a constructor. Occasionally, however, it is useful to separate those two operations. If each iterator in the range [result, result + (last - first)) points to uninitialized memory, then uninitialized_copy creates a copy of [first, last) in that range. That is, for each iterator i in the input, uninitialized_copy creates a copy of *i in the location pointed to by the corresponding iterator in the output range by ForwardIterator's value_type's copy constructor with *i as its argument.

The following code snippet demonstrates how to use uninitialized_copy to initialize a range of uninitialized memory.

#include <thrust/uninitialized_copy.h>
#include <thrust/device_malloc.h>
#include <thrust/device_vector.h>

struct Int
{
  __host__ __device__
  Int(int x) : val(x) {}
  int val;
};  
...
const int N = 137;

Int val(46);
thrust::device_vector<Int> input(N, val);
thrust::device_ptr<Int> array = thrust::device_malloc<Int>(N);
thrust::uninitialized_copy(input.begin(), input.end(), array);

// Int x = array[i];
// x.val == 46 for all 0 <= i < N

See also

copy

See also

device_new

See also

device_malloc

Parameters
  • first – The first element of the input range to copy from.

  • last – The last element of the input range to copy from.

  • result – The first element of the output range to copy to.

Template Parameters
  • InputIterator – is a model of Input Iterator.

  • ForwardIterator – is a model of Forward Iterator, ForwardIterator is mutable, and ForwardIterator's value_type has a constructor that takes a single argument whose type is InputIterator's value_type.

Returns

An iterator pointing to the last element of the output range.

Pre

first may equal result, but the range [first, last) and the range [result, result + (last - first)) shall not overlap otherwise.

template<typename DerivedPolicy, typename InputIterator, typename Size, typename ForwardIterator> __host__ __device__ ForwardIterator uninitialized_copy_n (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator first, Size n, ForwardIterator result)

In thrust, the function thrust::device_new allocates memory for an object and then creates an object at that location by calling a constructor. Occasionally, however, it is useful to separate those two operations. If each iterator in the range [result, result + n) points to uninitialized memory, then uninitialized_copy_n creates a copy of [first, first + n) in that range. That is, for each iterator i in the input, uninitialized_copy_n creates a copy of *i in the location pointed to by the corresponding iterator in the output range by InputIterator's value_type's copy constructor with *i as its argument.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use uninitialized_copy to initialize a range of uninitialized memory using the thrust::device execution policy for parallelization:

#include <thrust/uninitialized_copy.h>
#include <thrust/device_malloc.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>

struct Int
{
  __host__ __device__
  Int(int x) : val(x) {}
  int val;
};  
...
const int N = 137;

Int val(46);
thrust::device_vector<Int> input(N, val);
thrust::device_ptr<Int> array = thrust::device_malloc<Int>(N);
thrust::uninitialized_copy_n(thrust::device, input.begin(), N, array);

// Int x = array[i];
// x.val == 46 for all 0 <= i < N

See also

copy

See also

device_new

See also

device_malloc

Parameters
  • exec – The execution policy to use for parallelization.

  • first – The first element of the input range to copy from.

  • n – The number of elements to copy.

  • result – The first element of the output range to copy to.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – is a model of Input Iterator.

  • Size – is an integral type.

  • ForwardIterator – is a model of Forward Iterator, ForwardIterator is mutable, and ForwardIterator's value_type has a constructor that takes a single argument whose type is InputIterator's value_type.

Returns

An iterator pointing to the last element of the output range.

Pre

first may equal result, but the range [first, first + n) and the range [result, result + n) shall not overlap otherwise.

template<typename InputIterator, typename Size, typename ForwardIterator>
ForwardIterator uninitialized_copy_n(InputIterator first, Size n, ForwardIterator result)#

In thrust, the function thrust::device_new allocates memory for an object and then creates an object at that location by calling a constructor. Occasionally, however, it is useful to separate those two operations. If each iterator in the range [result, result + n) points to uninitialized memory, then uninitialized_copy_n creates a copy of [first, first + n) in that range. That is, for each iterator i in the input, uninitialized_copy_n creates a copy of *i in the location pointed to by the corresponding iterator in the output range by InputIterator's value_type's copy constructor with *i as its argument.

The following code snippet demonstrates how to use uninitialized_copy to initialize a range of uninitialized memory.

#include <thrust/uninitialized_copy.h>
#include <thrust/device_malloc.h>
#include <thrust/device_vector.h>

struct Int
{
  __host__ __device__
  Int(int x) : val(x) {}
  int val;
};  
...
const int N = 137;

Int val(46);
thrust::device_vector<Int> input(N, val);
thrust::device_ptr<Int> array = thrust::device_malloc<Int>(N);
thrust::uninitialized_copy_n(input.begin(), N, array);

// Int x = array[i];
// x.val == 46 for all 0 <= i < N

See also

copy

See also

device_new

See also

device_malloc

Parameters
  • first – The first element of the input range to copy from.

  • n – The number of elements to copy.

  • result – The first element of the output range to copy to.

Template Parameters
  • InputIterator – is a model of Input Iterator.

  • Size – is an integral type.

  • ForwardIterator – is a model of Forward Iterator, ForwardIterator is mutable, and ForwardIterator's value_type has a constructor that takes a single argument whose type is InputIterator's value_type.

Returns

An iterator pointing to the last element of the output range.

Pre

first may equal result, but the range [first, first + n) and the range [result, result + n) shall not overlap otherwise.

Functions

template<typename DerivedPolicy, typename InputIterator, typename RandomAccessIterator, typename OutputIterator> __host__ __device__ OutputIterator gather (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator map_first, InputIterator map_last, RandomAccessIterator input_first, OutputIterator result)

gather copies elements from a source array into a destination range according to a map. For each input iterator i in the range [map_first, map_last), the value input_first[*i] is assigned to *(result + (i - map_first)). RandomAccessIterator must permit random access.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use gather to reorder a range using the thrust::device execution policy for parallelization:

Remark

gather is the inverse of thrust::scatter.

#include <thrust/gather.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
// mark even indices with a 1; odd indices with a 0
int values[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::device_vector<int> d_values(values, values + 10);

// gather all even indices into the first half of the range
// and odd indices to the last half of the range
int map[10]   = {0, 2, 4, 6, 8, 1, 3, 5, 7, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10);
thrust::gather(thrust::device,
               d_map.begin(), d_map.end(),
               d_values.begin(),
               d_output.begin());
// d_output is now {1, 1, 1, 1, 1, 0, 0, 0, 0, 0}
Parameters
  • exec – The execution policy to use for parallelization.

  • map_first – Beginning of the range of gather locations.

  • map_last – End of the range of gather locations.

  • input_first – Beginning of the source range.

  • result – Beginning of the destination range.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator – must be a model of Input Iterator and InputIterator's value_type must be convertible to RandomAccessIterator's difference_type.

  • RandomAccessIterator – must be a model of Random Access Iterator and RandomAccessIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

Pre

The range [map_first, map_last) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The input data shall not overlap the range [result, result + (map_last - map_first)).

template<typename InputIterator, typename RandomAccessIterator, typename OutputIterator>
OutputIterator gather(InputIterator map_first, InputIterator map_last, RandomAccessIterator input_first, OutputIterator result)#

gather copies elements from a source array into a destination range according to a map. For each input iterator i in the range [map_first, map_last), the value input_first[*i] is assigned to *(result + (i - map_first)). RandomAccessIterator must permit random access.

The following code snippet demonstrates how to use gather to reorder a range.

Remark

gather is the inverse of thrust::scatter.

#include <thrust/gather.h>
#include <thrust/device_vector.h>
...
// mark even indices with a 1; odd indices with a 0
int values[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::device_vector<int> d_values(values, values + 10);

// gather all even indices into the first half of the range
// and odd indices to the last half of the range
int map[10]   = {0, 2, 4, 6, 8, 1, 3, 5, 7, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10);
thrust::gather(d_map.begin(), d_map.end(),
               d_values.begin(),
               d_output.begin());
// d_output is now {1, 1, 1, 1, 1, 0, 0, 0, 0, 0}
Parameters
  • map_first – Beginning of the range of gather locations.

  • map_last – End of the range of gather locations.

  • input_first – Beginning of the source range.

  • result – Beginning of the destination range.

Template Parameters
  • InputIterator – must be a model of Input Iterator and InputIterator's value_type must be convertible to RandomAccessIterator's difference_type.

  • RandomAccessIterator – must be a model of Random Access Iterator and RandomAccessIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

Pre

The range [map_first, map_last) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The input data shall not overlap the range [result, result + (map_last - map_first)).

template<typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename RandomAccessIterator, typename OutputIterator> __host__ __device__ OutputIterator gather_if (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator1 map_first, InputIterator1 map_last, InputIterator2 stencil, RandomAccessIterator input_first, OutputIterator result)

gather_if conditionally copies elements from a source array into a destination range according to a map. For each input iterator i in the range [map_first, map_last), such that the value of *(stencil + (i - map_first)) is true, the value input_first[*i] is assigned to *(result + (i - map_first)). RandomAccessIterator must permit random access.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use gather_if to gather selected values from an input range using the thrust::device execution policy:

Remark

gather_if is the inverse of scatter_if.

#include <thrust/gather.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...

int values[10] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
thrust::device_vector<int> d_values(values, values + 10);

// select elements at even-indexed locations
int stencil[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::device_vector<int> d_stencil(stencil, stencil + 10);

// map all even indices into the first half of the range
// and odd indices to the last half of the range
int map[10]   = {0, 2, 4, 6, 8, 1, 3, 5, 7, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10, 7);
thrust::gather_if(thrust::device,
                  d_map.begin(), d_map.end(),
                  d_stencil.begin(),
                  d_values.begin(),
                  d_output.begin());
// d_output is now {0, 7, 4, 7, 8, 7, 3, 7, 7, 7}
Parameters
  • exec – The execution policy to use for parallelization.

  • map_first – Beginning of the range of gather locations.

  • map_last – End of the range of gather locations.

  • stencil – Beginning of the range of predicate values.

  • input_first – Beginning of the source range.

  • result – Beginning of the destination range.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator1 – must be a model of Input Iterator and InputIterator1's value_type must be convertible to RandomAccessIterator's difference_type.

  • InputIterator2 – must be a model of Input Iterator and InputIterator2's value_type must be convertible to bool.

  • RandomAccessIterator – must be a model of Random Access iterator and RandomAccessIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

Pre

The range [map_first, map_last) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The range [stencil, stencil + (map_last - map_first)) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The input data shall not overlap the range [result, result + (map_last - map_first)).

template<typename InputIterator1, typename InputIterator2, typename RandomAccessIterator, typename OutputIterator>
OutputIterator gather_if(InputIterator1 map_first, InputIterator1 map_last, InputIterator2 stencil, RandomAccessIterator input_first, OutputIterator result)#

gather_if conditionally copies elements from a source array into a destination range according to a map. For each input iterator i in the range [map_first, map_last), such that the value of *(stencil + (i - map_first)) is true, the value input_first[*i] is assigned to *(result + (i - map_first)). RandomAccessIterator must permit random access.

The following code snippet demonstrates how to use gather_if to gather selected values from an input range.

Remark

gather_if is the inverse of scatter_if.

#include <thrust/gather.h>
#include <thrust/device_vector.h>
...

int values[10] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
thrust::device_vector<int> d_values(values, values + 10);

// select elements at even-indexed locations
int stencil[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::device_vector<int> d_stencil(stencil, stencil + 10);

// map all even indices into the first half of the range
// and odd indices to the last half of the range
int map[10]   = {0, 2, 4, 6, 8, 1, 3, 5, 7, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10, 7);
thrust::gather_if(d_map.begin(), d_map.end(),
                  d_stencil.begin(),
                  d_values.begin(),
                  d_output.begin());
// d_output is now {0, 7, 4, 7, 8, 7, 3, 7, 7, 7}
Parameters
  • map_first – Beginning of the range of gather locations.

  • map_last – End of the range of gather locations.

  • stencil – Beginning of the range of predicate values.

  • input_first – Beginning of the source range.

  • result – Beginning of the destination range.

Template Parameters
  • InputIterator1 – must be a model of Input Iterator and InputIterator1's value_type must be convertible to RandomAccessIterator's difference_type.

  • InputIterator2 – must be a model of Input Iterator and InputIterator2's value_type must be convertible to bool.

  • RandomAccessIterator – must be a model of Random Access iterator and RandomAccessIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

Pre

The range [map_first, map_last) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The range [stencil, stencil + (map_last - map_first)) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The input data shall not overlap the range [result, result + (map_last - map_first)).

template<typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename RandomAccessIterator, typename OutputIterator, typename Predicate> __host__ __device__ OutputIterator gather_if (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator1 map_first, InputIterator1 map_last, InputIterator2 stencil, RandomAccessIterator input_first, OutputIterator result, Predicate pred)

gather_if conditionally copies elements from a source array into a destination range according to a map. For each input iterator i in the range [map_first, map_last) such that the value of pred(*(stencil + (i - map_first))) is true, the value input_first[*i] is assigned to *(result + (i - map_first)). RandomAccessIterator must permit random access.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use gather_if to gather selected values from an input range based on an arbitrary selection function using the thrust::device execution policy for parallelization:

Remark

gather_if is the inverse of scatter_if.

#include <thrust/gather.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>

struct is_even
{
  __host__ __device__
  bool operator()(const int x)
  {
    return (x % 2) == 0;
  }
};
...

int values[10] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
thrust::device_vector<int> d_values(values, values + 10);

// we will select an element when our stencil is even
int stencil[10] = {0, 3, 4, 1, 4, 1, 2, 7, 8, 9};
thrust::device_vector<int> d_stencil(stencil, stencil + 10);

// map all even indices into the first half of the range
// and odd indices to the last half of the range
int map[10]   = {0, 2, 4, 6, 8, 1, 3, 5, 7, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10, 7);
thrust::gather_if(thrust::device,
                  d_map.begin(), d_map.end(),
                  d_stencil.begin(),
                  d_values.begin(),
                  d_output.begin(),
                  is_even());
// d_output is now {0, 7, 4, 7, 8, 7, 3, 7, 7, 7}
Parameters
  • exec – The execution policy to use for parallelization.

  • map_first – Beginning of the range of gather locations.

  • map_last – End of the range of gather locations.

  • stencil – Beginning of the range of predicate values.

  • input_first – Beginning of the source range.

  • result – Beginning of the destination range.

  • pred – Predicate to apply to the stencil values.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator1 – must be a model of Input Iterator and InputIterator1's value_type must be convertible to RandomAccessIterator's difference_type.

  • InputIterator2 – must be a model of Input Iterator and InputIterator2's value_type must be convertible to Predicate's argument_type.

  • RandomAccessIterator – must be a model of Random Access iterator and RandomAccessIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

  • Predicate – must be a model of Predicate.

Pre

The range [map_first, map_last) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The range [stencil, stencil + (map_last - map_first)) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The input data shall not overlap the range [result, result + (map_last - map_first)).

template<typename InputIterator1, typename InputIterator2, typename RandomAccessIterator, typename OutputIterator, typename Predicate>
OutputIterator gather_if(InputIterator1 map_first, InputIterator1 map_last, InputIterator2 stencil, RandomAccessIterator input_first, OutputIterator result, Predicate pred)#

gather_if conditionally copies elements from a source array into a destination range according to a map. For each input iterator i in the range [map_first, map_last) such that the value of pred(*(stencil + (i - map_first))) is true, the value input_first[*i] is assigned to *(result + (i - map_first)). RandomAccessIterator must permit random access.

The following code snippet demonstrates how to use gather_if to gather selected values from an input range based on an arbitrary selection function.

Remark

gather_if is the inverse of scatter_if.

#include <thrust/gather.h>
#include <thrust/device_vector.h>

struct is_even
{
  __host__ __device__
  bool operator()(const int x)
  {
    return (x % 2) == 0;
  }
};
...

int values[10] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
thrust::device_vector<int> d_values(values, values + 10);

// we will select an element when our stencil is even
int stencil[10] = {0, 3, 4, 1, 4, 1, 2, 7, 8, 9};
thrust::device_vector<int> d_stencil(stencil, stencil + 10);

// map all even indices into the first half of the range
// and odd indices to the last half of the range
int map[10]   = {0, 2, 4, 6, 8, 1, 3, 5, 7, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10, 7);
thrust::gather_if(d_map.begin(), d_map.end(),
                  d_stencil.begin(),
                  d_values.begin(),
                  d_output.begin(),
                  is_even());
// d_output is now {0, 7, 4, 7, 8, 7, 3, 7, 7, 7}
Parameters
  • map_first – Beginning of the range of gather locations.

  • map_last – End of the range of gather locations.

  • stencil – Beginning of the range of predicate values.

  • input_first – Beginning of the source range.

  • result – Beginning of the destination range.

  • pred – Predicate to apply to the stencil values.

Template Parameters
  • InputIterator1 – must be a model of Input Iterator and InputIterator1's value_type must be convertible to RandomAccessIterator's difference_type.

  • InputIterator2 – must be a model of Input Iterator and InputIterator2's value_type must be convertible to Predicate's argument_type.

  • RandomAccessIterator – must be a model of Random Access iterator and RandomAccessIterator's value_type must be convertible to OutputIterator's value_type.

  • OutputIterator – must be a model of Output Iterator.

  • Predicate – must be a model of Predicate.

Pre

The range [map_first, map_last) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The range [stencil, stencil + (map_last - map_first)) shall not overlap the range [result, result + (map_last - map_first)).

Pre

The input data shall not overlap the range [result, result + (map_last - map_first)).

Functions

template<typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename RandomAccessIterator> __host__ __device__ void scatter (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator1 first, InputIterator1 last, InputIterator2 map, RandomAccessIterator result)

scatter copies elements from a source range into an output array according to a map. For each iterator i in the range [first, last), the value *i is assigned to output[*(map + (i - first))]. The output iterator must permit random access. If the same index appears more than once in the range [map, map + (last - first)), the result is undefined.

The algorithm’s execution is parallelized as determined by exec.

The following code snippet demonstrates how to use scatter to reorder a range using the thrust::device execution policy for parallelization:

#include <thrust/scatter.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>
...
// mark even indices with a 1; odd indices with a 0
int values[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::device_vector<int> d_values(values, values + 10);

// scatter all even indices into the first half of the
// range, and odd indices vice versa
int map[10]   = {0, 5, 1, 6, 2, 7, 3, 8, 4, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10);
thrust::scatter(thrust::device,
                d_values.begin(), d_values.end(),
                d_map.begin(), d_output.begin());
// d_output is now {1, 1, 1, 1, 1, 0, 0, 0, 0, 0}

Note

scatter is the inverse of thrust::gather.

Parameters
  • exec – The execution policy to use for parallelization.

  • first – Beginning of the sequence of values to scatter.

  • last – End of the sequence of values to scatter.

  • map – Beginning of the sequence of output indices.

  • result – Destination of the source elements.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator1 – must be a model of Input Iterator and InputIterator1's value_type must be convertible to RandomAccessIterator's value_type.

  • InputIterator2 – must be a model of Input Iterator and InputIterator2's value_type must be convertible to RandomAccessIterator's difference_type.

  • RandomAccessIterator – must be a model of Random Access iterator.

Pre

The iterator result + i shall not refer to any element referenced by any iterator j in the range [first,last) for all iterators i in the range [map,map + (last - first)).

Pre

The iterator result + i shall not refer to any element referenced by any iterator j in the range [map,map + (last - first)) for all iterators i in the range [map,map + (last - first)).

Pre

The expression result[*i] shall be valid for all iterators in the range [map,map + (last - first)).

template<typename InputIterator1, typename InputIterator2, typename RandomAccessIterator>
void scatter(InputIterator1 first, InputIterator1 last, InputIterator2 map, RandomAccessIterator result)#

scatter copies elements from a source range into an output array according to a map. For each iterator i in the range [first, last), the value *i is assigned to output[*(map + (i - first))]. The output iterator must permit random access. If the same index appears more than once in the range [map, map + (last - first)), the result is undefined.

The following code snippet demonstrates how to use scatter to reorder a range.

#include <thrust/scatter.h>
#include <thrust/device_vector.h>
...
// mark even indices with a 1; odd indices with a 0
int values[10] = {1, 0, 1, 0, 1, 0, 1, 0, 1, 0};
thrust::device_vector<int> d_values(values, values + 10);

// scatter all even indices into the first half of the
// range, and odd indices vice versa
int map[10]   = {0, 5, 1, 6, 2, 7, 3, 8, 4, 9};
thrust::device_vector<int> d_map(map, map + 10);

thrust::device_vector<int> d_output(10);
thrust::scatter(d_values.begin(), d_values.end(),
                d_map.begin(), d_output.begin());
// d_output is now {1, 1, 1, 1, 1, 0, 0, 0, 0, 0}

Note

scatter is the inverse of thrust::gather.

Parameters
  • first – Beginning of the sequence of values to scatter.

  • last – End of the sequence of values to scatter.

  • map – Beginning of the sequence of output indices.

  • result – Destination of the source elements.

Template Parameters
  • InputIterator1 – must be a model of Input Iterator and InputIterator1's value_type must be convertible to RandomAccessIterator's value_type.

  • InputIterator2 – must be a model of Input Iterator and InputIterator2's value_type must be convertible to RandomAccessIterator's difference_type.

  • RandomAccessIterator – must be a model of Random Access iterator.

Pre

The iterator result + i shall not refer to any element referenced by any iterator j in the range [first,last) for all iterators i in the range [map,map + (last - first)).

Pre

The iterator result + i shall not refer to any element referenced by any iterator j in the range [map,map + (last - first)) for all iterators i in the range [map,map + (last - first)).

Pre

The expression result[*i] shall be valid for all iterators in the range [map,map + (last - first)).

template<typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename InputIterator3, typename RandomAccessIterator> __host__ __device__ void scatter_if (const thrust::detail::execution_policy_base< DerivedPolicy > &exec, InputIterator1 first, InputIterator1 last, InputIterator2 map, InputIterator3 stencil, RandomAccessIterator output)

scatter_if conditionally copies elements from a source range into an output array according to a map. For each iterator i in the range [first, last) such that *(stencil + (i - first)) is true, the value *i is assigned to output[*(map + (i - first))]. The output iterator must permit random access. If the same index appears more than once in the range [map, map + (last - first)) the result is undefined.

The algorithm’s execution is parallelized as determined by exec.

#include <thrust/scatter.h>
#include <thrust/execution_policy.h>
...
int V[8] = {10, 20, 30, 40, 50, 60, 70, 80};
int M[8] = {0, 5, 1, 6, 2, 7, 3, 4};
int S[8] = {1, 0, 1, 0, 1, 0, 1, 0};
int D[8] = {0, 0, 0, 0, 0, 0, 0, 0};

thrust::scatter_if(thrust::host, V, V + 8, M, S, D);

// D contains [10, 30, 50, 70, 0, 0, 0, 0];

Note

scatter_if is the inverse of thrust::gather_if.

Parameters
  • exec – The execution policy to use for parallelization.

  • first – Beginning of the sequence of values to scatter.

  • last – End of the sequence of values to scatter.

  • map – Beginning of the sequence of output indices.

  • stencil – Beginning of the sequence of predicate values.

  • output – Beginning of the destination range.

Template Parameters
  • DerivedPolicy – The name of the derived execution policy.

  • InputIterator1 – must be a model of Input Iterator and InputIterator1's value_type must be convertible to RandomAccessIterator's value_type.

  • InputIterator2 – must be a model of Input Iterator and InputIterator2's value_type must be convertible to RandomAccessIterator's difference_type.

  • InputIterator3 – must be a model of Input Iterator and InputIterator3's value_type must be convertible to bool.

  • RandomAccessIterator – must be a model of Random Access iterator.

Pre

The iterator result + i shall not refer to any element referenced by any iterator j in the range [first,last) for all iterators i in the range [map,map + (last - first)).

Pre

The iterator result + i shall not refer to any element referenced by any iterator j in the range [map,map + (last - first)) for all iterators i in the range [map,map + (last - first)).

Pre

The iterator result + i shall not refer to any element referenced by any iterator j in the range [stencil,stencil + (last - first)) for all iterators i in the range [map,map + (last - first)).

Pre

The expression result[*i] shall be valid for all iterators i in the range [map,map + (last - first)) for which the following condition holds: *(stencil + i) != false.

template<typename InputIterator1, typename InputIterator2, typename InputIterator3, typename RandomAccessIterator>
void scatter_if(InputIterator1 first, InputIterator1 last, InputIterator2 map, InputIterator3 stencil, RandomAccessIterator output)#

scatter_if conditionally copies elements from a source range into an output array according to a map. For each iterator i in the range [first, last) such that *(stencil + (i - first)) is true, the value *i is assigned to output[*(map + (i - first))]. The output iterator must permit random access. If the same index appears more than once in the range [map, map + (last - first)) the result is undefined.

#include <thrust/scatter.h>
...
int V[8] = {10, 20, 30, 40, 50, 60, 70, 80};
int M[8] = {0, 5, 1, 6, 2, 7, 3, 4};
int S[8] = {1, 0, 1, 0, 1</