Avoid C-style casts to convert iterators

miscco · miscco · commit 12d640f55aa3 · 2024-08-23T10:29:21.000+02:00
diff --git a/c/include/cccl/types.h b/c/include/cccl/types.h
@@ -45,7 +45,7 @@ struct cccl_type_info
 enum class cccl_op_kind_t
 {
   stateless = 0,
-  stateful = 1
+  stateful  = 1
 };
 
 struct cccl_op_t
@@ -61,7 +61,7 @@ struct cccl_op_t
 
 enum class cccl_iterator_kind_t
 {
-  pointer = 0,
+  pointer  = 0,
   iterator = 1
 };
 
diff --git a/libcudacxx/include/cuda/std/inplace_vector b/libcudacxx/include/cuda/std/inplace_vector
@@ -1047,7 +1047,7 @@ public:
   // [containers.sequences.inplace.vector.access], element access
   _CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY constexpr reference at(const size_type __pos)
   {
-    if (this->size() < __pos)
+    if (__pos > this->size())
     {
       _CUDA_VSTD::__throw_out_of_range("inplace_vector::at");
     }
@@ -1056,7 +1056,7 @@ public:
 
   _CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reference at(const size_type __pos) const
   {
-    if (this->size() < __pos)
+    if (__pos > this->size())
     {
       _CUDA_VSTD::__throw_out_of_range("inplace_vector::at");
     }
@@ -1168,43 +1168,44 @@ public:
   _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator
   insert(const_iterator __cpos, const size_type __count, const _Tp& __value)
   {
-    const iterator __pos = (iterator) __cpos;
-    const iterator __end = this->end();
-    if (this->size() + __count > _Capacity)
+    const auto __pos = static_cast<size_type>(__cpos - this->cbegin());
+    if (__count > _Capacity - this->size())
     {
       _CUDA_VSTD::__throw_bad_alloc();
     }
-    else if (__pos < this->begin() || __end < __pos)
+    else if (__pos > this->size())
     {
       _CUDA_VSTD::__throw_out_of_range("inplace_vector::insert(const_iterator, size_type, T)");
     }
 
+    const iterator __first = this->begin() + __pos;
     if (__count == 0)
     {
-      return __pos;
+      return __first;
     }
 
-    if (__pos == __end)
+    const iterator __end = this->end();
+    if (__pos == this->size())
     {
       this->__uninitialized_fill(__end, __end + __count, __value);
-      return __pos;
+      return __first;
     }
 
-    const iterator __middle = __pos + __count;
+    const iterator __middle = __first + __count;
     if (__end <= __middle)
     { // all existing elements are pushed into uninitialized storage
       this->__uninitialized_fill(__end, __middle, __value);
-      this->__uninitialized_move(__pos, __end, __middle);
-      _CUDA_VSTD::fill(__pos, __end, __value);
+      this->__uninitialized_move(__first, __end, __middle);
+      _CUDA_VSTD::fill(__first, __end, __value);
     }
     else
     { // some elements get copied into existing storage
       this->__uninitialized_move(__end - __count, __end, __end);
-      _CUDA_VSTD::move_backward(__pos, __end - __count, __end);
-      _CUDA_VSTD::fill(__pos, __middle, __value);
+      _CUDA_VSTD::move_backward(__first, __end - __count, __end);
+      _CUDA_VSTD::fill(__first, __middle, __value);
     }
 
-    return __pos;
+    return __first;
   }
 
   _LIBCUDACXX_TEMPLATE(class _Iter)
@@ -1219,99 +1220,99 @@ public:
       emplace_back(*__first);
     }
 
-    const iterator __pos = (iterator) __cpos;
-    _CUDA_VSTD::rotate(__pos, __old_end, this->end());
-    return __pos;
+    const iterator __res = this->begin() + (__cpos - this->cbegin());
+    _CUDA_VSTD::rotate(__res, __old_end, this->end());
+    return __res;
   }
 
   _LIBCUDACXX_TEMPLATE(class _Iter)
   _LIBCUDACXX_REQUIRES(__is_cpp17_forward_iterator<_Iter>::value)
   _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator insert(const_iterator __cpos, _Iter __first, _Iter __last)
   {
-    const iterator __pos = (iterator) __cpos;
-    const iterator __end = this->end();
-    const auto __count   = _CUDA_VSTD::distance(__first, __last);
-    if (this->size() + __count > _Capacity)
+    const auto __pos   = static_cast<size_type>(__cpos - this->cbegin());
+    const auto __count = static_cast<size_type>(_CUDA_VSTD::distance(__first, __last));
+    if (__count > _Capacity - this->size())
     {
       _CUDA_VSTD::__throw_bad_alloc();
     }
-    else if (__pos < this->begin() || __end < __pos)
+    else if (__pos > this->size())
     {
       _CUDA_VSTD::__throw_out_of_range("inplace_vector::insert(const_iterator, Iter, Iter)");
     }
 
+    const iterator __res = this->begin() + __pos;
     if (__count == 0)
     {
-      return __pos;
+      return __res;
     }
 
-    if (__pos == __end)
+    const iterator __end = this->end();
+    if (__pos == this->size())
     {
       this->__uninitialized_copy(__first, __last, __end);
-      return __pos;
+      return __res;
     }
 
-    const iterator __middle = __pos + __count;
-    const auto __to_copy    = __end - __pos;
+    const iterator __middle = __res + __count;
     if (__end <= __middle)
     { // all existing elements are pushed into uninitialized storage
-      _Iter __imiddle = _CUDA_VSTD::next(__first, __to_copy);
+      _Iter __imiddle = _CUDA_VSTD::next(__first, this->size() - __pos);
       this->__uninitialized_copy(__imiddle, __last, __end);
-      this->__uninitialized_move(__pos, __end, __middle);
-      _CUDA_VSTD::copy(__first, __imiddle, __pos);
+      this->__uninitialized_move(__res, __end, __middle);
+      _CUDA_VSTD::copy(__first, __imiddle, __res);
     }
     else
     { // all new elements get copied into existing storage
       this->__uninitialized_move(__end - __count, __end, __end);
-      _CUDA_VSTD::move_backward(__pos, __end - __count, __end);
-      _CUDA_VSTD::copy(__first, __last, __pos);
+      _CUDA_VSTD::move_backward(__res, __end - __count, __end);
+      _CUDA_VSTD::copy(__first, __last, __res);
     }
 
-    return __pos;
+    return __res;
   }
 
   _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator insert(const_iterator __cpos, initializer_list<_Tp> __ilist)
   {
-    const iterator __pos = (iterator) __cpos;
-    const iterator __end = this->end();
-    const auto __count   = __ilist.size();
-    if (this->size() + __count > _Capacity)
+    const auto __pos   = static_cast<size_type>(__cpos - this->cbegin());
+    const auto __count = __ilist.size();
+    if (__count > _Capacity - this->size())
     {
       _CUDA_VSTD::__throw_bad_alloc();
     }
-    else if (__pos < this->begin() || __end < __pos)
+    else if (__pos > this->size())
     {
       _CUDA_VSTD::__throw_out_of_range("inplace_vector::insert(const_iterator, initializer_list)");
     }
 
+    const iterator __res = this->begin() + __pos;
     if (__count == 0)
     {
-      return __pos;
+      return __res;
     }
 
-    if (__pos == __end)
+    const iterator __end = this->end();
+    if (__pos == this->size())
     {
       this->__uninitialized_copy(__ilist.begin(), __ilist.end(), __end);
-      return __pos;
+      return __res;
     }
 
-    const iterator __middle = __pos + __count;
-    const auto __to_copy    = __end - __pos;
+    const iterator __middle = __res + __count;
     if (__end <= __middle)
     { // all existing elements are pushed into uninitialized storage
-      auto __imiddel = __ilist.begin() + __to_copy;
+      auto __imiddel = __ilist.begin() + this->size() - __pos;
       this->__uninitialized_copy(__imiddel, __ilist.end(), __end);
-      this->__uninitialized_move(__pos, __end, __middle);
-      _CUDA_VSTD::copy(__ilist.begin(), __imiddel, __pos);
+      this->__uninitialized_move(__res, __end, __middle);
+      _CUDA_VSTD::copy(__ilist.begin(), __imiddel, __res);
     }
     else
     { // all new elements get copied into existing storage
       this->__uninitialized_move(__end - __count, __end, __end);
-      _CUDA_VSTD::move_backward(__pos, __end - __count, __end);
-      _CUDA_VSTD::copy(__ilist.begin(), __ilist.end(), __pos);
+      _CUDA_VSTD::move_backward(__res, __end - __count, __end);
+      _CUDA_VSTD::copy(__ilist.begin(), __ilist.end(), __res);
     }
 
-    return __pos;
+    return __res;
   }
 
 #  if _CCCL_STD_VER >= 2017 && !defined(_CCCL_COMPILER_MSVC_2017)
