utiltools/X86MAC64/cuda/include/thrust/device_reference.h

   1 /*
   2  *  Copyright 2008-2012 NVIDIA Corporation
   3  *
   4  *  Licensed under the Apache License, Version 2.0 (the "License");
   5  *  you may not use this file except in compliance with the License.
   6  *  You may obtain a copy of the License at
   7  *
   8  *      http://www.apache.org/licenses/LICENSE-2.0
   9  *
  10  *  Unless required by applicable law or agreed to in writing, software
  11  *  distributed under the License is distributed on an "AS IS" BASIS,
  12  *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  13  *  See the License for the specific language governing permissions and
  14  *  limitations under the License.
  15  */
  16
  17
  18 /*! \file device_reference.h
  19  *  \brief A reference to a variable which resides in the "device" system's memory space
  20  */
  21
  22 #pragma once
  23
  24 #include <thrust/detail/config.h>
  25 #include <thrust/device_ptr.h>
  26 #include <thrust/detail/type_traits.h>
  27 #include <thrust/detail/reference.h>
  28
  29 namespace thrust
  30 {
  31
  32 /*! \addtogroup memory_management_classes Memory Management Classes
  33  *  \ingroup memory_management
  34  *  \{
  35  */
  36
  37 /*! \p device_reference acts as a reference-like object to an object stored in device memory.
  38  *  \p device_reference is not intended to be used directly; rather, this type
  39  *  is the result of deferencing a \p device_ptr. Similarly, taking the address of
  40  *  a \p device_reference yields a \p device_ptr.
  41  *
  42  *  \p device_reference may often be used from host code in place of operations defined on
  43  *  its associated \c value_type. For example, when \p device_reference refers to an
  44  *  arithmetic type, arithmetic operations on it are legal:
  45  *
  46  *  \code
  47  *  #include <thrust/device_vector.h>
  48  *
  49  *  int main(void)
  50  *  {
  51  *    thrust::device_vector<int> vec(1, 13);
  52  *
  53  *    thrust::device_reference<int> ref_to_thirteen = vec[0];
  54  *
  55  *    int x = ref_to_thirteen + 1;
  56  *
  57  *    // x is 14
  58  *
  59  *    return 0;
  60  *  }
  61  *  \endcode
  62  *
  63  *  Similarly, we can print the value of \c ref_to_thirteen in the above code by using an
  64  *  \c iostream:
  65  *
  66  *  \code
  67  *  #include <thrust/device_vector.h>
  68  *  #include <iostream>
  69  *
  70  *  int main(void)
  71  *  {
  72  *    thrust::device_vector<int> vec(1, 13);
  73  *
  74  *    thrust::device_reference<int> ref_to_thirteen = vec[0];
  75  *
  76  *    std::cout << ref_to_thirteen << std::endl;
  77  *
  78  *    // 13 is printed
  79  *
  80  *    return 0;
  81  *  }
  82  *  \endcode
  83  *
  84  *  Of course, we needn't explicitly create a \p device_reference in the previous
  85  *  example, because one is returned by \p device_vector's bracket operator. A more natural
  86  *  way to print the value of a \p device_vector element might be:
  87  *
  88  *  \code
  89  *  #include <thrust/device_vector.h>
  90  *  #include <iostream>
  91  *
  92  *  int main(void)
  93  *  {
  94  *    thrust::device_vector<int> vec(1, 13);
  95  *
  96  *    std::cout << vec[0] << std::endl;
  97  *
  98  *    // 13 is printed
  99  *
 100  *    return 0;
 101  *  }
 102  *  \endcode
 103  *
 104  *  These kinds of operations should be used sparingly in performance-critical code, because
 105  *  they imply a potentially expensive copy between host and device space.
 106  *
 107  *  Some operations which are possible with regular objects are impossible with their
 108  *  corresponding \p device_reference objects due to the requirements of the C++ language. For
 109  *  example, because the member access operator cannot be overloaded, member variables and functions
 110  *  of a referent object cannot be directly accessed through its \p device_reference.
 111  *
 112  *  The following code, which generates a compiler error, illustrates:
 113  *
 114  *  \code
 115  *  #include <thrust/device_vector.h>
 116  *
 117  *  struct foo
 118  *  {
 119  *    int x;
 120  *  };
 121  *
 122  *  int main(void)
 123  *  {
 124  *    thrust::device_vector<foo> foo_vec(1);
 125  *
 126  *    thrust::device_reference<foo> foo_ref = foo_vec[0];
 127  *
 128  *    foo_ref.x = 13; // ERROR: x cannot be accessed through foo_ref
 129  *
 130  *    return 0;
 131  *  }
 132  *  \endcode
 133  *
 134  *  Instead, a host space copy must be created to access \c foo's \c x member:
 135  *
 136  *  \code
 137  *  #include <thrust/device_vector.h>
 138  *
 139  *  struct foo
 140  *  {
 141  *    int x;
 142  *  };
 143  *
 144  *  int main(void)
 145  *  {
 146  *    thrust::device_vector<foo> foo_vec(1);
 147  *
 148  *    // create a local host-side foo object
 149  *    foo host_foo;
 150  *    host_foo.x = 13;
 151  *
 152  *    thrust::device_reference<foo> foo_ref = foo_vec[0];
 153  *
 154  *    foo_ref = host_foo;
 155  *
 156  *    // foo_ref's x member is 13
 157  *
 158  *    return 0;
 159  *  }
 160  *  \endcode
 161  *
 162  *  Another common case where a \p device_reference cannot directly be used in place of
 163  *  its referent object occurs when passing them as parameters to functions like \c printf
 164  *  which have varargs parameters. Because varargs parameters must be Plain Old Data, a
 165  *  \p device_reference to a POD type requires a cast when passed to \c printf:
 166  *
 167  *  \code
 168  *  #include <stdio.h>
 169  *  #include <thrust/device_vector.h>
 170  *
 171  *  int main(void)
 172  *  {
 173  *    thrust::device_vector<int> vec(1,13);
 174  *
 175  *    // vec[0] must be cast to int when passing to printf
 176  *    printf("%d\n", (int) vec[0]);
 177  *
 178  *    return 0;
 179  *  }
 180  *  \endcode
 181  *
 182  *  \see device_ptr
 183  *  \see device_vector
 184  */
 185 template<typename T>
 186   class device_reference
 187     : public thrust::reference<
 188                T,
 189                thrust::device_ptr<T>,
 190                thrust::device_reference<T>
 191              >
 192 {
 193   private:
 194     typedef thrust::reference<
 195       T,
 196       thrust::device_ptr<T>,
 197       thrust::device_reference<T>
 198     > super_t;
 199
 200   public:
 201     /*! The type of the value referenced by this type of \p device_reference.
 202      */
 203     typedef typename super_t::value_type value_type;
 204
 205     /*! The type of the expression <tt>&ref</tt>, where <tt>ref</tt> is a \p device_reference.
 206      */
 207     typedef typename super_t::pointer    pointer;
 208
 209     /*! This copy constructor accepts a const reference to another
 210      *  \p device_reference. After this \p device_reference is constructed,
 211      *  it shall refer to the same object as \p other.
 212      *
 213      *  \param other A \p device_reference to copy from.
 214      *
 215      *  The following code snippet demonstrates the semantics of this
 216      *  copy constructor.
 217      *
 218      *  \code
 219      *  #include <thrust/device_vector.h>
 220      *  #include <assert.h>
 221      *  ...
 222      *  thrust::device_vector<int> v(1,0);
 223      *  thrust::device_reference<int> ref = v[0];
 224      *
 225      *  // ref equals the object at v[0]
 226      *  assert(ref == v[0]);
 227      *
 228      *  // the address of ref equals the address of v[0]
 229      *  assert(&ref == &v[0]);
 230      *
 231      *  // modifying v[0] modifies ref
 232      *  v[0] = 13;
 233      *  assert(ref == 13);
 234      *  \endcode
 235      *
 236      *  \note This constructor is templated primarily to allow initialization of
 237      *  <tt>device_reference<const T></tt> from <tt>device_reference<T></tt>.
 238      */
 239     template<typename OtherT>
 240     __host__ __device__
 241     device_reference(const device_reference<OtherT> &other,
 242                      typename thrust::detail::enable_if_convertible<
 243                        typename device_reference<OtherT>::pointer,
 244                        pointer
 245                      >::type * = 0)
 246       : super_t(other)
 247     {}
 248
 249     /*! This copy constructor initializes this \p device_reference
 250      *  to refer to an object pointed to by the given \p device_ptr. After
 251      *  this \p device_reference is constructed, it shall refer to the
 252      *  object pointed to by \p ptr.
 253      *
 254      *  \param ptr A \p device_ptr to copy from.
 255      *
 256      *  The following code snippet demonstrates the semantic of this
 257      *  copy constructor.
 258      *
 259      *  \code
 260      *  #include <thrust/device_vector.h>
 261      *  #include <assert.h>
 262      *  ...
 263      *  thrust::device_vector<int> v(1,0);
 264      *  thrust::device_ptr<int> ptr = &v[0];
 265      *  thrust::device_reference<int> ref(ptr);
 266      *
 267      *  // ref equals the object pointed to by ptr
 268      *  assert(ref == *ptr);
 269      *
 270      *  // the address of ref equals ptr
 271      *  assert(&ref == ptr);
 272      *
 273      *  // modifying *ptr modifies ref
 274      *  *ptr = 13;
 275      *  assert(ref == 13);
 276      *  \endcode
 277      */
 278     __host__ __device__
 279     explicit device_reference(const pointer &ptr)
 280       : super_t(ptr)
 281     {}
 282
 283     /*! This assignment operator assigns the value of the object referenced by
 284      *  the given \p device_reference to the object referenced by this
 285      *  \p device_reference.
 286      *
 287      *  \param other The \p device_reference to assign from.
 288      *  \return <tt>*this</tt>
 289      */
 290     template<typename OtherT>
 291     __host__ __device__
 292     device_reference &operator=(const device_reference<OtherT> &other);
 293
 294     /*! Assignment operator assigns the value of the given value to the
 295      *  value referenced by this \p device_reference.
 296      *
 297      *  \param x The value to assign from.
 298      *  \return <tt>*this</tt>
 299      */
 300     __host__ __device__
 301     device_reference &operator=(const value_type &x);
 302
 303 // declare these members for the purpose of Doxygenating them
 304 // they actually exist in a derived-from class
 305 #if 0
 306     /*! Address-of operator returns a \p device_ptr pointing to the object
 307      *  referenced by this \p device_reference. It does not return the
 308      *  address of this \p device_reference.
 309      *
 310      *  \return A \p device_ptr pointing to the object this
 311      *  \p device_reference references.
 312      */
 313     __host__ __device__
 314     pointer operator&(void) const;
 315
 316     /*! Conversion operator converts this \p device_reference to T
 317      *  by returning a copy of the object referenced by this
 318      *  \p device_reference.
 319      *
 320      *  \return A copy of the object referenced by this \p device_reference.
 321      */
 322     __host__ __device__
 323     operator value_type (void) const;
 324
 325     /*! swaps the value this \p device_reference references with another.
 326      *  \p other The other \p device_reference with which to swap.
 327      */
 328     __host__ __device__
 329     void swap(device_reference &other);
 330
 331     /*! Prefix increment operator increments the object referenced by this
 332      *  \p device_reference.
 333      *
 334      *  \return <tt>*this</tt>
 335      *
 336      *  The following code snippet demonstrates the semantics of
 337      *  \p device_reference's prefix increment operator.
 338      *
 339      *  \code
 340      *  #include <thrust/device_vector.h>
 341      *  #include <assert.h>
 342      *  ...
 343      *  thrust::device_vector<int> v(1,0);
 344      *  thrust::device_ptr<int> ptr = &v[0];
 345      *  thrust::device_reference<int> ref(ptr);
 346      *
 347      *  // ref equals 0
 348      *  assert(ref == 0);
 349      *
 350      *  // the object pointed to by ptr equals 1
 351      *  assert(*ptr == 1);
 352      *
 353      *  // v[0] equals 1
 354      *  assert(v[0] == 1);
 355      *
 356      *  // increment ref
 357      *  ++ref;
 358      *
 359      *  // ref equals 1
 360      *  assert(ref == 1);
 361      *
 362      *  // the object pointed to by ptr equals 1
 363      *  assert(*ptr == 1);
 364      *
 365      *  // v[0] equals 1
 366      *  assert(v[0] == 1);
 367      *  \endcode
 368      *
 369      *  \note The increment executes as if it were executed on the host.
 370      *  This may change in a later version.
 371      */
 372     device_reference &operator++(void);
 373
 374     /*! Postfix increment operator copies the object referenced by this
 375      *  \p device_reference, increments the object referenced by this
 376      *  \p device_reference, and returns the copy.
 377      *
 378      *  \return A copy of the object referenced by this \p device_reference
 379      *          before being incremented.
 380      *
 381      *  The following code snippet demonstrates the semantics of
 382      *  \p device_reference's postfix increment operator.
 383      *
 384      *  \code
 385      *  #include <thrust/device_vector.h>
 386      *  #include <assert.h>
 387      *  ...
 388      *  thrust::device_vector<int> v(1,0);
 389      *  thrust::device_ptr<int> ptr = &v[0];
 390      *  thrust::device_reference<int> ref(ptr);
 391      *
 392      *  // ref equals 0
 393      *  assert(ref == 0);
 394      *
 395      *  // the object pointed to by ptr equals 0
 396      *  assert(*ptr == 0);
 397      *
 398      *  // v[0] equals 0
 399      *  assert(v[0] == 0);
 400      *
 401      *  // increment ref
 402      *  int x = ref++;
 403      *
 404      *  // x equals 0
 405      *  assert(x == 0)
 406      *
 407      *  // ref equals 1
 408      *  assert(ref == 1);
 409      *
 410      *  // the object pointed to by ptr equals 1
 411      *  assert(*ptr == 1);
 412      *
 413      *  // v[0] equals 1
 414      *  assert(v[0] == 1);
 415      *  \endcode
 416      *
 417      *  \note The increment executes as if it were executed on the host.
 418      *  This may change in a later version.
 419      */
 420     value_type operator++(int);
 421
 422     /*! Addition assignment operator add-assigns the object referenced by this
 423      *  \p device_reference and returns this \p device_reference.
 424      *
 425      *  \param rhs The right hand side of the add-assignment.
 426      *  \return <tt>*this</tt>.
 427      *
 428      *  The following code snippet demonstrates the semantics of
 429      *  \p device_reference's addition assignment operator.
 430      *
 431      *  \code
 432      *  #include <thrust/device_vector.h>
 433      *  #include <assert.h>
 434      *  ...
 435      *  thrust::device_vector<int> v(1,0);
 436      *  thrust::device_ptr<int> ptr = &v[0];
 437      *  thrust::device_reference<int> ref(ptr);
 438      *
 439      *  // ref equals 0
 440      *  assert(ref == 0);
 441      *
 442      *  // the object pointed to by ptr equals 0
 443      *  assert(*ptr == 0);
 444      *
 445      *  // v[0] equals 0
 446      *  assert(v[0] == 0);
 447      *
 448      *  // add-assign ref
 449      *  ref += 5;
 450      *
 451      *  // ref equals 5
 452      *  assert(ref == 5);
 453      *
 454      *  // the object pointed to by ptr equals 5
 455      *  assert(*ptr == 5);
 456      *
 457      *  // v[0] equals 5
 458      *  assert(v[0] == 5);
 459      *  \endcode
 460      *
 461      *  \note The add-assignment executes as as if it were executed on the host.
 462      *  This may change in a later version.
 463      */
 464     device_reference &operator+=(const T &rhs);
 465
 466     /*! Prefix decrement operator decrements the object referenced by this
 467      *  \p device_reference.
 468      *
 469      *  \return <tt>*this</tt>
 470      *
 471      *  The following code snippet demonstrates the semantics of
 472      *  \p device_reference's prefix decrement operator.
 473      *
 474      *  \code
 475      *  #include <thrust/device_vector.h>
 476      *  #include <assert.h>
 477      *  ...
 478      *  thrust::device_vector<int> v(1,0);
 479      *  thrust::device_ptr<int> ptr = &v[0];
 480      *  thrust::device_reference<int> ref(ptr);
 481      *
 482      *  // ref equals 0
 483      *  assert(ref == 0);
 484      *
 485      *  // the object pointed to by ptr equals 0
 486      *  assert(*ptr == 0);
 487      *
 488      *  // v[0] equals 0
 489      *  assert(v[0] == 0);
 490      *
 491      *  // decrement ref
 492      *  --ref;
 493      *
 494      *  // ref equals -1
 495      *  assert(ref == -1);
 496      *
 497      *  // the object pointed to by ptr equals -1
 498      *  assert(*ptr == -1);
 499      *
 500      *  // v[0] equals -1
 501      *  assert(v[0] == -1);
 502      *  \endcode
 503      *
 504      *  \note The decrement executes as if it were executed on the host.
 505      *  This may change in a later version.
 506      */
 507     device_reference &operator--(void);
 508
 509     /*! Postfix decrement operator copies the object referenced by this
 510      *  \p device_reference, decrements the object referenced by this
 511      *  \p device_reference, and returns the copy.
 512      *
 513      *  \return A copy of the object referenced by this \p device_reference
 514      *          before being decremented.
 515      *
 516      *  The following code snippet demonstrates the semantics of
 517      *  \p device_reference's postfix decrement operator.
 518      *
 519      *  \code
 520      *  #include <thrust/device_vector.h>
 521      *  #include <assert.h>
 522      *  ...
 523      *  thrust::device_vector<int> v(1,0);
 524      *  thrust::device_ptr<int> ptr = &v[0];
 525      *  thrust::device_reference<int> ref(ptr);
 526      *
 527      *  // ref equals 0
 528      *  assert(ref == 0);
 529      *
 530      *  // the object pointed to by ptr equals 0
 531      *  assert(*ptr == 0);
 532      *
 533      *  // v[0] equals 0
 534      *  assert(v[0] == 0);
 535      *
 536      *  // decrement ref
 537      *  int x = ref--;
 538      *
 539      *  // x equals 0
 540      *  assert(x == 0)
 541      *
 542      *  // ref equals -1
 543      *  assert(ref == -1);
 544      *
 545      *  // the object pointed to by ptr equals -1
 546      *  assert(*ptr == -1);
 547      *
 548      *  // v[0] equals -1
 549      *  assert(v[0] == -1);
 550      *  \endcode
 551      *
 552      *  \note The decrement executes as if it were executed on the host.
 553      *  This may change in a later version.
 554      */
 555     value_type operator--(int);
 556
 557     /*! Subtraction assignment operator subtract-assigns the object referenced by this
 558      *  \p device_reference and returns this \p device_reference.
 559      *
 560      *  \param rhs The right hand side of the subtraction-assignment.
 561      *  \return <tt>*this</tt>.
 562      *
 563      *  The following code snippet demonstrates the semantics of
 564      *  \p device_reference's addition assignment operator.
 565      *
 566      *  \code
 567      *  #include <thrust/device_vector.h>
 568      *  #include <assert.h>
 569      *  ...
 570      *  thrust::device_vector<int> v(1,0);
 571      *  thrust::device_ptr<int> ptr = &v[0];
 572      *  thrust::device_reference<int> ref(ptr);
 573      *
 574      *  // ref equals 0
 575      *  assert(ref == 0);
 576      *
 577      *  // the object pointed to by ptr equals 0
 578      *  assert(*ptr == 0);
 579      *
 580      *  // v[0] equals 0
 581      *  assert(v[0] == 0);
 582      *
 583      *  // subtract-assign ref
 584      *  ref -= 5;
 585      *
 586      *  // ref equals -5
 587      *  assert(ref == -5);
 588      *
 589      *  // the object pointed to by ptr equals -5
 590      *  assert(*ptr == -5);
 591      *
 592      *  // v[0] equals -5
 593      *  assert(v[0] == -5);
 594      *  \endcode
 595      *
 596      *  \note The subtract-assignment executes as as if it were executed on the host.
 597      *  This may change in a later version.
 598      */
 599     device_reference &operator-=(const T &rhs);
 600
 601     /*! Multiplication assignment operator multiply-assigns the object referenced by this
 602      *  \p device_reference and returns this \p device_reference.
 603      *
 604      *  \param rhs The right hand side of the multiply-assignment.
 605      *  \return <tt>*this</tt>.
 606      *
 607      *  The following code snippet demonstrates the semantics of
 608      *  \p device_reference's multiply assignment operator.
 609      *
 610      *  \code
 611      *  #include <thrust/device_vector.h>
 612      *  #include <assert.h>
 613      *  ...
 614      *  thrust::device_vector<int> v(1,1);
 615      *  thrust::device_ptr<int> ptr = &v[0];
 616      *  thrust::device_reference<int> ref(ptr);
 617      *
 618      *  // ref equals 1
 619      *  assert(ref == 1);
 620      *
 621      *  // the object pointed to by ptr equals 1
 622      *  assert(*ptr == 1);
 623      *
 624      *  // v[0] equals 1
 625      *  assert(v[0] == 1);
 626      *
 627      *  // multiply-assign ref
 628      *  ref *= 5;
 629      *
 630      *  // ref equals 5
 631      *  assert(ref == 5);
 632      *
 633      *  // the object pointed to by ptr equals 5
 634      *  assert(*ptr == 5);
 635      *
 636      *  // v[0] equals 5
 637      *  assert(v[0] == 5);
 638      *  \endcode
 639      *
 640      *  \note The multiply-assignment executes as as if it were executed on the host.
 641      *  This may change in a later version.
 642      */
 643     device_reference &operator*=(const T &rhs);
 644
 645     /*! Division assignment operator divide-assigns the object referenced by this
 646      *  \p device_reference and returns this \p device_reference.
 647      *
 648      *  \param rhs The right hand side of the divide-assignment.
 649      *  \return <tt>*this</tt>.
 650      *
 651      *  The following code snippet demonstrates the semantics of
 652      *  \p device_reference's divide assignment operator.
 653      *
 654      *  \code
 655      *  #include <thrust/device_vector.h>
 656      *  #include <assert.h>
 657      *  ...
 658      *  thrust::device_vector<int> v(1,5);
 659      *  thrust::device_ptr<int> ptr = &v[0];
 660      *  thrust::device_reference<int> ref(ptr);
 661      *
 662      *  // ref equals 5
 663      *  assert(ref == 5);
 664      *
 665      *  // the object pointed to by ptr equals 5
 666      *  assert(*ptr == 5);
 667      *
 668      *  // v[0] equals 5
 669      *  assert(v[0] == 5);
 670      *
 671      *  // divide-assign ref
 672      *  ref /= 5;
 673      *
 674      *  // ref equals 1
 675      *  assert(ref == 1);
 676      *
 677      *  // the object pointed to by ptr equals 1
 678      *  assert(*ptr == 1);
 679      *
 680      *  // v[0] equals 1
 681      *  assert(v[0] == 1);
 682      *  \endcode
 683      *
 684      *  \note The divide-assignment executes as as if it were executed on the host.
 685      *  This may change in a later version.
 686      */
 687     device_reference &operator/=(const T &rhs);
 688
 689     /*! Modulation assignment operator modulus-assigns the object referenced by this
 690      *  \p device_reference and returns this \p device_reference.
 691      *
 692      *  \param rhs The right hand side of the divide-assignment.
 693      *  \return <tt>*this</tt>.
 694      *
 695      *  The following code snippet demonstrates the semantics of
 696      *  \p device_reference's divide assignment operator.
 697      *
 698      *  \code
 699      *  #include <thrust/device_vector.h>
 700      *  #include <assert.h>
 701      *  ...
 702      *  thrust::device_vector<int> v(1,5);
 703      *  thrust::device_ptr<int> ptr = &v[0];
 704      *  thrust::device_reference<int> ref(ptr);
 705      *
 706      *  // ref equals 5
 707      *  assert(ref == 5);
 708      *
 709      *  // the object pointed to by ptr equals 5
 710      *  assert(*ptr == 5);
 711      *
 712      *  // v[0] equals 5
 713      *  assert(v[0] == 5);
 714      *
 715      *  // modulus-assign ref
 716      *  ref %= 5;
 717      *
 718      *  // ref equals 0
 719      *  assert(ref == 0);
 720      *
 721      *  // the object pointed to by ptr equals 0
 722      *  assert(*ptr == 0);
 723      *
 724      *  // v[0] equals 0
 725      *  assert(v[0] == 0);
 726      *  \endcode
 727      *
 728      *  \note The modulus-assignment executes as as if it were executed on the host.
 729      *  This may change in a later version.
 730      */
 731     device_reference &operator%=(const T &rhs);
 732
 733     /*! Bitwise left shift assignment operator left shift-assigns the object referenced by this
 734      *  \p device_reference and returns this \p device_reference.
 735      *
 736      *  \param rhs The right hand side of the left shift-assignment.
 737      *  \return <tt>*this</tt>.
 738      *
 739      *  The following code snippet demonstrates the semantics of
 740      *  \p device_reference's left shift assignment operator.
 741      *
 742      *  \code
 743      *  #include <thrust/device_vector.h>
 744      *  #include <assert.h>
 745      *  ...
 746      *  thrust::device_vector<int> v(1,1);
 747      *  thrust::device_ptr<int> ptr = &v[0];
 748      *  thrust::device_reference<int> ref(ptr);
 749      *
 750      *  // ref equals 1
 751      *  assert(ref == 1);
 752      *
 753      *  // the object pointed to by ptr equals 1
 754      *  assert(*ptr == 1);
 755      *
 756      *  // v[0] equals 1
 757      *  assert(v[0] == 1);
 758      *
 759      *  // left shift-assign ref
 760      *  ref <<= 1;
 761      *
 762      *  // ref equals 2
 763      *  assert(ref == 2);
 764      *
 765      *  // the object pointed to by ptr equals 2
 766      *  assert(*ptr == 2);
 767      *
 768      *  // v[0] equals 2
 769      *  assert(v[0] == 2);
 770      *  \endcode
 771      *
 772      *  \note The left shift-assignment executes as as if it were executed on the host.
 773      *  This may change in a later version.
 774      */
 775     device_reference &operator<<=(const T &rhs);
 776
 777     /*! Bitwise right shift assignment operator right shift-assigns the object referenced by this
 778      *  \p device_reference and returns this \p device_reference.
 779      *
 780      *  \param rhs The right hand side of the right shift-assignment.
 781      *  \return <tt>*this</tt>.
 782      *
 783      *  The following code snippet demonstrates the semantics of
 784      *  \p device_reference's right shift assignment operator.
 785      *
 786      *  \code
 787      *  #include <thrust/device_vector.h>
 788      *  #include <assert.h>
 789      *  ...
 790      *  thrust::device_vector<int> v(1,2);
 791      *  thrust::device_ptr<int> ptr = &v[0];
 792      *  thrust::device_reference<int> ref(ptr);
 793      *
 794      *  // ref equals 2
 795      *  assert(ref == 2);
 796      *
 797      *  // the object pointed to by ptr equals 2
 798      *  assert(*ptr == 2);
 799      *
 800      *  // v[0] equals 2
 801      *  assert(v[0] == 2);
 802      *
 803      *  // right shift-assign ref
 804      *  ref >>= 1;
 805      *
 806      *  // ref equals 1
 807      *  assert(ref == 1);
 808      *
 809      *  // the object pointed to by ptr equals 1
 810      *  assert(*ptr == 1);
 811      *
 812      *  // v[0] equals 1
 813      *  assert(v[0] == 1);
 814      *  \endcode
 815      *
 816      *  \note The right shift-assignment executes as as if it were executed on the host.
 817      *  This may change in a later version.
 818      */
 819     device_reference &operator>>=(const T &rhs);
 820
 821     /*! Bitwise AND assignment operator AND-assigns the object referenced by this
 822      *  \p device_reference and returns this \p device_reference.
 823      *
 824      *  \param rhs The right hand side of the AND-assignment.
 825      *  \return <tt>*this</tt>.
 826      *
 827      *  The following code snippet demonstrates the semantics of
 828      *  \p device_reference's AND assignment operator.
 829      *
 830      *  \code
 831      *  #include <thrust/device_vector.h>
 832      *  #include <assert.h>
 833      *  ...
 834      *  thrust::device_vector<int> v(1,1);
 835      *  thrust::device_ptr<int> ptr = &v[0];
 836      *  thrust::device_reference<int> ref(ptr);
 837      *
 838      *  // ref equals 1
 839      *  assert(ref == 1);
 840      *
 841      *  // the object pointed to by ptr equals 1
 842      *  assert(*ptr == 1);
 843      *
 844      *  // v[0] equals 1
 845      *  assert(v[0] == 1);
 846      *
 847      *  // right AND-assign ref
 848      *  ref &= 0;
 849      *
 850      *  // ref equals 0
 851      *  assert(ref == 0);
 852      *
 853      *  // the object pointed to by ptr equals 0
 854      *  assert(*ptr == 0);
 855      *
 856      *  // v[0] equals 0
 857      *  assert(v[0] == 0);
 858      *  \endcode
 859      *
 860      *  \note The AND-assignment executes as as if it were executed on the host.
 861      *  This may change in a later version.
 862      */
 863     device_reference &operator&=(const T &rhs);
 864
 865     /*! Bitwise OR assignment operator OR-assigns the object referenced by this
 866      *  \p device_reference and returns this \p device_reference.
 867      *
 868      *  \param rhs The right hand side of the OR-assignment.
 869      *  \return <tt>*this</tt>.
 870      *
 871      *  The following code snippet demonstrates the semantics of
 872      *  \p device_reference's OR assignment operator.
 873      *
 874      *  \code
 875      *  #include <thrust/device_vector.h>
 876      *  #include <assert.h>
 877      *  ...
 878      *  thrust::device_vector<int> v(1,0);
 879      *  thrust::device_ptr<int> ptr = &v[0];
 880      *  thrust::device_reference<int> ref(ptr);
 881      *
 882      *  // ref equals 0
 883      *  assert(ref == 0);
 884      *
 885      *  // the object pointed to by ptr equals 0
 886      *  assert(*ptr == 0);
 887      *
 888      *  // v[0] equals 0
 889      *  assert(v[0] == 0);
 890      *
 891      *  // right OR-assign ref
 892      *  ref |= 1;
 893      *
 894      *  // ref equals 1
 895      *  assert(ref == 1);
 896      *
 897      *  // the object pointed to by ptr equals 1
 898      *  assert(*ptr == 1);
 899      *
 900      *  // v[0] equals 1
 901      *  assert(v[0] == 1);
 902      *  \endcode
 903      *
 904      *  \note The OR-assignment executes as as if it were executed on the host.
 905      *  This may change in a later version.
 906      */
 907     device_reference &operator|=(const T &rhs);
 908
 909     /*! Bitwise XOR assignment operator XOR-assigns the object referenced by this
 910      *  \p device_reference and returns this \p device_reference.
 911      *
 912      *  \param rhs The right hand side of the XOR-assignment.
 913      *  \return <tt>*this</tt>.
 914      *
 915      *  The following code snippet demonstrates the semantics of
 916      *  \p device_reference's XOR assignment operator.
 917      *
 918      *  \code
 919      *  #include <thrust/device_vector.h>
 920      *  #include <assert.h>
 921      *  ...
 922      *  thrust::device_vector<int> v(1,1);
 923      *  thrust::device_ptr<int> ptr = &v[0];
 924      *  thrust::device_reference<int> ref(ptr);
 925      *
 926      *  // ref equals 1
 927      *  assert(ref == 1);
 928      *
 929      *  // the object pointed to by ptr equals 1
 930      *  assert(*ptr == 1);
 931      *
 932      *  // v[0] equals 1
 933      *  assert(v[0] == 1);
 934      *
 935      *  // right XOR-assign ref
 936      *  ref ^= 1;
 937      *
 938      *  // ref equals 0
 939      *  assert(ref == 0);
 940      *
 941      *  // the object pointed to by ptr equals 0
 942      *  assert(*ptr == 0);
 943      *
 944      *  // v[0] equals 0
 945      *  assert(v[0] == 0);
 946      *  \endcode
 947      *
 948      *  \note The XOR-assignment executes as as if it were executed on the host.
 949      *  This may change in a later version.
 950      */
 951     device_reference &operator^=(const T &rhs);
 952 #endif // end doxygen-only members
 953 }; // end device_reference
 954
 955 /*! swaps the value of one \p device_reference with another.
 956  *  \p x The first \p device_reference of interest.
 957  *  \p y The second \p device_reference of interest.
 958  */
 959 template<typename T>
 960 __host__ __device__
 961 void swap(device_reference<T> &x, device_reference<T> &y);
 962
 963 /*! \}
 964  */
 965
 966 } // end thrust
 967
 968 #include <thrust/detail/device_reference.inl>
 969