remove __cccl_lib_mdspan

NVIDIA · Mar 6, 2025 · ba45344 · ba45344
1 parent fc8ade9
commit ba45344
Show file tree

Hide file tree

Showing 7 changed files with 68 additions and 104 deletions.
diff --git a/cub/cub/detail/mdspan_utils.cuh b/cub/cub/detail/mdspan_utils.cuh
@@ -37,15 +37,13 @@
 #  pragma system_header
 #endif // no system header
 
-#if __cccl_lib_mdspan
+#include <cub/detail/fast_modulo_division.cuh> // fast_div_mod
 
-#  include <cub/detail/fast_modulo_division.cuh> // fast_div_mod
-
-#  include <cuda/std/array> // cuda::std::array
-#  include <cuda/std/cstddef> // size_t
-#  include <cuda/std/mdspan>
-#  include <cuda/std/type_traits> // make_unsigned_t
-#  include <cuda/std/utility> // ::cuda::std::index_sequence
+#include <cuda/std/array> // cuda::std::array
+#include <cuda/std/cstddef> // size_t
+#include <cuda/std/mdspan>
+#include <cuda/std/type_traits> // make_unsigned_t
+#include <cuda/std/utility> // ::cuda::std::index_sequence
 
 CUB_NAMESPACE_BEGIN
 
@@ -118,5 +116,3 @@ _CCCL_NODISCARD _CCCL_HOST_DEVICE _CCCL_FORCEINLINE constexpr bool is_sub_size_s
 } // namespace detail
 
 CUB_NAMESPACE_END
-
-#endif // if __cccl_lib_mdspan
diff --git a/cub/cub/detail/type_traits.cuh b/cub/cub/detail/type_traits.cuh
@@ -48,11 +48,9 @@
 _CCCL_SUPPRESS_DEPRECATED_PUSH
 #include <cuda/std/functional>
 _CCCL_SUPPRESS_DEPRECATED_POP
-#include <cuda/std/array>
-#if __cccl_lib_mdspan
-#  include <cuda/std/mdspan>
-#endif // __cccl_lib_mdspan
 #include <cuda/std/__concepts/concept_macros.h> // IWYU pragma: keep
+#include <cuda/std/array>
+#include <cuda/std/mdspan>
 #include <cuda/std/span>
 #include <cuda/std/type_traits>
 
@@ -103,16 +101,12 @@ template <typename T, ::cuda::std::size_t N>
 struct is_fixed_size_random_access_range<::cuda::std::array<T, N>, void> : ::cuda::std::true_type
 {};
 
-#if __cccl_lib_mdspan
-
 template <typename T, typename E, typename L, typename A>
 struct is_fixed_size_random_access_range<
   ::cuda::std::mdspan<T, E, L, A>,
   ::cuda::std::enable_if_t<E::rank == 1 && E::static_extent(0) != ::cuda::std::dynamic_extent>> : ::cuda::std::true_type
 {};
 
-#endif // __cccl_lib_mdspan
-
 template <typename T>
 using is_fixed_size_random_access_range_t = typename is_fixed_size_random_access_range<T>::type;
 
@@ -143,16 +137,12 @@ template <typename T, ::cuda::std::size_t N>
 struct static_size<::cuda::std::span<T, N>, void> : ::cuda::std::integral_constant<int, N>
 {};
 
-#if __cccl_lib_mdspan
-
 template <typename T, typename E, typename L, typename A>
 struct static_size<::cuda::std::mdspan<T, E, L, A>,
                    ::cuda::std::enable_if_t<E::rank == 1 && E::static_extent(0) != ::cuda::std::dynamic_extent>>
     : ::cuda::std::integral_constant<int, E::static_extent(1)>
 {};
 
-#endif // __cccl_lib_mdspan
-
 template <typename T>
 _CCCL_NODISCARD _CCCL_HOST_DEVICE _CCCL_FORCEINLINE constexpr ::cuda::std::size_t static_size_v()
 {

diff --git a/cub/cub/device/device_for.cuh b/cub/cub/device/device_for.cuh
@@ -47,13 +47,10 @@
 #include <thrust/type_traits/is_contiguous_iterator.h>
 #include <thrust/type_traits/unwrap_contiguous_iterator.h>
 
+#include <cuda/std/__mdspan/extents.h>
 #include <cuda/std/iterator>
 #include <cuda/std/type_traits>
 
-#if __cccl_lib_mdspan
-#  include <cuda/std/__mdspan/extents.h>
-#endif // __cccl_lib_mdspan
-
 CUB_NAMESPACE_BEGIN
 
 namespace detail
@@ -843,8 +840,6 @@ public:
    * ForEachInExtents
    ********************************************************************************************************************/
 
-#if __cccl_lib_mdspan
-
   //! @rst
   //! Overview
   //! +++++++++++++++++++++++++++++++++++++++++++++
@@ -990,7 +985,6 @@ public:
     CUB_DETAIL_NVTX_RANGE_SCOPE("cub::DeviceFor::ForEachInExtents");
     return detail::for_each_in_extents::dispatch_t<extents_type, OpType>::dispatch(extents, op, stream);
   }
