diff --git a/cub/cub/detail/mdspan_utils.cuh b/cub/cub/detail/mdspan_utils.cuh index eaa4186b871..2cf97db1604 100644 --- 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 // fast_div_mod -# include // fast_div_mod - -# include // cuda::std::array -# include // size_t -# include -# include // make_unsigned_t -# include // ::cuda::std::index_sequence +#include // cuda::std::array +#include // size_t +#include +#include // make_unsigned_t +#include // ::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 index 463281de6c0..bfc84f48c36 100644 --- a/cub/cub/detail/type_traits.cuh +++ b/cub/cub/detail/type_traits.cuh @@ -48,11 +48,9 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH #include _CCCL_SUPPRESS_DEPRECATED_POP -#include -#if __cccl_lib_mdspan -# include -#endif // __cccl_lib_mdspan #include // IWYU pragma: keep +#include +#include #include #include @@ -103,16 +101,12 @@ template struct is_fixed_size_random_access_range<::cuda::std::array, void> : ::cuda::std::true_type {}; -#if __cccl_lib_mdspan - template struct is_fixed_size_random_access_range< ::cuda::std::mdspan, ::cuda::std::enable_if_t> : ::cuda::std::true_type {}; -#endif // __cccl_lib_mdspan - template using is_fixed_size_random_access_range_t = typename is_fixed_size_random_access_range::type; @@ -143,16 +137,12 @@ template struct static_size<::cuda::std::span, void> : ::cuda::std::integral_constant {}; -#if __cccl_lib_mdspan - template struct static_size<::cuda::std::mdspan, ::cuda::std::enable_if_t> : ::cuda::std::integral_constant {}; -#endif // __cccl_lib_mdspan - template _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 index 5c49c5e2d72..d580eb0d999 100644 --- a/cub/cub/device/device_for.cuh +++ b/cub/cub/device/device_for.cuh @@ -47,13 +47,10 @@ #include #include +#include #include #include -#if __cccl_lib_mdspan -# include -#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::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 index 4fbd04a9bfd..1b8f4bc5b6e 100644 --- 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 // size(extent) -# include -# include -# include -# include - -# include - -# include // cuda::std::integral_constant -# include // cuda::std::index_sequence -# include // cuda::std::array -# include // size_t - -# define _CUB_RETURN_IF_ERROR(STATUS) \ - { \ - cudaError_t error_ = CubDebug(STATUS); \ - if (error_ != cudaSuccess) \ - { \ - return error_; \ - } \ - } +#include // size(extent) +#include +#include +#include +#include + +#include + +#include // cuda::std::integral_constant +#include // cuda::std::index_sequence +#include // cuda::std::array +#include // 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 index 4890ddcae29..f54f141c1f6 100644 --- 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 // fast_div_mod +#include // is_sub_size_static +#include // implicit_prom_t -# include // fast_div_mod -# include // is_sub_size_static -# include // implicit_prom_t - -# include // size_t -# include // dynamic_extent -# include // enable_if -# include // index_sequence +#include // size_t +#include // dynamic_extent +#include // enable_if +#include // 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 -#if __cccl_lib_mdspan && !_CCCL_COMPILER(MSVC) +#include -# include +#include +#include +#include +#include -# include -# include -# include -# include +#include +#include +#include -# include -# include -# include - -# include +#include // 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 index f604d82776c..7fd7df9c9a1 100644 --- 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 -#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