util/X86LINUX64/cuda-6.5/include/thrust/detail/vector_base.inl

   1 /*
   2  *  Copyright 2008-2013 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 vector_base.inl
  19  *  \brief Inline file for vector_base.h.
  20  */
  21
  22 #include <thrust/detail/vector_base.h>
  23 #include <thrust/detail/copy.h>
  24 #include <thrust/detail/overlapped_copy.h>
  25 #include <thrust/equal.h>
  26 #include <thrust/distance.h>
  27 #include <thrust/advance.h>
  28 #include <thrust/detail/type_traits.h>
  29 #include <thrust/detail/minmax.h>
  30 #include <thrust/iterator/iterator_traits.h>
  31 #include <thrust/detail/temporary_array.h>
  32
  33 #include <stdexcept>
  34
  35 namespace thrust
  36 {
  37
  38 namespace detail
  39 {
  40
  41 template<typename T, typename Alloc>
  42   vector_base<T,Alloc>
  43     ::vector_base(void)
  44       :m_storage(),
  45        m_size(0)
  46 {
  47   ;
  48 } // end vector_base::vector_base()
  49
  50 template<typename T, typename Alloc>
  51   vector_base<T,Alloc>
  52     ::vector_base(size_type n)
  53       :m_storage(),
  54        m_size(0)
  55 {
  56   default_init(n);
  57 } // end vector_base::vector_base()
  58
  59 template<typename T, typename Alloc>
  60   vector_base<T,Alloc>
  61     ::vector_base(size_type n, const value_type &value)
  62       :m_storage(),
  63        m_size(0)
  64 {
  65   fill_init(n,value);
  66 } // end vector_base::vector_base()
  67
  68 template<typename T, typename Alloc>
  69   vector_base<T,Alloc>
  70     ::vector_base(const vector_base &v)
  71       :m_storage(),
  72        m_size(0)
  73 {
  74   range_init(v.begin(), v.end());
  75 } // end vector_base::vector_base()
  76
  77 template<typename T, typename Alloc>
  78   vector_base<T,Alloc> &
  79     vector_base<T,Alloc>
  80       ::operator=(const vector_base &v)
  81 {
  82   if(this != &v)
  83   {
  84     assign(v.begin(), v.end());
  85   } // end if
  86
  87   return *this;
  88 } // end vector_base::operator=()
  89
  90 template<typename T, typename Alloc>
  91   template<typename OtherT, typename OtherAlloc>
  92     vector_base<T,Alloc>
  93       ::vector_base(const vector_base<OtherT,OtherAlloc> &v)
  94         :m_storage(),
  95          m_size(0)
  96 {
  97   range_init(v.begin(), v.end());
  98 } // end vector_base::vector_base()
  99
 100 template<typename T, typename Alloc>
 101   template<typename OtherT, typename OtherAlloc>
 102     vector_base<T,Alloc> &
 103       vector_base<T,Alloc>
 104         ::operator=(const vector_base<OtherT,OtherAlloc> &v)
 105 {
 106   assign(v.begin(), v.end());
 107
 108   return *this;
 109 } // end vector_base::operator=()
 110
 111 template<typename T, typename Alloc>
 112   template<typename OtherT, typename OtherAlloc>
 113     vector_base<T,Alloc>
 114       ::vector_base(const std::vector<OtherT,OtherAlloc> &v)
 115         :m_storage(),
 116          m_size(0)
 117 {
 118   range_init(v.begin(), v.end());
 119 } // end vector_base::vector_base()
 120
 121 template<typename T, typename Alloc>
 122   template<typename OtherT, typename OtherAlloc>
 123     vector_base<T,Alloc> &
 124       vector_base<T,Alloc>
 125         ::operator=(const std::vector<OtherT,OtherAlloc> &v)
 126 {
 127   assign(v.begin(), v.end());
 128
 129   return *this;
 130 } // end vector_base::operator=()
 131
 132 template<typename T, typename Alloc>
 133   template<typename IteratorOrIntegralType>
 134     void vector_base<T,Alloc>
 135       ::init_dispatch(IteratorOrIntegralType n,
 136                       IteratorOrIntegralType value,
 137                       true_type)
 138 {
 139   fill_init(n,value);
 140 } // end vector_base::init_dispatch()
 141
 142 template<typename T, typename Alloc>
 143   void vector_base<T,Alloc>
 144     ::default_init(size_type n)
 145 {
 146   if(n > 0)
 147   {
 148     m_storage.allocate(n);
 149     m_size = n;
 150
 151     m_storage.default_construct_n(begin(), size());
 152   } // end if
 153 } // end vector_base::default_init()
 154
 155 template<typename T, typename Alloc>
 156   void vector_base<T,Alloc>
 157     ::fill_init(size_type n, const T &x)
 158 {
 159   if(n > 0)
 160   {
 161     m_storage.allocate(n);
 162     m_size = n;
 163
 164     m_storage.uninitialized_fill_n(begin(), size(), x);
 165   } // end if
 166 } // end vector_base::fill_init()
 167
 168 template<typename T, typename Alloc>
 169   template<typename InputIterator>
 170     void vector_base<T,Alloc>
 171       ::init_dispatch(InputIterator first,
 172                       InputIterator last,
 173                       false_type)
 174 {
 175   range_init(first, last);
 176 } // end vector_base::init_dispatch()
 177
 178 template<typename T, typename Alloc>
 179   template<typename InputIterator>
 180     void vector_base<T,Alloc>
 181       ::range_init(InputIterator first,
 182                    InputIterator last)
 183 {
 184   range_init(first, last,
 185     typename thrust::iterator_traversal<InputIterator>::type());
 186 } // end vector_base::range_init()
 187
 188 template<typename T, typename Alloc>
 189   template<typename InputIterator>
 190     void vector_base<T,Alloc>
 191       ::range_init(InputIterator first,
 192                    InputIterator last,
 193                    thrust::incrementable_traversal_tag)
 194 {
 195   for(; first != last; ++first)
 196     push_back(*first);
 197 } // end vector_base::range_init()
 198
 199 template<typename T, typename Alloc>
 200   template<typename ForwardIterator>
 201     void vector_base<T,Alloc>
 202       ::range_init(ForwardIterator first,
 203                    ForwardIterator last,
 204                    thrust::random_access_traversal_tag)
 205 {
 206   size_type new_size = thrust::distance(first, last);
 207
 208   allocate_and_copy(new_size, first, last, m_storage);
 209   m_size    = new_size;
 210 } // end vector_base::range_init()
 211
 212 template<typename T, typename Alloc>
 213   template<typename InputIterator>
 214     vector_base<T,Alloc>
 215       ::vector_base(InputIterator first,
 216                     InputIterator last)
 217         :m_storage(),
 218          m_size(0)
 219 {
 220   // check the type of InputIterator: if it's an integral type,
 221   // we need to interpret this call as (size_type, value_type)
 222   typedef thrust::detail::is_integral<InputIterator> Integer;
 223
 224   init_dispatch(first, last, Integer());
 225 } // end vector_basee::vector_base()
 226
 227 template<typename T, typename Alloc>
 228   void vector_base<T,Alloc>
 229     ::resize(size_type new_size)
 230 {
 231   if(new_size < size())
 232   {
 233     iterator new_end = begin();
 234     thrust::advance(new_end, new_size);
 235     erase(new_end, end());
 236   } // end if
 237   else
 238   {
 239     append(new_size - size());
 240   } // end else
 241 } // end vector_base::resize()
 242
 243 template<typename T, typename Alloc>
 244   void vector_base<T,Alloc>
 245     ::resize(size_type new_size, const value_type &x)
 246 {
 247   if(new_size < size())
 248   {
 249     iterator new_end = begin();
 250     thrust::advance(new_end, new_size);
 251     erase(new_end, end());
 252   } // end if
 253   else
 254   {
 255     insert(end(), new_size - size(), x);
 256   } // end else
 257 } // end vector_base::resize()
 258
 259 template<typename T, typename Alloc>
 260   typename vector_base<T,Alloc>::size_type
 261     vector_base<T,Alloc>
 262       ::size(void) const
 263 {
 264   return m_size;
 265 } // end vector_base::size()
 266
 267 template<typename T, typename Alloc>
 268   typename vector_base<T,Alloc>::size_type
 269     vector_base<T,Alloc>
 270       ::max_size(void) const
 271 {
 272   return m_storage.max_size();
 273 } // end vector_base::max_size()
 274
 275 template<typename T, typename Alloc>
 276   void vector_base<T,Alloc>
 277     ::reserve(size_type n)
 278 {
 279   if(n > capacity())
 280   {
 281     allocate_and_copy(n, begin(), end(), m_storage);
 282   } // end if
 283 } // end vector_base::reserve()
 284
 285 template<typename T, typename Alloc>
 286   typename vector_base<T,Alloc>::size_type
 287     vector_base<T,Alloc>
 288       ::capacity(void) const
 289 {
 290   return m_storage.size();
 291 } // end vector_base::capacity()
 292
 293 template<typename T, typename Alloc>
 294   void vector_base<T,Alloc>
 295     ::shrink_to_fit(void)
 296 {
 297   // use the swap trick
 298   vector_base(*this).swap(*this);
 299 } // end vector_base::shrink_to_fit()
 300
 301 template<typename T, typename Alloc>
 302   typename vector_base<T,Alloc>::reference
 303     vector_base<T,Alloc>
 304       ::operator[](const size_type n)
 305 {
 306   return m_storage[n];
 307 } // end vector_base::operator[]
 308
 309 template<typename T, typename Alloc>
 310   typename vector_base<T,Alloc>::const_reference
 311     vector_base<T,Alloc>
 312       ::operator[](const size_type n) const
 313 {
 314   return m_storage[n];
 315 } // end vector_base::operator[]
 316
 317 template<typename T, typename Alloc>
 318   typename vector_base<T,Alloc>::iterator
 319     vector_base<T,Alloc>
 320       ::begin(void)
 321 {
 322   return m_storage.begin();
 323 } // end vector_base::begin()
 324
 325 template<typename T, typename Alloc>
 326   typename vector_base<T,Alloc>::const_iterator
 327     vector_base<T,Alloc>
 328       ::begin(void) const
 329 {
 330   return m_storage.begin();
 331 } // end vector_base::begin()
 332
 333 template<typename T, typename Alloc>
 334   typename vector_base<T,Alloc>::const_iterator
 335     vector_base<T,Alloc>
 336       ::cbegin(void) const
 337 {
 338   return begin();
 339 } // end vector_base::cbegin()
 340
 341 template<typename T, typename Alloc>
 342   typename vector_base<T,Alloc>::reverse_iterator
 343     vector_base<T,Alloc>
 344       ::rbegin(void)
 345 {
 346   return reverse_iterator(end());
 347 } // end vector_base::rbegin()
 348
 349 template<typename T, typename Alloc>
 350   typename vector_base<T,Alloc>::const_reverse_iterator
 351     vector_base<T,Alloc>
 352       ::rbegin(void) const
 353 {
 354   return const_reverse_iterator(end());
 355 } // end vector_base::rbegin()
 356
 357 template<typename T, typename Alloc>
 358   typename vector_base<T,Alloc>::const_reverse_iterator
 359     vector_base<T,Alloc>
 360       ::crbegin(void) const
 361 {
 362   return rbegin();
 363 } // end vector_base::crbegin()
 364
 365 template<typename T, typename Alloc>
 366   typename vector_base<T,Alloc>::iterator
 367     vector_base<T,Alloc>
 368       ::end(void)
 369 {
 370   iterator result = begin();
 371   thrust::advance(result, size());
 372   return result;
 373 } // end vector_base::end()
 374
 375 template<typename T, typename Alloc>
 376   typename vector_base<T,Alloc>::const_iterator
 377     vector_base<T,Alloc>
 378       ::end(void) const
 379 {
 380   const_iterator result = begin();
 381   thrust::advance(result, size());
 382   return result;
 383 } // end vector_base::end()
 384
 385 template<typename T, typename Alloc>
 386   typename vector_base<T,Alloc>::const_iterator
 387     vector_base<T,Alloc>
 388       ::cend(void) const
 389 {
 390   return end();
 391 } // end vector_base::cend()
 392
 393 template<typename T, typename Alloc>
 394   typename vector_base<T,Alloc>::reverse_iterator
 395     vector_base<T,Alloc>
 396       ::rend(void)
 397 {
 398   return reverse_iterator(begin());
 399 } // end vector_base::rend()
 400
 401 template<typename T, typename Alloc>
 402   typename vector_base<T,Alloc>::const_reverse_iterator
 403     vector_base<T,Alloc>
 404       ::rend(void) const
 405 {
 406   return const_reverse_iterator(begin());
 407 } // end vector_base::rend()
 408
 409 template<typename T, typename Alloc>
 410   typename vector_base<T,Alloc>::const_reverse_iterator
 411     vector_base<T,Alloc>
 412       ::crend(void) const
 413 {
 414   return rend();
 415 } // end vector_base::crend()
 416
 417 template<typename T, typename Alloc>
 418   typename vector_base<T,Alloc>::const_reference
 419     vector_base<T,Alloc>
 420       ::front(void) const
 421 {
 422   return *begin();
 423 } // end vector_base::front()
 424
 425 template<typename T, typename Alloc>
 426   typename vector_base<T,Alloc>::reference
 427     vector_base<T,Alloc>
 428       ::front(void)
 429 {
 430   return *begin();
 431 } // end vector_base::front()
 432
 433 template<typename T, typename Alloc>
 434   typename vector_base<T,Alloc>::const_reference
 435     vector_base<T,Alloc>
 436       ::back(void) const
 437 {
 438   const_iterator ptr_to_back = end();
 439   --ptr_to_back;
 440   return *ptr_to_back;
 441 } // end vector_base::vector_base
 442
 443 template<typename T, typename Alloc>
 444   typename vector_base<T,Alloc>::reference
 445     vector_base<T,Alloc>
 446       ::back(void)
 447 {
 448   iterator ptr_to_back = end();
 449   --ptr_to_back;
 450   return *ptr_to_back;
 451 } // end vector_base::vector_base
 452
 453 template<typename T, typename Alloc>
 454   typename vector_base<T,Alloc>::pointer
 455     vector_base<T,Alloc>
 456       ::data(void)
 457 {
 458   return &front();
 459 } // end vector_base::data()
 460
 461 template<typename T, typename Alloc>
 462   typename vector_base<T,Alloc>::const_pointer
 463     vector_base<T,Alloc>
 464       ::data(void) const
 465 {
 466   return &front();
 467 } // end vector_base::data()
 468
 469 template<typename T, typename Alloc>
 470   vector_base<T,Alloc>
 471     ::~vector_base(void)
 472 {
 473   // destroy every living thing
 474   m_storage.destroy(begin(),end());
 475 } // end vector_base::~vector_base()
 476
 477 template<typename T, typename Alloc>
 478   void vector_base<T,Alloc>
 479     ::clear(void)
 480 {
 481   resize(0);
 482 } // end vector_base::~vector_dev()
 483
 484 template<typename T, typename Alloc>
 485   bool vector_base<T,Alloc>
 486     ::empty(void) const
 487 {
 488   return size() == 0;
 489 } // end vector_base::empty();
 490
 491 template<typename T, typename Alloc>
 492   void vector_base<T,Alloc>
 493     ::push_back(const value_type &x)
 494 {
 495   insert(end(), x);
 496 } // end vector_base::push_back()
 497
 498 template<typename T, typename Alloc>
 499   void vector_base<T,Alloc>
 500     ::pop_back(void)
 501 {
 502   iterator e = end();
 503   iterator ptr_to_back = e;
 504   --ptr_to_back;
 505   m_storage.destroy(ptr_to_back, e);
 506   --m_size;
 507 } // end vector_base::pop_back()
 508
 509 template<typename T, typename Alloc>
 510   typename vector_base<T,Alloc>::iterator vector_base<T,Alloc>
 511     ::erase(iterator pos)
 512 {
 513   iterator end = pos;
 514   ++end;
 515   return erase(pos,end);
 516 } // end vector_base::erase()
 517
 518 template<typename T, typename Alloc>
 519   typename vector_base<T,Alloc>::iterator vector_base<T,Alloc>
 520     ::erase(iterator first, iterator last)
 521 {
 522   // overlap copy the range [last,end()) to first
 523   // XXX this copy only potentially overlaps
 524   iterator i = thrust::detail::overlapped_copy(last, end(), first);
 525
 526   // destroy everything after i
 527   m_storage.destroy(i, end());
 528
 529   // modify our size
 530   m_size -= (last - first);
 531
 532   // return an iterator pointing to the position of the first element
 533   // following the erased range
 534   return first;
 535 } // end vector_base::erase()
 536
 537 template<typename T, typename Alloc>
 538   void vector_base<T,Alloc>
 539     ::swap(vector_base &v)
 540 {
 541   thrust::swap(m_storage,  v.m_storage);
 542   thrust::swap(m_size,     v.m_size);
 543 } // end vector_base::swap()
 544
 545 template<typename T, typename Alloc>
 546   void vector_base<T,Alloc>
 547     ::assign(size_type n, const T &x)
 548 {
 549   fill_assign(n, x);
 550 } // end vector_base::assign()
 551
 552 template<typename T, typename Alloc>
 553   template<typename InputIterator>
 554     void vector_base<T,Alloc>
 555       ::assign(InputIterator first, InputIterator last)
 556 {
 557   // we could have received assign(n, x), so disambiguate on the
 558   // type of InputIterator
 559   typedef typename thrust::detail::is_integral<InputIterator> integral;
 560
 561   assign_dispatch(first, last, integral());
 562 } // end vector_base::assign()
 563
 564 template<typename T, typename Alloc>
 565   typename vector_base<T,Alloc>::allocator_type
 566     vector_base<T,Alloc>
 567       ::get_allocator(void) const
 568 {
 569   return m_storage.get_allocator();
 570 } // end vector_base::get_allocator()
 571
 572 template<typename T, typename Alloc>
 573   typename vector_base<T,Alloc>::iterator
 574     vector_base<T,Alloc>
 575       ::insert(iterator position, const T &x)
 576 {
 577   // find the index of the insertion
 578   size_type index = thrust::distance(begin(), position);
 579
 580   // make the insertion
 581   insert(position, 1, x);
 582
 583   // return an iterator pointing back to position
 584   iterator result = begin();
 585   thrust::advance(result, index);
 586   return result;
 587 } // end vector_base::insert()
 588
 589 template<typename T, typename Alloc>
 590   void vector_base<T,Alloc>
 591     ::insert(iterator position, size_type n, const T &x)
 592 {
 593   fill_insert(position, n, x);
 594 } // end vector_base::insert()
 595
 596 template<typename T, typename Alloc>
 597   template<typename InputIterator>
 598     void vector_base<T,Alloc>
 599       ::insert(iterator position, InputIterator first, InputIterator last)
 600 {
 601   // we could have received insert(position, n, x), so disambiguate on the
 602   // type of InputIterator
 603   typedef typename thrust::detail::is_integral<InputIterator> integral;
 604
 605   insert_dispatch(position, first, last, integral());
 606 } // end vector_base::insert()
 607
 608 template<typename T, typename Alloc>
 609   template<typename InputIterator>
 610     void vector_base<T,Alloc>
 611       ::assign_dispatch(InputIterator first, InputIterator last, false_type)
 612 {
 613   range_assign(first, last);
 614 } // end vector_base::assign_dispatch()
 615
 616 template<typename T, typename Alloc>
 617   template<typename Integral>
 618     void vector_base<T,Alloc>
 619       ::assign_dispatch(Integral n, Integral x, true_type)
 620 {
 621   fill_assign(n, x);
 622 } // end vector_base::assign_dispatch()
 623
 624 template<typename T, typename Alloc>
 625   template<typename InputIterator>
 626     void vector_base<T,Alloc>
 627       ::insert_dispatch(iterator position, InputIterator first, InputIterator last, false_type)
 628 {
 629   copy_insert(position, first, last);
 630 } // end vector_base::insert_dispatch()
 631
 632 template<typename T, typename Alloc>
 633   template<typename Integral>
 634     void vector_base<T,Alloc>
 635       ::insert_dispatch(iterator position, Integral n, Integral x, true_type)
 636 {
 637   fill_insert(position, n, x);
 638 } // end vector_base::insert_dispatch()
 639
 640 template<typename T, typename Alloc>
 641   template<typename ForwardIterator>
 642     void vector_base<T,Alloc>
 643       ::copy_insert(iterator position,
 644                     ForwardIterator first,
 645                     ForwardIterator last)
 646 {
 647   if(first != last)
 648   {
 649     // how many new elements will we create?
 650     const size_type num_new_elements = thrust::distance(first, last);
 651     if(capacity() - size() >= num_new_elements)
 652     {
 653       // we've got room for all of them
 654       // how many existing elements will we displace?
 655       const size_type num_displaced_elements = end() - position;
 656       iterator old_end = end();
 657
 658       if(num_displaced_elements > num_new_elements)
 659       {
 660         // construct copy n displaced elements to new elements
 661         // following the insertion
 662         m_storage.uninitialized_copy(end() - num_new_elements, end(), end());
 663
 664         // extend the size
 665         m_size += num_new_elements;
 666
 667         // copy num_displaced_elements - num_new_elements elements to existing elements
 668         // this copy overlaps
 669         const size_type copy_length = (old_end - num_new_elements) - position;
 670         thrust::detail::overlapped_copy(position, old_end - num_new_elements, old_end - copy_length);
 671
 672         // finally, copy the range to the insertion point
 673         thrust::copy(first, last, position);
 674       } // end if
 675       else
 676       {
 677         ForwardIterator mid = first;
 678         thrust::advance(mid, num_displaced_elements);
 679
 680         // construct copy new elements at the end of the vector
 681         m_storage.uninitialized_copy(mid, last, end());
 682
 683         // extend the size
 684         m_size += num_new_elements - num_displaced_elements;
 685
 686         // construct copy the displaced elements
 687         m_storage.uninitialized_copy(position, old_end, end());
 688
 689         // extend the size
 690         m_size += num_displaced_elements;
 691
 692         // copy to elements which already existed
 693         thrust::copy(first, mid, position);
 694       } // end else
 695     } // end if
 696     else
 697     {
 698       const size_type old_size = size();
 699
 700       // compute the new capacity after the allocation
 701       size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION (old_size, num_new_elements);
 702
 703       // allocate exponentially larger new storage
 704       new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION <size_type>(new_capacity, 2 * capacity());
 705
 706       // do not exceed maximum storage
 707       new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION <size_type>(new_capacity, max_size());
 708
 709       if(new_capacity > max_size())
 710       {
 711         throw std::length_error("insert(): insertion exceeds max_size().");
 712       } // end if
 713
 714       storage_type new_storage(new_capacity);
 715
 716       // record how many constructors we invoke in the try block below
 717       iterator new_end = new_storage.begin();
 718
 719       try
 720       {
 721         // construct copy elements before the insertion to the beginning of the newly
 722         // allocated storage
 723         new_end = m_storage.uninitialized_copy(begin(), position, new_storage.begin());
 724
 725         // construct copy elements to insert
 726         new_end = m_storage.uninitialized_copy(first, last, new_end);
 727
 728         // construct copy displaced elements from the old storage to the new storage
 729         // remember [position, end()) refers to the old storage
 730         new_end = m_storage.uninitialized_copy(position, end(), new_end);
 731       } // end try
 732       catch(...)
 733       {
 734         // something went wrong, so destroy & deallocate the new storage
 735         m_storage.destroy(new_storage.begin(), new_end);
 736         new_storage.deallocate();
 737
 738         // rethrow
 739         throw;
 740       } // end catch
 741
 742       // call destructors on the elements in the old storage
 743       m_storage.destroy(begin(), end());
 744
 745       // record the vector's new state
 746       m_storage.swap(new_storage);
 747       m_size = old_size + num_new_elements;
 748     } // end else
 749   } // end if
 750 } // end vector_base::copy_insert()
 751
 752 template<typename T, typename Alloc>
 753   void vector_base<T,Alloc>
 754     ::append(size_type n)
 755 {
 756   if(n != 0)
 757   {
 758     if(capacity() - size() >= n)
 759     {
 760       // we've got room for all of them
 761
 762       // default construct new elements at the end of the vector
 763       m_storage.default_construct_n(end(), n);
 764
 765       // extend the size
 766       m_size += n;
 767     } // end if
 768     else
 769     {
 770       const size_type old_size = size();
 771
 772       // compute the new capacity after the allocation
 773       size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION (old_size, n);
 774
 775       // allocate exponentially larger new storage
 776       new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION <size_type>(new_capacity, 2 * capacity());
 777
 778       // do not exceed maximum storage
 779       new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION <size_type>(new_capacity, max_size());
 780
 781       // create new storage
 782       storage_type new_storage(new_capacity);
 783
 784       // record how many constructors we invoke in the try block below
 785       iterator new_end = new_storage.begin();
 786
 787       try
 788       {
 789         // construct copy all elements into the newly allocated storage
 790         new_end = m_storage.uninitialized_copy(begin(), end(), new_storage.begin());
 791
 792         // construct new elements to insert
 793         m_storage.default_construct_n(new_end, n);
 794         new_end += n;
 795       } // end try
 796       catch(...)
 797       {
 798         // something went wrong, so destroy & deallocate the new storage
 799         m_storage.destroy(new_storage.begin(), new_end);
 800         new_storage.deallocate();
 801
 802         // rethrow
 803         throw;
 804       } // end catch
 805
 806       // call destructors on the elements in the old storage
 807       m_storage.destroy(begin(), end());
 808
 809       // record the vector's new state
 810       m_storage.swap(new_storage);
 811       m_size    = old_size + n;
 812     } // end else
 813   } // end if
 814 } // end vector_base::append()
 815
 816 template<typename T, typename Alloc>
 817   void vector_base<T,Alloc>
 818     ::fill_insert(iterator position, size_type n, const T &x)
 819 {
 820   if(n != 0)
 821   {
 822     if(capacity() - size() >= n)
 823     {
 824       // we've got room for all of them
 825       // how many existing elements will we displace?
 826       const size_type num_displaced_elements = end() - position;
 827       iterator old_end = end();
 828
 829       if(num_displaced_elements > n)
 830       {
 831         // construct copy n displaced elements to new elements
 832         // following the insertion
 833         m_storage.uninitialized_copy(end() - n, end(), end());
 834
 835         // extend the size
 836         m_size += n;
 837
 838         // copy num_displaced_elements - n elements to existing elements
 839         // this copy overlaps
 840         const size_type copy_length = (old_end - n) - position;
 841         thrust::detail::overlapped_copy(position, old_end - n, old_end - copy_length);
 842
 843         // finally, fill the range to the insertion point
 844         thrust::fill_n(position, n, x);
 845       } // end if
 846       else
 847       {
 848         // construct new elements at the end of the vector
 849         m_storage.uninitialized_fill_n(end(), n - num_displaced_elements, x);
 850
 851         // extend the size
 852         m_size += n - num_displaced_elements;
 853
 854         // construct copy the displaced elements
 855         m_storage.uninitialized_copy(position, old_end, end());
 856
 857         // extend the size
 858         m_size += num_displaced_elements;
 859
 860         // fill to elements which already existed
 861         thrust::fill(position, old_end, x);
 862       } // end else
 863     } // end if
 864     else
 865     {
 866       const size_type old_size = size();
 867
 868       // compute the new capacity after the allocation
 869       size_type new_capacity = old_size + thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION (old_size, n);
 870
 871       // allocate exponentially larger new storage
 872       new_capacity = thrust::max THRUST_PREVENT_MACRO_SUBSTITUTION <size_type>(new_capacity, 2 * capacity());
 873
 874       // do not exceed maximum storage
 875       new_capacity = thrust::min THRUST_PREVENT_MACRO_SUBSTITUTION <size_type>(new_capacity, max_size());
 876
 877       if(new_capacity > max_size())
 878       {
 879         throw std::length_error("insert(): insertion exceeds max_size().");
 880       } // end if
 881
 882       storage_type new_storage(new_capacity);
 883
 884       // record how many constructors we invoke in the try block below
 885       iterator new_end = new_storage.begin();
 886
 887       try
 888       {
 889         // construct copy elements before the insertion to the beginning of the newly
 890         // allocated storage
 891         new_end = m_storage.uninitialized_copy(begin(), position, new_storage.begin());
 892
 893         // construct new elements to insert
 894         m_storage.uninitialized_fill_n(new_end, n, x);
 895         new_end += n;
 896
 897         // construct copy displaced elements from the old storage to the new storage
 898         // remember [position, end()) refers to the old storage
 899         new_end = m_storage.uninitialized_copy(position, end(), new_end);
 900       } // end try
 901       catch(...)
 902       {
 903         // something went wrong, so destroy & deallocate the new storage
 904         m_storage.destroy(new_storage.begin(), new_end);
 905         new_storage.deallocate();
 906
 907         // rethrow
 908         throw;
 909       } // end catch
 910
 911       // call destructors on the elements in the old storage
 912       m_storage.destroy(begin(), end());
 913
 914       // record the vector's new state
 915       m_storage.swap(new_storage);
 916       m_size    = old_size + n;
 917     } // end else
 918   } // end if
 919 } // end vector_base::fill_insert()
 920
 921 template<typename T, typename Alloc>
 922   template<typename InputIterator>
 923     void vector_base<T,Alloc>
 924       ::range_assign(InputIterator first,
 925                      InputIterator last)
 926 {
 927   // dispatch on traversal
 928   range_assign(first, last,
 929     typename thrust::iterator_traversal<InputIterator>::type());
 930 } // end range_assign()
 931
 932 template<typename T, typename Alloc>
 933   template<typename InputIterator>
 934     void vector_base<T,Alloc>
 935       ::range_assign(InputIterator first,
 936                      InputIterator last,
 937                      thrust::incrementable_traversal_tag)
 938 {
 939   iterator current(begin());
 940
 941   // assign to elements which already exist
 942   for(; first != last && current != end(); ++current, ++first)
 943   {
 944     *current = *first;
 945   } // end for
 946
 947   // either just the input was exhausted or both
 948   // the input and vector elements were exhausted
 949   if(first == last)
 950   {
 951     // if we exhausted the input, erase leftover elements
 952     erase(current, end());
 953   } // end if
 954   else
 955   {
 956     // insert the rest of the input at the end of the vector
 957     insert(end(), first, last);
 958   } // end else
 959 } // end vector_base::range_assign()
 960
 961 template<typename T, typename Alloc>
 962   template<typename RandomAccessIterator>
 963     void vector_base<T,Alloc>
 964       ::range_assign(RandomAccessIterator first,
 965                      RandomAccessIterator last,
 966                      thrust::random_access_traversal_tag)
 967 {
 968   const size_type n = thrust::distance(first, last);
 969
 970   if(n > capacity())
 971   {
 972     storage_type new_storage;
 973     allocate_and_copy(n, first, last, new_storage);
 974
 975     // call destructors on the elements in the old storage
 976     m_storage.destroy(begin(), end());
 977
 978     // record the vector's new state
 979     m_storage.swap(new_storage);
 980     m_size = n;
 981   } // end if
 982   else if(size() >= n)
 983   {
 984     // we can already accomodate the new range
 985     iterator new_end = thrust::copy(first, last, begin());
 986
 987     // destroy the elements we don't need
 988     m_storage.destroy(new_end, end());
 989
 990     // update size
 991     m_size = n;
 992   } // end else if
 993   else
 994   {
 995     // range fits inside allocated storage, but some elements
 996     // have not been constructed yet
 997
 998     // XXX TODO we could possibly implement this with one call
 999     // to transform rather than copy + uninitialized_copy
1000
1001     // copy to elements which already exist
1002     RandomAccessIterator mid = first;
1003     thrust::advance(mid, size());
1004     thrust::copy(first, mid, begin());
1005
1006     // uninitialize_copy to elements which must be constructed
1007     m_storage.uninitialized_copy(mid, last, end());
1008
1009     // update size
1010     m_size = n;
1011   } // end else
1012 } // end vector_base::assign()
1013
1014 template<typename T, typename Alloc>
1015   void vector_base<T,Alloc>
1016     ::fill_assign(size_type n, const T &x)
1017 {
1018   if(n > capacity())
1019   {
1020     // XXX we should also include a copy of the allocator:
1021     // vector_base<T,Alloc> temp(n, x, get_allocator());
1022     vector_base<T,Alloc> temp(n, x);
1023     temp.swap(*this);
1024   } // end if
1025   else if(n > size())
1026   {
1027     // fill to existing elements
1028     thrust::fill(begin(), end(), x);
1029
1030     // construct uninitialized elements
1031     m_storage.uninitialized_fill_n(end(), n - size(), x);
1032
1033     // adjust size
1034     m_size += (n - size());
1035   } // end else if
1036   else
1037   {
1038     // fill to existing elements
1039     iterator new_end = thrust::fill_n(begin(), n, x);
1040
1041     // erase the elements after the fill
1042     erase(new_end, end());
1043   } // end else
1044 } // end vector_base::fill_assign()
1045
1046 template<typename T, typename Alloc>
1047   template<typename ForwardIterator>
1048     void vector_base<T,Alloc>
1049       ::allocate_and_copy(size_type requested_size,
1050                           ForwardIterator first, ForwardIterator last,
1051                           storage_type &new_storage)
1052 {
1053   if(requested_size == 0)
1054   {
1055     new_storage.deallocate();
1056     return;
1057   } // end if
1058
1059   // allocate exponentially larger new storage
1060   size_type allocated_size = thrust::max<size_type>(requested_size, 2 * capacity());
1061
1062   // do not exceed maximum storage
1063   allocated_size = thrust::min<size_type>(allocated_size, max_size());
1064
1065   if(requested_size > allocated_size)
1066   {
1067     throw std::length_error("assignment exceeds max_size().");
1068   } // end if
1069
1070   new_storage.allocate(allocated_size);
1071
1072   try
1073   {
1074     // construct the range to the newly allocated storage
1075     m_storage.uninitialized_copy(first, last, new_storage.begin());
1076   } // end try
1077   catch(...)
1078   {
1079     // something went wrong, so destroy & deallocate the new storage
1080     // XXX seems like this destroys too many elements -- should just be last - first instead of requested_size
1081     iterator new_storage_end = new_storage.begin();
1082     thrust::advance(new_storage_end, requested_size);
1083     m_storage.destroy(new_storage.begin(), new_storage_end);
1084     new_storage.deallocate();
1085
1086     // rethrow
1087     throw;
1088   } // end catch
1089 } // end vector_base::allocate_and_copy()
1090
1091
1092 } // end detail
1093
1094 template<typename T, typename Alloc>
1095   void swap(detail::vector_base<T,Alloc> &a,
1096             detail::vector_base<T,Alloc> &b)
1097 {
1098   a.swap(b);
1099 } // end swap()
1100
1101
1102
1103 namespace detail
1104 {
1105
1106 // iterator tags match
1107 template <typename InputIterator1, typename InputIterator2>
1108 bool vector_equal(InputIterator1 first1, InputIterator1 last1,
1109                   InputIterator2 first2,
1110                   thrust::detail::true_type)
1111 {
1112   return thrust::equal(first1, last1, first2);
1113 }
1114
1115 // iterator tags differ
1116 template <typename InputIterator1, typename InputIterator2>
1117 bool vector_equal(InputIterator1 first1, InputIterator1 last1,
1118                   InputIterator2 first2,
1119                   thrust::detail::false_type)
1120 {
1121   typename thrust::iterator_difference<InputIterator1>::type n = thrust::distance(first1,last1);
1122
1123   typedef typename thrust::iterator_system<InputIterator1>::type FromSystem1;
1124   typedef typename thrust::iterator_system<InputIterator2>::type FromSystem2;
1125
1126   // bring both ranges to the host system
1127   // note that these copies are no-ops if the range is already convertible to the host system
1128   FromSystem1 from_system1;
1129   FromSystem2 from_system2;
1130   thrust::host_system_tag to_system;
1131   thrust::detail::move_to_system<InputIterator1, FromSystem1, thrust::host_system_tag> rng1(from_system1, to_system, first1, last1);
1132   thrust::detail::move_to_system<InputIterator2, FromSystem2, thrust::host_system_tag> rng2(from_system2, to_system, first2, first2 + n);
1133
1134   return thrust::equal(rng1.begin(), rng1.end(), rng2.begin());
1135 }
1136
1137 template <typename InputIterator1, typename InputIterator2>
1138 bool vector_equal(InputIterator1 first1, InputIterator1 last1,
1139                   InputIterator2 first2)
1140 {
1141   typedef typename thrust::iterator_system<InputIterator1>::type system1;
1142   typedef typename thrust::iterator_system<InputIterator2>::type system2;
1143
1144   // dispatch on the sameness of the two systems
1145   return vector_equal(first1, last1, first2,
1146     thrust::detail::is_same<system1,system2>());
1147 }
1148
1149 } // end namespace detail
1150
1151
1152
1153
1154 template<typename T1, typename Alloc1,
1155          typename T2, typename Alloc2>
1156 bool operator==(const detail::vector_base<T1,Alloc1>& lhs,
1157                 const detail::vector_base<T2,Alloc2>& rhs)
1158 {
1159     return lhs.size() == rhs.size() && detail::vector_equal(lhs.begin(), lhs.end(), rhs.begin());
1160 }
1161
1162 template<typename T1, typename Alloc1,
1163          typename T2, typename Alloc2>
1164 bool operator==(const detail::vector_base<T1,Alloc1>& lhs,
1165                 const std::vector<T2,Alloc2>&         rhs)
1166 {
1167     return lhs.size() == rhs.size() && detail::vector_equal(lhs.begin(), lhs.end(), rhs.begin());
1168 }
1169
1170 template<typename T1, typename Alloc1,
1171          typename T2, typename Alloc2>
1172 bool operator==(const std::vector<T1,Alloc1>&         lhs,
1173                 const detail::vector_base<T2,Alloc2>& rhs)
1174 {
1175     return lhs.size() == rhs.size() && detail::vector_equal(lhs.begin(), lhs.end(), rhs.begin());
1176 }
1177
1178 template<typename T1, typename Alloc1,
1179          typename T2, typename Alloc2>
1180 bool operator!=(const detail::vector_base<T1,Alloc1>& lhs,
1181                 const detail::vector_base<T2,Alloc2>& rhs)
1182 {
1183     return !(lhs == rhs);
1184 }
1185
1186 template<typename T1, typename Alloc1,
1187          typename T2, typename Alloc2>
1188 bool operator!=(const detail::vector_base<T1,Alloc1>& lhs,
1189                 const std::vector<T2,Alloc2>&         rhs)
1190 {
1191     return !(lhs == rhs);
1192 }
1193
1194 template<typename T1, typename Alloc1,
1195          typename T2, typename Alloc2>
1196 bool operator!=(const std::vector<T1,Alloc1>&         lhs,
1197                 const detail::vector_base<T2,Alloc2>& rhs)
1198 {
1199     return !(lhs == rhs);
1200 }
1201
1202 } // end thrust
1203