-#endif // __cccl_lib_mdspan
 };
 
 CUB_NAMESPACE_END
diff --git a/cub/cub/device/dispatch/dispatch_for_each_in_extents.cuh b/cub/cub/device/dispatch/dispatch_for_each_in_extents.cuh
@@ -37,39 +37,37 @@
 #  pragma system_header
 #endif // no system header
 
-#if __cccl_lib_mdspan
-
-#  include <cub/detail/mdspan_utils.cuh> // size(extent)
-#  include <cub/device/dispatch/kernels/for_each_in_extents_kernel.cuh>
-#  include <cub/device/dispatch/tuning/tuning_for.cuh>
-#  include <cub/util_device.cuh>
-#  include <cub/util_namespace.cuh>
-
-#  include <thrust/system/cuda/detail/core/triple_chevron_launch.h>
-
-#  include <cuda/std/__type_traits/integral_constant.h> // cuda::std::integral_constant
-#  include <cuda/std/__utility/integer_sequence.h> // cuda::std::index_sequence
-#  include <cuda/std/array> // cuda::std::array
-#  include <cuda/std/cstddef> // size_t
-
-#  define _CUB_RETURN_IF_ERROR(STATUS)       \
-    {                                        \
-      cudaError_t error_ = CubDebug(STATUS); \
-      if (error_ != cudaSuccess)             \
-      {                                      \
-        return error_;                       \
-      }                                      \
-    }
+#include <cub/detail/mdspan_utils.cuh> // size(extent)
+#include <cub/device/dispatch/kernels/for_each_in_extents_kernel.cuh>
+#include <cub/device/dispatch/tuning/tuning_for.cuh>
+#include <cub/util_device.cuh>
+#include <cub/util_namespace.cuh>
+
+#include <thrust/system/cuda/detail/core/triple_chevron_launch.h>
+
+#include <cuda/std/__type_traits/integral_constant.h> // cuda::std::integral_constant
+#include <cuda/std/__utility/integer_sequence.h> // cuda::std::index_sequence
+#include <cuda/std/array> // cuda::std::array
+#include <cuda/std/cstddef> // size_t
+
+#define _CUB_RETURN_IF_ERROR(STATUS)       \
+  {                                        \
+    cudaError_t error_ = CubDebug(STATUS); \
+    if (error_ != cudaSuccess)             \
+    {                                      \
+      return error_;                       \
+    }                                      \
+  }
 
-#  define _CUB_RETURN_IF_STREAM_ERROR(STREAM)                         \
-    {                                                                 \
-      cudaError_t error_ = CubDebug(detail::DebugSyncStream(STREAM)); \
-      if (error_ != cudaSuccess)                                      \
-      {                                                               \
-        CubDebug(error_ = SyncStream(STREAM));                        \
-        return error_;                                                \
-      }                                                               \
-    }
+#define _CUB_RETURN_IF_STREAM_ERROR(STREAM)                         \
+  {                                                                 \
+    cudaError_t error_ = CubDebug(detail::DebugSyncStream(STREAM)); \
+    if (error_ != cudaSuccess)                                      \
+    {                                                               \
+      CubDebug(error_ = SyncStream(STREAM));                        \
+      return error_;                                                \
+    }                                                               \
+  }
 
 CUB_NAMESPACE_BEGIN
 
@@ -115,13 +113,13 @@ public:
     constexpr unsigned items_per_thread = ActivePolicyT::for_policy_t::items_per_thread;
     unsigned num_cta                    = ::cuda::ceil_div(_size, block_threads * items_per_thread);
 
-#  ifdef CUB_DEBUG_LOG
+#ifdef CUB_DEBUG_LOG
     _CubLog("Invoking detail::for_each_in_extents::static_kernel<<<%u, %u, 0, %p>>>(), items_per_thread: %u\n",
             num_cta,
             block_threads,
             _stream,
             items_per_thread);
-#  endif
+#endif
     auto status =
       THRUST_NS_QUALIFIER::cuda_cub::detail::triple_chevron(num_cta, block_threads, 0, _stream)
         .doit(detail::for_each_in_extents::
@@ -153,13 +151,13 @@ public:
     _CUB_RETURN_IF_ERROR(status)
     unsigned num_cta = ::cuda::ceil_div(_size, block_threads * items_per_thread);
 
-#  ifdef CUB_DEBUG_LOG
+#ifdef CUB_DEBUG_LOG
     _CubLog("Invoking detail::for_each_in_extents::dynamic_kernel<<<%u, %u, 0, %p>>>(), items_per_thread: %u\n",
             num_cta,
             block_threads,
             _stream,
             items_per_thread);
-#  endif
+#endif
     status = THRUST_NS_QUALIFIER::cuda_cub::detail::triple_chevron(num_cta, block_threads, 0, _stream)
                .doit(kernel, _op, _ext, sub_sizes_div_array, extents_div_array);
     _CUB_RETURN_IF_ERROR(status)
@@ -204,5 +202,3 @@ private:
 } // namespace detail::for_each_in_extents
 
 CUB_NAMESPACE_END
