thrust
Public Types | Public Member Functions | List of all members
thrust::device_reference< T > Class Template Reference

#include <thrust/device_reference.h>

Inherits reference< T, thrust::device_ptr< T >, thrust::device_reference< T > >.

Public Types

typedef super_t::value_type value_type
 
typedef super_t::pointer pointer
 

Public Member Functions

template<typename OtherT >
__host__ __device__ device_reference (const device_reference< OtherT > &other, typename thrust::detail::enable_if_convertible< typename device_reference< OtherT >::pointer, pointer >::type *=0)
 
__host__ __device__ device_reference (const pointer &ptr)
 
template<typename OtherT >
__host__ __device__ device_referenceoperator= (const device_reference< OtherT > &other)
 
__host__ __device__ device_referenceoperator= (const value_type &x)
 

Detailed Description

template<typename T>
class 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:

int main(void)
{
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 <iostream>
int main(void)
{
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 <iostream>
int main(void)
{
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:

struct foo
{
int x;
};
int main(void)
{
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:

struct foo
{
int x;
};
int main(void)
{
// 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>
int main(void)
{
// vec[0] must be cast to int when passing to printf
printf("%d\n", (int) vec[0]);
return 0;
}
See also
device_ptr
device_vector

The documentation for this class was generated from the following files: