/*
 *  Copyright 2008-2010 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */


/*! \file device_reference.h
 *  \brief Defines the interface to a reference to
 *         a variable which resides on a CUDA device.
 */

#pragma once

#include <thrust/detail/config.h>
#include <thrust/device_ptr.h>
#include <thrust/detail/type_traits.h>
#include <thrust/iterator/iterator_traits.h>

namespace thrust
{

/*! \addtogroup memory_management_classes Memory Management Classes
 *  \ingroup memory_management
 *  \{
 */

/*! \p device_reference acts as a reference-like object to an object stored in device memory.
 *  \p device_reference is not intended to be used directly; rather, this type
 *  is the result of deferencing a \p device_ptr. Similarly, taking the address of
 *  a \p device_reference yields a \p device_ptr.
 *  
 *  \p device_reference may often be used from host code in place of operations defined on
 *  its associated \c value_type. For example, when \p device_reference refers to an
 *  arithmetic type, arithmetic operations on it are legal:
 *
 *  \code
 *  #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;
 *  }
 *  \endcode
 *
 *  Similarly, we can print the value of \c ref_to_thirteen in the above code by using an
 *  \c iostream:
 *
 *  \code
 *  #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;
 *  }
 *  \endcode
 *
 *  Of course, we needn't explicitly create a \p device_reference in the previous
 *  example, because one is returned by \p device_vector's bracket operator. A more natural
 *  way to print the value of a \p device_vector element might be:
 *
 *  \code
 *  #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;
 *  }
 *  \endcode
 *
 *  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 \p 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 \p device_reference.
 *
 *  The following code, which generates a compiler error, illustrates:
 *
 *  \code
 *  #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;
 *  }
 *  \endcode
 *
 *  Instead, a host space copy must be created to access \c foo's \c x member:
 *
 *  \code
 *  #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;
 *  }
 *  \endcode
 *  
 *  Another common case where a \p device_reference cannot directly be used in place of
 *  its referent object occurs when passing them as parameters to functions like \c printf
 *  which have varargs parameters. Because varargs parameters must be Plain Old Data, a
 *  \p device_reference to a POD type requires a cast when passed to \c printf:
 *
 *  \code
 *  #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;
 *  }
 *  \endcode
 *
 *  \see device_ptr
 *  \see device_vector
 */
template<typename T>
  class device_reference
{
  public:
    typedef device_ptr<T> pointer;
    typedef typename detail::remove_const<T>::type value_type;

    /*! This copy constructor accepts a const reference to another
     *  \p device_reference. After this \p device_reference is constructed,
     *  it shall refer to the same object as \p ref.
     *  
     *  \param ref A \p device_reference to copy from.
     *
     *  The following code snippet demonstrates the semantics of this
     *  copy constructor.
     *
     *  \code
     *  #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(ref1 == 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);
     *  \endcode
     *
     *  \note This constructor is templated to initialization of const
     *  device_references to const T from device_references to T.
     */
    template<typename OtherT>
    device_reference(const device_reference<OtherT> &ref

// XXX msvc screws this up
#ifndef _MSC_VER
                     , typename
                     detail::enable_if<
                       detail::is_convertible<
                         typename device_reference<OtherT>::pointer,
                         pointer
                       >::value
                     >::type *dummy = 0
#endif // __MSC_VER
                    );

    /*! This copy constructor initializes this \p device_reference
     *  to refer to an object pointed to by the given \p device_ptr. After
     *  this \p device_reference is constructed, it shall refer to the
     *  object pointed to by \p ptr.
     *
     *  \param ptr A \p device_ptr to copy from.
     *
     *  The following code snippet demonstrates the semantic of this
     *  copy constructor.
     *
     *  \code
     *  #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);
     *  \endcode
     */
    explicit device_reference(const pointer &ptr);

    /*! Address-of operator returns a \p device_ptr pointing to the object
     *  referenced by this \p device_reference. It does not return the
     *  address of this \p device_reference.
     *
     *  \return A \p device_ptr pointing to the object this
     *  \p device_reference references.
     */
    pointer operator&(void) const;

    /*! Assignment operator copies the value of the given object to the
     *  object referenced by this \p device_reference.
     *  
     *  \param v The value to copy from.
     *  \return This \p device_reference.
     */
    device_reference &operator=(const T &v);

