Template Class device_reference¶
Defined in File device_ptr.h
Inheritance Relationships¶
Base Type¶
public thrust::reference< T, thrust::device_ptr< T >, thrust::device_reference< T > >
Class Documentation¶
-
template<typename
T
>
classdevice_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 adevice_ptr
. Similarly, taking the address of adevice_reference
yields adevice_ptr
.device_reference
may often be used from host code in place of operations defined on its associatedvalue_type
. For example, whendevice_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 aniostream:
#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 bydevice_vector's
bracket operator. A more natural way to print the value of adevice_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 itsdevice_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 likeprintf
which have varargs parameters. Because varargs parameters must be Plain Old Data, adevice_reference
to a POD type requires a cast when passed toprintf:
#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
- See
Public Types
-
typedef super_t::value_type
value_type
¶ The type of the value referenced by this type of
device_reference
.
-
typedef super_t::pointer
pointer
¶ The type of the expression
&ref
, whereref
is adevice_reference
.
Public Functions
-
template<typename OtherT>__host__ __device__ thrust::device_reference::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 thisdevice_reference
is constructed, it shall refer to the same object asother
.The following code snippet demonstrates the semantics of this copy constructor.
- Parameters
other
: Adevice_reference
to copy from.
#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>
fromdevice_reference<T>
.
-
__host__ __device__ thrust::device_reference::device_reference(const pointer & ptr)
This copy constructor initializes this
device_reference
to refer to an object pointed to by the givendevice_ptr
. After thisdevice_reference
is constructed, it shall refer to the object pointed to byptr
.The following code snippet demonstrates the semantic of this copy constructor.
- Parameters
ptr
: Adevice_ptr
to copy from.
#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);
-
template<typename OtherT>__host__ __device__ device_reference& thrust::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 thisdevice_reference
.- Return
*this
- Parameters
other
: Thedevice_reference
to assign from.
-
__host__ __device__ device_reference& thrust::device_reference::operator=(const value_type & x)
Assignment operator assigns the value of the given value to the value referenced by this
device_reference
.- Return
*this
- Parameters
x
: The value to assign from.