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

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.

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

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 device_reference &operator=(const value_type &x)

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