NVRTC

NVIDIA · Jan 22, 2025 · a270cf2 · a270cf2
1 parent 75fe5db
commit a270cf2
Show file tree

Hide file tree

Showing 5 changed files with 51 additions and 27 deletions.
diff --git a/thrust/thrust/detail/numeric_traits.h b/thrust/thrust/detail/numeric_traits.h
@@ -27,10 +27,9 @@
 #endif // no system header
 #include <thrust/detail/type_traits.h>
 
+#include <cuda/std/limits>
 #include <cuda/std/type_traits>
 
-#include <limits>
-
 // #include <stdint.h> // for intmax_t (not provided on MSVS 2005)
 
 THRUST_NAMESPACE_BEGIN
@@ -42,7 +41,7 @@ namespace detail
 using intmax_t = long long;
 
 template <typename Number>
-struct is_signed : integral_constant<bool, std::numeric_limits<Number>::is_signed>
+struct is_signed : integral_constant<bool, ::cuda::std::numeric_limits<Number>::is_signed>
 {}; // end is_signed
 
 template <typename T>

diff --git a/thrust/thrust/iterator/detail/iterator_facade_category.h b/thrust/thrust/iterator/detail/iterator_facade_category.h
@@ -94,6 +94,7 @@ struct iterator_facade_default_category;
 // Instead, it simply assumes that if is_convertible<Traversal, single_pass_traversal_tag>,
 // then the category is input_iterator_tag
 
+#if !_CCCL_COMPILER(NVRTC)
 // this is the function for standard system iterators
 template <typename Traversal, typename ValueParam, typename Reference>
 struct iterator_facade_default_category_std
@@ -110,6 +111,7 @@ struct iterator_facade_default_category_std
           thrust::detail::identity_<std::input_iterator_tag>,
           thrust::detail::identity_<Traversal>>>
 {}; // end iterator_facade_default_category_std
+#endif // !_CCCL_COMPILER(NVRTC)
 
 // this is the function for host system iterators
 template <typename Traversal, typename ValueParam, typename Reference>
@@ -127,7 +129,6 @@ struct iterator_facade_default_category_host
           thrust::detail::identity_<thrust::input_host_iterator_tag>,
           thrust::detail::identity_<Traversal>>>
 {}; // end iterator_facade_default_category_host
-
 // this is the function for device system iterators
 template <typename Traversal, typename ValueParam, typename Reference>
 struct iterator_facade_default_category_device
@@ -146,6 +147,7 @@ struct iterator_facade_default_category_device
           thrust::detail::identity_<Traversal>>>
 {}; // end iterator_facade_default_category_device
 
+#if !_CCCL_COMPILER(NVRTC)
 // this is the function for any system iterators
 template <typename Traversal, typename ValueParam, typename Reference>
 struct iterator_facade_default_category_any
@@ -155,29 +157,41 @@ struct iterator_facade_default_category_any
     thrust::any_system_tag,
     Traversal>;
 }; // end iterator_facade_default_category_any
+#endif // !_CCCL_COMPILER(NVRTC)
 
 template <typename System, typename Traversal, typename ValueParam, typename Reference>
 struct iterator_facade_default_category
+    :
+#if !_CCCL_COMPILER(NVRTC)
     // check for any system
