Class thrust::device_reference
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;
}
Inherits From: thrust::reference< T, thrust::device_ptr< T >, thrust::device_reference< T > >
See:
#include <thrust/device_reference.h>
template <typename T> class thrust::device_reference { public: typedef see below value_type;
typedef see below pointer;
template <typename OtherT> _CCCL_HOST_DEVICE device_reference(const device_reference< OtherT > & other, typename thrust::detail::enable_if_convertible< typename device_reference< OtherT >::pointer, pointer >::type * = 0);
explicit _CCCL_HOST_DEVICE device_reference(const pointer & ptr);
template <typename OtherT> _CCCL_HOST_DEVICE device_reference & operator=(const device_reference< OtherT > & other);
_CCCL_HOST_DEVICE device_reference & operator=(const value_type & x); };
Member Types
Typedef thrust::device_reference::value_type
typedef super_t::value_typevalue_type;
The type of the value referenced by this type of device_reference
.
Typedef thrust::device_reference::pointer
typedef super_t::pointerpointer;
The type of the expression &ref
, where ref
is a device_reference
.
Member Functions
Function thrust::device_reference::device_reference
template <typename OtherT> _CCCL_HOST_DEVICE device_reference(const device_reference< OtherT > & other, typename thrust::detail::enable_if_convertible< typename device_reference< OtherT >::pointer, pointer >::type * = 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>
.
Function Parameters: other
: A device_reference
to copy from.
Function thrust::device_reference::device_reference
explicit _CCCL_HOST_DEVICE 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);
Function Parameters: ptr
: A device_ptr
to copy from.
Function thrust::device_reference::operator=
template <typename OtherT> _CCCL_HOST_DEVICE 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
.
Function Parameters: other
: The device_reference
to assign from.
Returns: *this
Function thrust::device_reference::operator=
_CCCL_HOST_DEVICE device_reference & operator=(const value_type & x);
Assignment operator assigns the value of the given value to the value referenced by this device_reference
.
Function Parameters: x
: The value to assign from.
Returns: *this