@@ -1329,7 +1330,7 @@ public:
       emplace_back(*__first);
     }
 
-    const iterator __pos = (iterator) __cpos;
+    const auto __pos = this->begin() + static_cast<size_type>(__cpos - this->cbegin());
     _CUDA_VSTD::rotate(__pos, __old_end, this->end());
     return __pos;
   }
@@ -1369,30 +1370,30 @@ public:
   template <class... _Args>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator emplace(const_iterator __cpos, _Args&&... __args)
   {
-    const iterator __pos = (iterator) __cpos;
-    const iterator __end = this->end();
+    const auto __pos = static_cast<size_type>(__cpos - this->cbegin());
     if (this->size() == _Capacity)
     {
       _CUDA_VSTD::__throw_bad_alloc();
     }
-    else if (__pos < this->begin() || __end < __pos)
+    else if (__pos > this->size())
     {
       _CUDA_VSTD::__throw_out_of_range("inplace_vector::emplace(const_iterator, Args...)");
     }
 
-    if (__pos == __end)
+    const iterator __res = this->begin() + __pos;
+    if (__pos == this->size())
     {
       this->unchecked_emplace_back(_CUDA_VSTD::forward<_Args>(__args)...);
-    }
-    else
-    {
-      _Tp __temp{_CUDA_VSTD::forward<_Args>(__args)...};
-      this->unchecked_emplace_back(_CUDA_VSTD::move(*(__end - 1)));
-      _CUDA_VSTD::move_backward(__pos, __end - 1, __end);
-      *__pos = _CUDA_VSTD::move(__temp);
+      return __res;
     }
 
-    return __pos;
+    const iterator __end = this->end();
+    _Tp __temp{_CUDA_VSTD::forward<_Args>(__args)...};
+    this->unchecked_emplace_back(_CUDA_VSTD::move(*(__end - 1)));
+    _CUDA_VSTD::move_backward(__res, __end - 1, __end);
+    *__res = _CUDA_VSTD::move(__temp);
+
+    return __res;
   }
 
   template <class... _Args>
