thrust::device_reference

Defined in thrust/device_reference.h

template<typename T> class device_reference : public thrust::reference<T, thrust::device_ptr<T>, thrust::device_reference<T>>

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 dereferencing 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()
{
  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()
{
  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()
{
  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()
{
  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()
{
  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()
{
  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_vector

Public Types

using value_type = typename super_t::value_type: The type of the value referenced by this type of device_reference.

using pointer = typename super_t::pointer: The type of the expression &ref, where ref is a device_reference.

Public Functions

device_reference(const device_reference &other) = default

template<typename OtherT> inline device_reference(const device_reference<OtherT> &other, thrust::detail::enable_if_convertible_t<typename device_reference<OtherT>::pointer, pointer>* = 0)

This copy constructor accepts a const reference to another device_reference. After this device_reference is constructed, it shall refer to the same object as other.

The following code snippet demonstrates the semantics of this copy constructor.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_reference<int> ref = v[0];

// ref equals the object at v[0]
assert(ref == v[0]);

// the address of ref equals the address of v[0]
assert(&ref == &v[0]);

// modifying v[0] modifies ref
v[0] = 13;
assert(ref == 13);

Note

This constructor is templated primarily to allow initialization of device_reference<const T> from device_reference<T>.

Parameters: other – A device_reference to copy from.

inline explicit device_reference(const pointer &ptr)

This copy constructor initializes this device_reference to refer to an object pointed to by the given device_ptr. After this device_reference is constructed, it shall refer to the object pointed to by ptr.

The following code snippet demonstrates the semantic of this copy constructor.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals the object pointed to by ptr
assert(ref == *ptr);

// the address of ref equals ptr
assert(&ref == ptr);

// modifying *ptr modifies ref
*ptr = 13;
assert(ref == 13);

Parameters: ptr – A device_ptr to copy from.

inline const device_reference &operator=(const device_reference &other) const

template<typename OtherT> inline const device_reference &operator=(const device_reference<OtherT> &other) const

This assignment operator assigns the value of the object referenced by the given device_reference to the object referenced by this device_reference.

Parameters: other – The device_reference to assign from.
Returns: *this

inline const device_reference &operator=(const value_type &x) const

Assignment operator assigns the value of the given value to the value referenced by this device_reference.

Parameters: x – The value to assign from.
Returns: *this

pointer operator&(void) const

Address-of operator returns a device_ptr pointing to the object referenced by this device_reference. It does not return the address of this device_reference.

Returns: A device_ptr pointing to the object this device_reference references.

operator value_type(void) const

Conversion operator converts this device_reference to T by returning a copy of the object referenced by this device_reference.

Returns: A copy of the object referenced by this device_reference.

void swap(device_reference other): swaps the value this device_reference references with another. other The other device_reference with which to swap.

device_reference &operator++(void)

Prefix increment operator increments the object referenced by this device_reference.

The following code snippet demonstrates the semantics of device_reference's prefix increment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

// increment ref
++ref;

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

Note

The increment executes as if it were executed on the host. This may change in a later version.

Returns: *this

value_type operator++(int)

Postfix increment operator copies the object referenced by this device_reference, increments the object referenced by this device_reference, and returns the copy.

The following code snippet demonstrates the semantics of device_reference's postfix increment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

// increment ref
int x = ref++;

// x equals 0
assert(x == 0)

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

Note

The increment executes as if it were executed on the host. This may change in a later version.

Returns: A copy of the object referenced by this device_reference before being incremented.

device_reference &operator+=(const T &rhs)

Addition assignment operator add-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's addition assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

// add-assign ref
ref += 5;

// ref equals 5
assert(ref == 5);

// the object pointed to by ptr equals 5
assert(*ptr == 5);

// v[0] equals 5
assert(v[0] == 5);

Note

The add-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the add-assignment.
Returns: *this.

device_reference &operator--(void)

Prefix decrement operator decrements the object referenced by this device_reference.

The following code snippet demonstrates the semantics of device_reference's prefix decrement operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

// decrement ref
--ref;

// ref equals -1
assert(ref == -1);

// the object pointed to by ptr equals -1
assert(*ptr == -1);

// v[0] equals -1
assert(v[0] == -1);

Note

The decrement executes as if it were executed on the host. This may change in a later version.

Returns: *this

value_type operator--(int)

Postfix decrement operator copies the object referenced by this device_reference, decrements the object referenced by this device_reference, and returns the copy.

The following code snippet demonstrates the semantics of device_reference's postfix decrement operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

// decrement ref
int x = ref--;

// x equals 0
assert(x == 0)

// ref equals -1
assert(ref == -1);

// the object pointed to by ptr equals -1
assert(*ptr == -1);

// v[0] equals -1
assert(v[0] == -1);

Note

The decrement executes as if it were executed on the host. This may change in a later version.

Returns: A copy of the object referenced by this device_reference before being decremented.

device_reference &operator-=(const T &rhs)

Subtraction assignment operator subtract-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's addition assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

// subtract-assign ref
ref -= 5;

// ref equals -5
assert(ref == -5);

// the object pointed to by ptr equals -5
assert(*ptr == -5);

// v[0] equals -5
assert(v[0] == -5);

Note

The subtract-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the subtraction-assignment.
Returns: *this.

device_reference &operator*=(const T &rhs)

Multiplication assignment operator multiply-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's multiply assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,1);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

// multiply-assign ref
ref *= 5;

// ref equals 5
assert(ref == 5);

// the object pointed to by ptr equals 5
assert(*ptr == 5);

// v[0] equals 5
assert(v[0] == 5);

Note

The multiply-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the multiply-assignment.
Returns: *this.

device_reference &operator/=(const T &rhs)

Division assignment operator divide-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's divide assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,5);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 5
assert(ref == 5);

// the object pointed to by ptr equals 5
assert(*ptr == 5);

// v[0] equals 5
assert(v[0] == 5);

// divide-assign ref
ref /= 5;

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

Note

The divide-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the divide-assignment.
Returns: *this.

device_reference &operator%=(const T &rhs)

Modulation assignment operator modulus-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's divide assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,5);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 5
assert(ref == 5);

// the object pointed to by ptr equals 5
assert(*ptr == 5);

// v[0] equals 5
assert(v[0] == 5);

// modulus-assign ref
ref %= 5;

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

Note

The modulus-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the divide-assignment.
Returns: *this.

device_reference &operator<<=(const T &rhs)

Bitwise left shift assignment operator left shift-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's left shift assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,1);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

// left shift-assign ref
ref <<= 1;

// ref equals 2
assert(ref == 2);

// the object pointed to by ptr equals 2
assert(*ptr == 2);

// v[0] equals 2
assert(v[0] == 2);

Note

The left shift-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the left shift-assignment.
Returns: *this.

device_reference &operator>>=(const T &rhs)

Bitwise right shift assignment operator right shift-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's right shift assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,2);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 2
assert(ref == 2);

// the object pointed to by ptr equals 2
assert(*ptr == 2);

// v[0] equals 2
assert(v[0] == 2);

// right shift-assign ref
ref >>= 1;

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

Note

The right shift-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the right shift-assignment.
Returns: *this.

device_reference &operator&=(const T &rhs)

Bitwise AND assignment operator AND-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's AND assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,1);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

// right AND-assign ref
ref &= 0;

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

Note

The AND-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the AND-assignment.
Returns: *this.

device_reference &operator|=(const T &rhs)

Bitwise OR assignment operator OR-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's OR assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,0);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

// right OR-assign ref
ref |= 1;

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

Note

The OR-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the OR-assignment.
Returns: *this.

device_reference &operator^=(const T &rhs)

Bitwise XOR assignment operator XOR-assigns the object referenced by this device_reference and returns this device_reference.

The following code snippet demonstrates the semantics of device_reference's XOR assignment operator.

#include <thrust/device_vector.h>
#include <assert.h>
...
thrust::device_vector<int> v(1,1);
thrust::device_ptr<int> ptr = &v[0];
thrust::device_reference<int> ref(ptr);

// ref equals 1
assert(ref == 1);

// the object pointed to by ptr equals 1
assert(*ptr == 1);

// v[0] equals 1
assert(v[0] == 1);

// right XOR-assign ref
ref ^= 1;

// ref equals 0
assert(ref == 0);

// the object pointed to by ptr equals 0
assert(*ptr == 0);

// v[0] equals 0
assert(v[0] == 0);

Note

The XOR-assignment executes as as if it were executed on the host. This may change in a later version.

Parameters: rhs – The right hand side of the XOR-assignment.
Returns: *this.