    /*! Assignment operator copies the value of the object referenced by
     *  the given \p device_reference to the object referenced by this
     *  \p device_reference.
     *
     *  \param ref The \p device_reference to copy from.
     *  \return This \p device_reference.
     */
    device_reference &operator=(const device_reference &ref);

    /*! Assignment operator copies the value of the object referenced by
     *  the given \p device_reference to the object referenced by this
     *  \p device_reference.
     *
     *  This operator is distinguished from the previous assignment operator
     *  by being templated on the type of the other device_reference to copy from,
     *  allowing copy from a convertible type.
     *
     *  \param ref The \p device_reference to copy from.
     *  \return This \p device_reference.
     */
    template<typename OtherT>
    device_reference &operator=(const device_reference<OtherT> &ref);

    /*! Prefix increment operator increments the object referenced by this
     *  \p device_reference.
     *
     *  \return <tt>*this</tt>
     *  
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's prefix increment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The increment executes as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator++(void);

    /*! Postfix increment operator copies the object referenced by this
     *  \p device_reference, increments the object referenced by this
     *  \p device_reference, and returns the copy.
     *
     *  \return A copy of the object referenced by this \p device_reference
     *          before being incremented.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's postfix increment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The increment executes as if it were executed on the host.
     *  This may change in a later version.
     */
    value_type operator++(int);

    /*! Addition assignment operator add-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the add-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's addition assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The add-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator+=(const T &rhs);

    /*! Prefix decrement operator decrements the object referenced by this
     *  \p device_reference.
     *
     *  \return <tt>*this</tt>
     *  
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's prefix decrement operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The decrement executes as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator--(void);

    /*! Postfix decrement operator copies the object referenced by this
     *  \p device_reference, decrements the object referenced by this
     *  \p device_reference, and returns the copy.
     *
     *  \return A copy of the object referenced by this \p device_reference
     *          before being decremented.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's postfix decrement operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The decrement executes as if it were executed on the host.
     *  This may change in a later version.
     */
    value_type operator--(int);

    /*! Subtraction assignment operator subtract-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the subtraction-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's addition assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The subtract-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator-=(const T &rhs);

    /*! Multiplication assignment operator multiply-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the multiply-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's multiply assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The multiply-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator*=(const T &rhs);

    /*! Division assignment operator divide-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the divide-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's divide assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The divide-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator/=(const T &rhs);

    /*! Modulation assignment operator modulus-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the divide-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's divide assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The modulus-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator%=(const T &rhs);

    /*! Bitwise left shift assignment operator left shift-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the left shift-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's left shift assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The left shift-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator<<=(const T &rhs);

    /*! Bitwise right shift assignment operator right shift-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the right shift-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's right shift assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The right shift-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator>>=(const T &rhs);

    /*! Bitwise AND assignment operator AND-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the AND-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's AND assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The AND-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator&=(const T &rhs);

    /*! Bitwise OR assignment operator OR-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the OR-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's OR assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The OR-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator|=(const T &rhs);

    /*! Bitwise XOR assignment operator XOR-assigns the object referenced by this
     *  \p device_reference and returns this \p device_reference.
     *
     *  \param rhs The right hand side of the XOR-assignment.
     *  \return <tt>*this</tt>.
     *
     *  The following code snippet demonstrates the semantics of
     *  \p device_reference's XOR assignment operator.
     *
     *  \code
     *  #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);
     *  \endcode
     *
     *  \note The XOR-assignment executes as as if it were executed on the host.
     *  This may change in a later version.
     */
    device_reference &operator^=(const T &rhs);

    /*! Conversion operator converts this \p device_reference to T
     *  by returning a copy of the object referenced by this
     *  \p device_reference.
     *
     *  \return A copy of the object referenced by this \p device_reference.
     */
    operator value_type (void) const;

  private:
    // allow access to mPtr for other device_references
    template <typename OtherT> friend class device_reference;

    // the pointer
    const pointer mPtr;

    template<typename Pointer>
    inline void assign_from(Pointer src, thrust::detail::true_type spaces_are_interoperable);

    template<typename Pointer>
    inline void assign_from(Pointer src, thrust::detail::false_type spaces_are_not_interoperable);

    inline value_type convert(thrust::detail::true_type spaces_are_interoperable) const;

    inline value_type convert(thrust::detail::false_type spaces_are_not_interoperable) const;
}; // end device_reference

/*! \}
 */

} // end thrust

#include <thrust/detail/device_reference.inl>