@@ -1532,21 +1533,22 @@ public:
   _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator
   erase(const_iterator __cpos) noexcept(_CCCL_TRAIT(is_nothrow_move_assignable, _Tp))
   {
-    const iterator __pos = (iterator) __cpos;
-    const iterator __end = this->end();
-    if (__pos == __end)
+    const auto __pos = static_cast<size_type>(__cpos - this->cbegin());
+    if (__pos > this->size())
     {
-      return __pos;
+      _CUDA_VSTD_NOVERSION::terminate();
     }
 
-    if (this->size() == 0 || __pos < this->begin() || __end < __pos)
+    const iterator __res = this->begin() + __pos;
+    if (__pos == this->size())
     {
-      _CUDA_VSTD_NOVERSION::terminate();
+      return __res;
     }
 
-    _CUDA_VSTD::move(__pos + 1, __end, __pos);
+    const iterator __end = this->end();
+    _CUDA_VSTD::move(__res + 1, __end, __res);
     this->__destroy(__end - 1, __end);
-    return __pos;
+    return __res;
   }
 
   _LIBCUDACXX_INLINE_VISIBILITY constexpr iterator
diff --git a/thrust/thrust/system/cuda/detail/scan.h b/thrust/thrust/system/cuda/detail/scan.h
@@ -71,7 +71,8 @@ _CCCL_HOST_DEVICE OutputIt inclusive_scan_n_impl(
   cudaError_t status;
 
   // Negative number of items are normalized to `0`
-  if(thrust::detail::is_negative(num_items)){
+  if (thrust::detail::is_negative(num_items))
+  {
     num_items = 0;
   }
 
@@ -126,7 +127,8 @@ _CCCL_HOST_DEVICE OutputIt exclusive_scan_n_impl(
   cudaError_t status;
 
   // Negative number of items are normalized to `0`
-  if(thrust::detail::is_negative(num_items)){
+  if (thrust::detail::is_negative(num_items))
+  {
     num_items = 0;
   }
 

Original file line number	Diff line number	Diff line change
`@@ -45,7 +45,7 @@ struct cccl_type_info`
`45`	`45`	`enum class cccl_op_kind_t`
`46`	`46`	`{`
`47`	`47`	`stateless = 0,`
`48`		`- stateful = 1`
	`48`	`+ stateful = 1`
`49`	`49`	`};`
`50`	`50`
`51`	`51`	`struct cccl_op_t`
`@@ -61,7 +61,7 @@ struct cccl_op_t`
`61`	`61`
`62`	`62`	`enum class cccl_iterator_kind_t`
`63`	`63`	`{`
`64`		`- pointer = 0,`
	`64`	`+ pointer = 0,`
`65`	`65`	`iterator = 1`
`66`	`66`	`};`
`67`	`67`
Original file line number	Diff line number	Diff line change
`@@ -1047,7 +1047,7 @@ public:`
`1047`	`1047`	`// [containers.sequences.inplace.vector.access], element access`
`1048`	`1048`	`_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY constexpr reference at(const size_type __pos)`
`1049`	`1049`	`{`
`1050`		`- if (this->size() < __pos)`
	`1050`	`+ if (__pos > this->size())`
`1051`	`1051`	`{`
`1052`	`1052`	`_CUDA_VSTD::__throw_out_of_range("inplace_vector::at");`
`1053`	`1053`	`}`
`@@ -1056,7 +1056,7 @@ public:`
`1056`	`1056`
`1057`	`1057`	`_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY constexpr const_reference at(const size_type __pos) const`
`1058`	`1058`	`{`
`1059`		`- if (this->size() < __pos)`
	`1059`	`+ if (__pos > this->size())`
`1060`	`1060`	`{`
`1061`	`1061`	`_CUDA_VSTD::__throw_out_of_range("inplace_vector::at");`
`1062`	`1062`	`}`
`@@ -1168,43 +1168,44 @@ public:`
`1168`	`1168`	`_LIBCUDACXX_INLINE_VISIBILITY constexpr iterator`
`1169`	`1169`	`insert(const_iterator __cpos, const size_type __count, const _Tp& __value)`
`1170`	`1170`	`{`
`1171`		`- const iterator __pos = (iterator) __cpos;`
`1172`		`- const iterator __end = this->end();`
`1173`		`- if (this->size() + __count > _Capacity)`
	`1171`	`+ const auto __pos = static_cast<size_type>(__cpos - this->cbegin());`
	`1172`	`+ if (__count > _Capacity - this->size())`
`1174`	`1173`	`{`
`1175`	`1174`	`_CUDA_VSTD::__throw_bad_alloc();`
`1176`	`1175`	`}`
`1177`		`- else if (__pos < this->begin() \|\| __end < __pos)`
	`1176`	`+ else if (__pos > this->size())`
`1178`	`1177`	`{`
`1179`	`1178`	`_CUDA_VSTD::__throw_out_of_range("inplace_vector::insert(const_iterator, size_type, T)");`
`1180`	`1179`	`}`
`1181`	`1180`
	`1181`	`+ const iterator __first = this->begin() + __pos;`
`1182`	`1182`	`if (__count == 0)`
`1183`	`1183`	`{`
`1184`		`- return __pos;`
	`1184`	`+ return __first;`
`1185`	`1185`	`}`
`1186`	`1186`
`1187`		`- if (__pos == __end)`
	`1187`	`+ const iterator __end = this->end();`
	`1188`	`+ if (__pos == this->size())`
`1188`	`1189`	`{`
`1189`	`1190`	`this->__uninitialized_fill(__end, __end + __count, __value);`
`1190`		`- return __pos;`
	`1191`	`+ return __first;`
`1191`	`1192`	`}`
`1192`	`1193`
`1193`		`- const iterator __middle = __pos + __count;`
	`1194`	`+ const iterator __middle = __first + __count;`
`1194`	`1195`	`if (__end <= __middle)`
`1195`	`1196`	`{ // all existing elements are pushed into uninitialized storage`
`1196`	`1197`	`this->__uninitialized_fill(__end, __middle, __value);`
`1197`		`- this->__uninitialized_move(__pos, __end, __middle);`
`1198`		`- _CUDA_VSTD::fill(__pos, __end, __value);`
	`1198`	`+ this->__uninitialized_move(__first, __end, __middle);`
	`1199`	`+ _CUDA_VSTD::fill(__first, __end, __value);`
`1199`	`1200`	`}`
`1200`	`1201`	`else`
`1201`	`1202`	`{ // some elements get copied into existing storage`
`1202`	`1203`	`this->__uninitialized_move(__end - __count, __end, __end);`
`1203`		`- _CUDA_VSTD::move_backward(__pos, __end - __count, __end);`
`1204`		`- _CUDA_VSTD::fill(__pos, __middle, __value);`
	`1204`	`+ _CUDA_VSTD::move_backward(__first, __end - __count, __end);`
	`1205`	`+ _CUDA_VSTD::fill(__first, __middle, __value);`
`1205`	`1206`	`}`
`1206`	`1207`
`1207`		`- return __pos;`
	`1208`	`+ return __first;`
`1208`	`1209`	`}`
`1209`	`1210`
`1210`	`1211`	`_LIBCUDACXX_TEMPLATE(class _Iter)`
`@@ -1219,99 +1220,99 @@ public:`
`1219`	`1220`	`emplace_back(*__first);`
`1220`	`1221`	`}`
`1221`	`1222`
`1222`		`- const iterator __pos = (iterator) __cpos;`
`1223`		`- _CUDA_VSTD::rotate(__pos, __old_end, this->end());`
`1224`		`- return __pos;`
	`1223`	`+ const iterator __res = this->begin() + (__cpos - this->cbegin());`
	`1224`	`+ _CUDA_VSTD::rotate(__res, __old_end, this->end());`
	`1225`	`+ return __res;`
`1225`	`1226`	`}`
`1226`	`1227`
`1227`	`1228`	`_LIBCUDACXX_TEMPLATE(class _Iter)`
`1228`	`1229`	`_LIBCUDACXX_REQUIRES(__is_cpp17_forward_iterator<_Iter>::value)`
`1229`	`1230`	`_LIBCUDACXX_INLINE_VISIBILITY constexpr iterator insert(const_iterator __cpos, _Iter __first, _Iter __last)`
`1230`	`1231`	`{`
`1231`		`- const iterator __pos = (iterator) __cpos;`
`1232`		`- const iterator __end = this->end();`
`1233`		`- const auto __count = _CUDA_VSTD::distance(__first, __last);`
`1234`		`- if (this->size() + __count > _Capacity)`
	`1232`	`+ const auto __pos = static_cast<size_type>(__cpos - this->cbegin());`
	`1233`	`+ const auto __count = static_cast<size_type>(_CUDA_VSTD::distance(__first, __last));`
	`1234`	`+ if (__count > _Capacity - this->size())`
`1235`	`1235`	`{`
`1236`	`1236`	`_CUDA_VSTD::__throw_bad_alloc();`
`1237`	`1237`	`}`
`1238`		`- else if (__pos < this->begin() \|\| __end < __pos)`
	`1238`	`+ else if (__pos > this->size())`
`1239`	`1239`	`{`
`1240`	`1240`	`_CUDA_VSTD::__throw_out_of_range("inplace_vector::insert(const_iterator, Iter, Iter)");`
`1241`	`1241`	`}`
`1242`	`1242`
	`1243`	`+ const iterator __res = this->begin() + __pos;`
`1243`	`1244`	`if (__count == 0)`
`1244`	`1245`	`{`
`1245`		`- return __pos;`
	`1246`	`+ return __res;`
`1246`	`1247`	`}`
`1247`	`1248`
`1248`		`- if (__pos == __end)`
	`1249`	`+ const iterator __end = this->end();`
	`1250`	`+ if (__pos == this->size())`
`1249`	`1251`	`{`
`1250`	`1252`	`this->__uninitialized_copy(__first, __last, __end);`
`1251`		`- return __pos;`
	`1253`	`+ return __res;`
`1252`	`1254`	`}`
`1253`	`1255`
`1254`		`- const iterator __middle = __pos + __count;`
`1255`		`- const auto __to_copy = __end - __pos;`
	`1256`	`+ const iterator __middle = __res + __count;`
`1256`	`1257`	`if (__end <= __middle)`
`1257`	`1258`	`{ // all existing elements are pushed into uninitialized storage`
`1258`		`- _Iter __imiddle = _CUDA_VSTD::next(__first, __to_copy);`
	`1259`	`+ _Iter __imiddle = _CUDA_VSTD::next(__first, this->size() - __pos);`
`1259`	`1260`	`this->__uninitialized_copy(__imiddle, __last, __end);`
`1260`		`- this->__uninitialized_move(__pos, __end, __middle);`
`1261`		`- _CUDA_VSTD::copy(__first, __imiddle, __pos);`
	`1261`	`+ this->__uninitialized_move(__res, __end, __middle);`
	`1262`	`+ _CUDA_VSTD::copy(__first, __imiddle, __res);`
`1262`	`1263`	`}`
`1263`	`1264`	`else`
`1264`	`1265`	`{ // all new elements get copied into existing storage`
`1265`	`1266`	`this->__uninitialized_move(__end - __count, __end, __end);`
`1266`		`- _CUDA_VSTD::move_backward(__pos, __end - __count, __end);`
`1267`		`- _CUDA_VSTD::copy(__first, __last, __pos);`
	`1267`	`+ _CUDA_VSTD::move_backward(__res, __end - __count, __end);`
	`1268`	`+ _CUDA_VSTD::copy(__first, __last, __res);`
`1268`	`1269`	`}`
`1269`	`1270`
`1270`		`- return __pos;`
	`1271`	`+ return __res;`
`1271`	`1272`	`}`
`1272`	`1273`
`1273`	`1274`	`_LIBCUDACXX_INLINE_VISIBILITY constexpr iterator insert(const_iterator __cpos, initializer_list<_Tp> __ilist)`
`1274`	`1275`	`{`
`1275`		`- const iterator __pos = (iterator) __cpos;`
`1276`		`- const iterator __end = this->end();`
`1277`		`- const auto __count = __ilist.size();`
`1278`		`- if (this->size() + __count > _Capacity)`
	`1276`	`+ const auto __pos = static_cast<size_type>(__cpos - this->cbegin());`
	`1277`	`+ const auto __count = __ilist.size();`
	`1278`	`+ if (__count > _Capacity - this->size())`
`1279`	`1279`	`{`
`1280`	`1280`	`_CUDA_VSTD::__throw_bad_alloc();`
`1281`	`1281`	`}`
`1282`		`- else if (__pos < this->begin() \|\| __end < __pos)`
	`1282`	`+ else if (__pos > this->size())`
`1283`	`1283`	`{`
`1284`	`1284`	`_CUDA_VSTD::__throw_out_of_range("inplace_vector::insert(const_iterator, initializer_list)");`
`1285`	`1285`	`}`
`1286`	`1286`
	`1287`	`+ const iterator __res = this->begin() + __pos;`
`1287`	`1288`	`if (__count == 0)`
`1288`	`1289`	`{`
`1289`		`- return __pos;`
	`1290`	`+ return __res;`
`1290`	`1291`	`}`
`1291`	`1292`
`1292`		`- if (__pos == __end)`
	`1293`	`+ const iterator __end = this->end();`
	`1294`	`+ if (__pos == this->size())`
`1293`	`1295`	`{`
`1294`	`1296`	`this->__uninitialized_copy(__ilist.begin(), __ilist.end(), __end);`
`1295`		`- return __pos;`
	`1297`	`+ return __res;`
`1296`	`1298`	`}`
`1297`	`1299`
`1298`		`- const iterator __middle = __pos + __count;`
`1299`		`- const auto __to_copy = __end - __pos;`
	`1300`	`+ const iterator __middle = __res + __count;`
`1300`	`1301`	`if (__end <= __middle)`
`1301`	`1302`	`{ // all existing elements are pushed into uninitialized storage`
`1302`		`- auto __imiddel = __ilist.begin() + __to_copy;`
	`1303`	`+ auto __imiddel = __ilist.begin() + this->size() - __pos;`
`1303`	`1304`	`this->__uninitialized_copy(__imiddel, __ilist.end(), __end);`
`1304`		`- this->__uninitialized_move(__pos, __end, __middle);`
`1305`		`- _CUDA_VSTD::copy(__ilist.begin(), __imiddel, __pos);`
	`1305`	`+ this->__uninitialized_move(__res, __end, __middle);`
	`1306`	`+ _CUDA_VSTD::copy(__ilist.begin(), __imiddel, __res);`
`1306`	`1307`	`}`
`1307`	`1308`	`else`
`1308`	`1309`	`{ // all new elements get copied into existing storage`
`1309`	`1310`	`this->__uninitialized_move(__end - __count, __end, __end);`
`1310`		`- _CUDA_VSTD::move_backward(__pos, __end - __count, __end);`
`1311`		`- _CUDA_VSTD::copy(__ilist.begin(), __ilist.end(), __pos);`
	`1311`	`+ _CUDA_VSTD::move_backward(__res, __end - __count, __end);`
	`1312`	`+ _CUDA_VSTD::copy(__ilist.begin(), __ilist.end(), __res);`
`1312`	`1313`	`}`
`1313`	`1314`
`1314`		`- return __pos;`
	`1315`	`+ return __res;`
`1315`	`1316`	`}`
`1316`	`1317`
`1317`	`1318`	`# if _CCCL_STD_VER >= 2017 && !defined(_CCCL_COMPILER_MSVC_2017)`
`@@ -1329,7 +1330,7 @@ public:`
`1329`	`1330`	`emplace_back(*__first);`
`1330`	`1331`	`}`
`1331`	`1332`
`1332`		`- const iterator __pos = (iterator) __cpos;`
	`1333`	`+ const auto __pos = this->begin() + static_cast<size_type>(__cpos - this->cbegin());`
`1333`	`1334`	`_CUDA_VSTD::rotate(__pos, __old_end, this->end());`
`1334`	`1335`	`return __pos;`
`1335`	`1336`	`}`
`@@ -1369,30 +1370,30 @@ public:`
`1369`	`1370`	`template <class... _Args>`
`1370`	`1371`	`_LIBCUDACXX_INLINE_VISIBILITY constexpr iterator emplace(const_iterator __cpos, _Args&&... __args)`
`1371`	`1372`	`{`
`1372`		`- const iterator __pos = (iterator) __cpos;`
`1373`		`- const iterator __end = this->end();`
	`1373`	`+ const auto __pos = static_cast<size_type>(__cpos - this->cbegin());`
`1374`	`1374`	`if (this->size() == _Capacity)`
`1375`	`1375`	`{`
`1376`	`1376`	`_CUDA_VSTD::__throw_bad_alloc();`
`1377`	`1377`	`}`
`1378`		`- else if (__pos < this->begin() \|\| __end < __pos)`
	`1378`	`+ else if (__pos > this->size())`
`1379`	`1379`	`{`
`1380`	`1380`	`_CUDA_VSTD::__throw_out_of_range("inplace_vector::emplace(const_iterator, Args...)");`
`1381`	`1381`	`}`
`1382`	`1382`
`1383`		`- if (__pos == __end)`
	`1383`	`+ const iterator __res = this->begin() + __pos;`
	`1384`	`+ if (__pos == this->size())`
`1384`	`1385`	`{`
`1385`	`1386`	`this->unchecked_emplace_back(_CUDA_VSTD::forward<_Args>(__args)...);`
`1386`		`- }`
`1387`		`- else`
`1388`		`- {`
`1389`		`- _Tp __temp{_CUDA_VSTD::forward<_Args>(__args)...};`
`1390`		`- this->unchecked_emplace_back(_CUDA_VSTD::move(*(__end - 1)));`
`1391`		`- _CUDA_VSTD::move_backward(__pos, __end - 1, __end);`
`1392`		`- *__pos = _CUDA_VSTD::move(__temp);`
	`1387`	`+ return __res;`
`1393`	`1388`	`}`
`1394`	`1389`
`1395`		`- return __pos;`
	`1390`	`+ const iterator __end = this->end();`
	`1391`	`+ _Tp __temp{_CUDA_VSTD::forward<_Args>(__args)...};`
	`1392`	`+ this->unchecked_emplace_back(_CUDA_VSTD::move(*(__end - 1)));`
	`1393`	`+ _CUDA_VSTD::move_backward(__res, __end - 1, __end);`
	`1394`	`+ *__res = _CUDA_VSTD::move(__temp);`
	`1395`	`+`
	`1396`	`+ return __res;`
`1396`	`1397`	`}`
`1397`	`1398`
`1398`	`1399`	`template <class... _Args>`
`@@ -1532,21 +1533,22 @@ public:`
`1532`	`1533`	`_LIBCUDACXX_INLINE_VISIBILITY constexpr iterator`
`1533`	`1534`	`erase(const_iterator __cpos) noexcept(_CCCL_TRAIT(is_nothrow_move_assignable, _Tp))`
`1534`	`1535`	`{`
`1535`		`- const iterator __pos = (iterator) __cpos;`
`1536`		`- const iterator __end = this->end();`
`1537`		`- if (__pos == __end)`
	`1536`	`+ const auto __pos = static_cast<size_type>(__cpos - this->cbegin());`
	`1537`	`+ if (__pos > this->size())`
`1538`	`1538`	`{`
`1539`		`- return __pos;`
	`1539`	`+ _CUDA_VSTD_NOVERSION::terminate();`
`1540`	`1540`	`}`
`1541`	`1541`
`1542`		`- if (this->size() == 0 \|\| __pos < this->begin() \|\| __end < __pos)`
	`1542`	`+ const iterator __res = this->begin() + __pos;`
	`1543`	`+ if (__pos == this->size())`
`1543`	`1544`	`{`
`1544`		`- _CUDA_VSTD_NOVERSION::terminate();`
	`1545`	`+ return __res;`
`1545`	`1546`	`}`
`1546`	`1547`
`1547`		`- _CUDA_VSTD::move(__pos + 1, __end, __pos);`
	`1548`	`+ const iterator __end = this->end();`
	`1549`	`+ _CUDA_VSTD::move(__res + 1, __end, __res);`
`1548`	`1550`	`this->__destroy(__end - 1, __end);`
`1549`		`- return __pos;`
	`1551`	`+ return __res;`
`1550`	`1552`	`}`
`1551`	`1553`
`1552`	`1554`	`_LIBCUDACXX_INLINE_VISIBILITY constexpr iterator`