-
-#endif // __cccl_lib_mdspan
diff --git a/cub/cub/device/dispatch/kernels/for_each_in_extents_kernel.cuh b/cub/cub/device/dispatch/kernels/for_each_in_extents_kernel.cuh
@@ -37,16 +37,14 @@
 #  pragma system_header
 #endif // no system header
 
-#if __cccl_lib_mdspan
+#include <cub/detail/fast_modulo_division.cuh> // fast_div_mod
+#include <cub/detail/mdspan_utils.cuh> // is_sub_size_static
+#include <cub/detail/type_traits.cuh> // implicit_prom_t
 
-#  include <cub/detail/fast_modulo_division.cuh> // fast_div_mod
-#  include <cub/detail/mdspan_utils.cuh> // is_sub_size_static
-#  include <cub/detail/type_traits.cuh> // implicit_prom_t
-
-#  include <cuda/std/cstddef> // size_t
-#  include <cuda/std/mdspan> // dynamic_extent
-#  include <cuda/std/type_traits> // enable_if
-#  include <cuda/std/utility> // index_sequence
+#include <cuda/std/cstddef> // size_t
+#include <cuda/std/mdspan> // dynamic_extent
+#include <cuda/std/type_traits> // enable_if
+#include <cuda/std/utility> // index_sequence
 
 CUB_NAMESPACE_BEGIN
 
@@ -131,11 +129,11 @@ computation(IndexType id, IndexType, Func func, ExtendType, FastDivModArrayType,
  **********************************************************************************************************************/
 
 // GCC6/7/8/9 raises unused parameter warning
-#  if _CCCL_COMPILER(GCC, <, 10)
-#    define _CUB_UNUSED_ATTRIBUTE __attribute__((unused))
-#  else
-#    define _CUB_UNUSED_ATTRIBUTE
-#  endif
+#if _CCCL_COMPILER(GCC, <, 10)
+#  define _CUB_UNUSED_ATTRIBUTE __attribute__((unused))
+#else
+#  define _CUB_UNUSED_ATTRIBUTE
+#endif
 
 template <typename ChainedPolicyT,
           typename Func,
@@ -186,5 +184,3 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void dynamic_kernel(
 } // namespace detail
 
 CUB_NAMESPACE_END
-
-#endif // __cccl_lib_mdspan
diff --git a/cub/test/catch2_test_device_for_each_in_extents_api.cu b/cub/test/catch2_test_device_for_each_in_extents_api.cu
@@ -30,20 +30,18 @@
 
 #include <cub/config.cuh>
 
-#if __cccl_lib_mdspan && !_CCCL_COMPILER(MSVC)
+#include <cub/device/device_for.cuh>
 
-#  include <cub/device/device_for.cuh>
+#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/device_vector.h>
+#include <thrust/fill.h>
+#include <thrust/host_vector.h>
 
-#  include <thrust/detail/raw_pointer_cast.h>
-#  include <thrust/device_vector.h>
-#  include <thrust/fill.h>
-#  include <thrust/host_vector.h>
+#include <cuda/std/array>
+#include <cuda/std/mdspan>
+#include <cuda/std/span>
 
-#  include <cuda/std/array>
-#  include <cuda/std/mdspan>
-#  include <cuda/std/span>
-
-#  include <c2h/catch2_test_helper.h>
+#include <c2h/catch2_test_helper.h>
 
 // example-begin for-each-in-extents-op
 struct linear_store_3D
@@ -87,5 +85,3 @@ C2H_TEST("Device ForEachInExtents", "[ForEachInExtents][device]")
   // example-end for-each-in-extents-example
 }
 // clang-format on
-
-#endif // __cccl_lib_mdspan && !_CCCL_COMPILER(MSVC)
diff --git a/cub/test/internal/catch2_test_fast_div_mod.cu b/cub/test/internal/catch2_test_fast_div_mod.cu
@@ -26,10 +26,8 @@
  ******************************************************************************/
 #include <cub/config.cuh>
 
-#if __cccl_lib_mdspan
-
-#  include "c2h/catch2_test_helper.h"
-#  include "c2h/utility.h"
+#include "c2h/catch2_test_helper.h"
+#include "c2h/utility.h"
 
 /***********************************************************************************************************************
  * TEST CASES
@@ -42,11 +40,11 @@ using index_types =
                  uint16_t,
                  int32_t,
                  uint32_t
-#  if _CCCL_HAS_INT128()
+#if _CCCL_HAS_INT128()
                  ,
                  int64_t,
                  uint64_t
-#  endif
+#endif
                  >;
 
 C2H_TEST("FastDivMod random", "[FastDivMod][Random]", index_types)
@@ -79,5 +77,3 @@ C2H_TEST("FastDivMod edge cases", "[FastDivMod][EdgeCases]", index_types)
   REQUIRE(0 == div_mod_min(0).quotient);
   REQUIRE(0 == div_mod_min(0).remainder);
 }
-
-#endif // __cccl_lib_mdspan