-    : thrust::detail::eval_if<
-        ::cuda::std::is_convertible<System, thrust::any_system_tag>::value,
-        iterator_facade_default_category_any<Traversal, ValueParam, Reference>,
-
-        // check for host system
-        thrust::detail::eval_if<
-          ::cuda::std::is_convertible<System, thrust::host_system_tag>::value,
-          iterator_facade_default_category_host<Traversal, ValueParam, Reference>,
-
-          // check for device system
-          thrust::detail::eval_if<::cuda::std::is_convertible<System, thrust::device_system_tag>::value,
-                                  iterator_facade_default_category_device<Traversal, ValueParam, Reference>,
-
-                                  // if we don't recognize the system, get a standard iterator category
-                                  // and combine it with System & Traversal
-                                  thrust::detail::identity_<thrust::detail::iterator_category_with_system_and_traversal<
-                                    typename iterator_facade_default_category_std<Traversal, ValueParam, Reference>::type,
-                                    System,
-                                    Traversal>>>>>
+    thrust::detail::eval_if<
+      ::cuda::std::is_convertible<System, thrust::any_system_tag>::value,
+      iterator_facade_default_category_any<Traversal, ValueParam, Reference>,
+#endif // !_CCCL_COMPILER(NVRTC)
+       // check for host system
+      thrust::detail::eval_if<
+        ::cuda::std::is_convertible<System, thrust::host_system_tag>::value,
+        iterator_facade_default_category_host<Traversal, ValueParam, Reference>,
+
+        // check for device system
+        thrust::detail::eval_if<::cuda::std::is_convertible<System, thrust::device_system_tag>::value,
+                                iterator_facade_default_category_device<Traversal, ValueParam, Reference>,
+
+                                // if we don't recognize the system, get a standard iterator category
+                                // and combine it with System & Traversal (for NVRTC fall back to host categories)
+                                thrust::detail::identity_<thrust::detail::iterator_category_with_system_and_traversal<
+                                  typename
+#if _CCCL_COMPILER(NVRTC)
+                                  iterator_facade_default_category_host
+#else // _CCCL_COMPILER(NVRTC)
+                                iterator_facade_default_category_std
+#endif // _CCCL_COMPILER(NVRTC)
+                                  <Traversal, ValueParam, Reference>::type,
+                                  System,
+                                  Traversal>>>>
+#if !_CCCL_COMPILER(NVRTC)
+      >
+#endif // !_CCCL_COMPILER(NVRTC)
 {};
 
 template <typename System, typename Traversal, typename ValueParam, typename Reference>

diff --git a/thrust/thrust/iterator/iterator_adaptor.h b/thrust/thrust/iterator/iterator_adaptor.h
@@ -55,6 +55,8 @@ THRUST_NAMESPACE_BEGIN
  *  \{
  */
 
+class iterator_core_access;
+
 /*! \p iterator_adaptor is an iterator which adapts an existing type of iterator to create a new type of
  *  iterator. Most of Thrust's fancy iterators are defined via inheritance from \p iterator_adaptor.
  *  While composition of these existing Thrust iterators is often sufficient for expressing the desired
@@ -132,7 +134,7 @@ class _CCCL_DECLSPEC_EMPTY_BASES iterator_adaptor
   /*! \cond
    */
 
-  friend class thrust::iterator_core_access;
+  friend class iterator_core_access;
 
 protected:
   using super_t =

diff --git a/thrust/thrust/iterator/iterator_facade.h b/thrust/thrust/iterator/iterator_facade.h
@@ -306,7 +306,7 @@ template <typename Derived,
           typename System,
           typename Traversal,
           typename Reference,
-          typename Difference = std::ptrdiff_t>
+          typename Difference = ::cuda::std::ptrdiff_t>
 class iterator_facade
 {
 private:

diff --git a/thrust/thrust/iterator/iterator_traits.h b/thrust/thrust/iterator/iterator_traits.h
@@ -39,15 +39,24 @@
 #  pragma system_header
 #endif // no system header
 
-#include <iterator>
+#if _CCCL_COMPILER(NVRTC)
+#  include <cuda/std/iterator>
+#else // _CCCL_COMPILER(NVRTC)
+#  include <iterator>
+#endif // _CCCL_COMPILER(NVRTC)
 
 THRUST_NAMESPACE_BEGIN
 
 /*! \p iterator_traits is a type trait class that provides a uniform
  *  interface for querying the properties of iterators at compile-time.
  */
 template <typename T>
-struct iterator_traits : std::iterator_traits<T>
+struct iterator_traits
+    :
+#if _CCCL_COMPILER(NVRTC)
+    ::cuda::
+#endif // _CCCL_COMPILER(NVRTC)
+      std::iterator_traits<T>
 {};
 
 template <typename Iterator>