diff --git a/libcudacxx/cmake/LibcudacxxPublicHeaderTestingHost.cmake b/libcudacxx/cmake/LibcudacxxPublicHeaderTestingHost.cmake index 26e3c1fe1609..d37367332d54 100644 --- a/libcudacxx/cmake/LibcudacxxPublicHeaderTestingHost.cmake +++ b/libcudacxx/cmake/LibcudacxxPublicHeaderTestingHost.cmake @@ -6,6 +6,7 @@ # Meta target for all configs' header builds: add_custom_target(libcudacxx.test.public_headers_host_only) +add_custom_target(libcudacxx.test.public_headers_host_only_with_ctk) # Grep all public headers file( @@ -13,18 +14,8 @@ file( LIST_DIRECTORIES false RELATIVE "${libcudacxx_SOURCE_DIR}/include" CONFIGURE_DEPENDS + "${libcudacxx_SOURCE_DIR}/include/cuda/*" "${libcudacxx_SOURCE_DIR}/include/cuda/std/*" - # Add some files we expect to work in host only compilation - "${libcudacxx_SOURCE_DIR}/include/cuda/bit" - "${libcudacxx_SOURCE_DIR}/include/cuda/cmath" - "${libcudacxx_SOURCE_DIR}/include/cuda/functional" - "${libcudacxx_SOURCE_DIR}/include/cuda/iterator" - "${libcudacxx_SOURCE_DIR}/include/cuda/mdspan" - "${libcudacxx_SOURCE_DIR}/include/cuda/memory" - "${libcudacxx_SOURCE_DIR}/include/cuda/numeric" - "${libcudacxx_SOURCE_DIR}/include/cuda/type_traits" - "${libcudacxx_SOURCE_DIR}/include/cuda/utility" - "${libcudacxx_SOURCE_DIR}/include/cuda/version" ) set(public_host_header_cxx_compile_options) @@ -63,6 +54,38 @@ function(libcudacxx_create_public_header_test_host header_name headertest_src) ) endfunction() +function( + libcudacxx_create_public_header_test_host_with_ctk + header_name + headertest_src +) + # Create the default target for that file + add_library( + public_headers_host_only_with_ctk_${header_name} + SHARED + "${headertest_src}.cpp" + ) + cccl_configure_target(public_headers_host_only_with_ctk_${header_name}) + target_compile_definitions( + public_headers_host_only_with_ctk_${header_name} + PRIVATE # + ${public_host_header_cxx_compile_definitions} + _CCCL_HEADER_TEST + ) + target_compile_options( + public_headers_host_only_with_ctk_${header_name} + PRIVATE ${public_host_header_cxx_compile_options} + ) + target_link_libraries( + public_headers_host_only_with_ctk_${header_name} + PUBLIC libcudacxx.compiler_interface CUDA::cudart + ) + add_dependencies( + libcudacxx.test.public_headers_host_only_with_ctk + public_headers_host_only_with_ctk_${header_name} + ) +endfunction() + function(libcudacxx_add_public_headers_host_only header) # ${header} contains the "/" from the subfolder, replace by "_" for actual names string(REPLACE "/" "_" header_name "${header}") @@ -76,6 +99,7 @@ function(libcudacxx_add_public_headers_host_only header) # Create the default target for that file libcudacxx_create_public_header_test_host(${header_name} ${headertest_src}) + libcudacxx_create_public_header_test_host_with_ctk(${header_name} ${headertest_src}) endfunction() foreach (header IN LISTS public_headers_host_only) diff --git a/libcudacxx/include/cuda/__container/buffer.h b/libcudacxx/include/cuda/__container/buffer.h index 2f1f50e3e796..9b91010f3fa1 100644 --- a/libcudacxx/include/cuda/__container/buffer.h +++ b/libcudacxx/include/cuda/__container/buffer.h @@ -21,35 +21,37 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER() -# include -#endif // _CCCL_HAS_CUDA_COMPILER() - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# if _CCCL_HAS_CUDA_COMPILER() +# include +# endif // _CCCL_HAS_CUDA_COMPILER() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include //! @file The \c buffer class provides a container of contiguous memory _CCCL_BEGIN_NAMESPACE_CUDA @@ -309,7 +311,7 @@ class buffer __buf_.size()); } -#ifndef _CCCL_DOXYGEN_INVOKED // doxygen conflates the overloads +# ifndef _CCCL_DOXYGEN_INVOKED // doxygen conflates the overloads _CCCL_TEMPLATE(class _Range, class _Resource, class _Env = ::cuda::std::execution::env<>) _CCCL_REQUIRES( ::cuda::mr::synchronous_resource<::cuda::std::decay_t<_Resource>> _CCCL_AND __compatible_range<_Range> @@ -332,7 +334,7 @@ class buffer __unwrapped_begin(), __buf_.size()); } -#endif // _CCCL_DOXYGEN_INVOKED +# endif // _CCCL_DOXYGEN_INVOKED //! @} //! @addtogroup iterators @@ -447,7 +449,7 @@ class buffer return __buf_.data(); } -#ifndef _CCCL_DOXYGEN_INVOKED +# ifndef _CCCL_DOXYGEN_INVOKED //! @brief Returns a pointer to the first element of the buffer. If the buffer //! is empty, the returned pointer will be null. [[nodiscard]] _CCCL_HIDE_FROM_ABI pointer __unwrapped_begin() noexcept @@ -477,7 +479,7 @@ class buffer { return __buf_.data() + __buf_.size(); } -#endif // _CCCL_DOXYGEN_INVOKED +# endif // _CCCL_DOXYGEN_INVOKED //! @} @@ -683,14 +685,14 @@ __fill_n(cuda::stream_ref __stream, _Tp* __first, ::cuda::std::size_t __count, c } else { -#if _CCCL_HAS_CUDA_COMPILER() +# if _CCCL_HAS_CUDA_COMPILER() ::cuda::__ensure_current_context __guard(__stream); ::cub::DeviceTransform::Fill(__first, __count, __value, __stream.get()); -#else +# else static_assert(0, "CUDA compiler is required to initialize a buffer with " "elements larger than 4 bytes"); -#endif +# endif } } } @@ -887,6 +889,8 @@ auto make_buffer(stream_ref __stream, _Resource&& __mr, _Range&& __range, const } _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___CONTAINER_BUFFER_H diff --git a/libcudacxx/include/cuda/__container/heterogeneous_iterator.h b/libcudacxx/include/cuda/__container/heterogeneous_iterator.h index e33c40f20165..9a62fc277fbe 100644 --- a/libcudacxx/include/cuda/__container/heterogeneous_iterator.h +++ b/libcudacxx/include/cuda/__container/heterogeneous_iterator.h @@ -21,17 +21,19 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include //! @file The \c heterogeneous_iterator class is an iterator that provides typed execution space safety. _CCCL_BEGIN_NAMESPACE_CUDA @@ -290,7 +292,7 @@ class heterogeneous_iterator return __temp; } -#ifndef _CCCL_DOXYGEN_INVOKED // Do not document +# ifndef _CCCL_DOXYGEN_INVOKED // Do not document //! @brief Advance a \c heterogeneous_iterator //! @param __count The number of elements to advance. //! @param __other A heterogeneous_iterator. @@ -301,7 +303,7 @@ class heterogeneous_iterator __other += __count; return __other; } -#endif // _CCCL_DOXYGEN_INVOKED +# endif // _CCCL_DOXYGEN_INVOKED //! @brief Advance a \c heterogeneous_iterator by the negative value of \p __count //! @param __count The number of elements to advance. @@ -330,7 +332,7 @@ class heterogeneous_iterator return static_cast(this->__ptr_ - __other.__ptr_); } -#ifndef _CCCL_DOXYGEN_INVOKED // Do not document +# ifndef _CCCL_DOXYGEN_INVOKED // Do not document //! @brief Equality comparison between two heterogeneous_iterator //! @param __lhs A heterogeneous_iterator. //! @param __rhs Another heterogeneous_iterator. @@ -340,7 +342,7 @@ class heterogeneous_iterator { return __lhs.__ptr_ == __rhs.__ptr_; } -# if _CCCL_STD_VER <= 2017 +# if _CCCL_STD_VER <= 2017 //! @brief Inequality comparison between two heterogeneous_iterator //! @param __lhs A heterogeneous_iterator. //! @param __rhs Another heterogeneous_iterator. @@ -350,15 +352,15 @@ class heterogeneous_iterator { return __lhs.__ptr_ != __rhs.__ptr_; } -# endif // _CCCL_STD_VER <= 2017 +# endif // _CCCL_STD_VER <= 2017 -# if _LIBCUDACXX_HAS_SPACESHIP_OPERATOR() +# if _LIBCUDACXX_HAS_SPACESHIP_OPERATOR() [[nodiscard]] _CCCL_API friend constexpr ::cuda::std::strong_ordering operator<=>(const heterogeneous_iterator& __lhs, const heterogeneous_iterator& __rhs) noexcept { return __lhs.__ptr_ <=> __rhs.__ptr_; } -# else // ^^^ _LIBCUDACXX_HAS_SPACESHIP_OPERATOR() ^^^ / vvv !_LIBCUDACXX_HAS_SPACESHIP_OPERATOR() vvv +# else // ^^^ _LIBCUDACXX_HAS_SPACESHIP_OPERATOR() ^^^ / vvv !_LIBCUDACXX_HAS_SPACESHIP_OPERATOR() vvv //! @brief Less than relation between two heterogeneous_iterator //! @param __lhs A heterogeneous_iterator. //! @param __rhs Another heterogeneous_iterator. @@ -399,8 +401,8 @@ class heterogeneous_iterator { return __lhs.__ptr_ >= __rhs.__ptr_; } -# endif // !_LIBCUDACXX_HAS_SPACESHIP_OPERATOR() -#endif // _CCCL_DOXYGEN_INVOKED +# endif // !_LIBCUDACXX_HAS_SPACESHIP_OPERATOR() +# endif // _CCCL_DOXYGEN_INVOKED _CCCL_API constexpr pointer __unwrap() const noexcept { @@ -431,6 +433,8 @@ struct pointer_traits<::cuda::heterogeneous_iterator<_Tp, _Properties...>> _CCCL_END_NAMESPACE_CUDA_STD -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //__CUDAX__CONTAINERS_HETEROGENEOUS_ITERATOR_CUH diff --git a/libcudacxx/include/cuda/__container/uninitialized_async_buffer.h b/libcudacxx/include/cuda/__container/uninitialized_async_buffer.h index d897bb7d57c3..c9bcb8ab1e0d 100644 --- a/libcudacxx/include/cuda/__container/uninitialized_async_buffer.h +++ b/libcudacxx/include/cuda/__container/uninitialized_async_buffer.h @@ -21,19 +21,21 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include //! @file //! The \c __uninitialized_async_buffer class provides a typed buffer allocated @@ -138,7 +140,7 @@ class __uninitialized_async_buffer return {__self.__get_data(), __self.size()}; } -#ifndef _CCCL_DOXYGEN_INVOKED +# ifndef _CCCL_DOXYGEN_INVOKED // This is needed to ensure that we do not do a deep copy in // __replace_allocation struct __fake_resource_ref @@ -179,7 +181,7 @@ class __uninitialized_async_buffer _CCCL_REQUIRES(::cuda::std::__is_included_in_v<_Property, _Properties...>) _CCCL_HIDE_FROM_ABI friend constexpr void get_property(const __fake_resource_ref&, _Property) noexcept {} }; -#endif // _CCCL_DOXYGEN_INVOKED +# endif // _CCCL_DOXYGEN_INVOKED public: using value_type = _Tp; @@ -411,6 +413,8 @@ template using uninitialized_async_device_buffer = __uninitialized_async_buffer<_Tp, ::cuda::mr::device_accessible>; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //__CUDAX__CONTAINERS_UNINITIALIZED_ASYNC_BUFFER_H diff --git a/libcudacxx/include/cuda/__device/arch_traits.h b/libcudacxx/include/cuda/__device/arch_traits.h index f406c3782ca1..ef410120a281 100644 --- a/libcudacxx/include/cuda/__device/arch_traits.h +++ b/libcudacxx/include/cuda/__device/arch_traits.h @@ -21,17 +21,15 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CTK() - -# include -# include -# include -# include -# include -# include -# include +#include +#include +#include +#include +#include +#include +#include -# include +#include _CCCL_BEGIN_NAMESPACE_CUDA @@ -507,7 +505,11 @@ template <> case arch_id::sm_121a: return ::cuda::arch_traits(); default: +#if _CCCL_HAS_CTK() ::cuda::__throw_cuda_error(::cudaErrorInvalidValue, "Traits requested for an unknown architecture"); +#else // ^^^ _CCCL_HAS_CTK() ^^^ / vvv !_CCCL_HAS_CTK() vvv + ::cuda::__throw_cuda_error(/*cudaErrorInvalidValue*/ 1, "Traits requested for an unknown architecture"); +#endif // ^^^ !_CCCL_HAS_CTK() ^^^ break; } } @@ -522,7 +524,7 @@ template <> _CCCL_END_NAMESPACE_CUDA -# if _CCCL_CUDA_COMPILATION() +#if _CCCL_CUDA_COMPILATION() _CCCL_BEGIN_NAMESPACE_CUDA_DEVICE @@ -535,19 +537,17 @@ _CCCL_BEGIN_NAMESPACE_CUDA_DEVICE template [[nodiscard]] _CCCL_DEVICE_API inline _CCCL_TARGET_CONSTEXPR ::cuda::arch_traits_t current_arch_traits() noexcept { -# if _CCCL_DEVICE_COMPILATION() +# if _CCCL_DEVICE_COMPILATION() return ::cuda::arch_traits_for(::cuda::device::current_arch_id<_Dummy>()); -# else // ^^^ _CCCL_DEVICE_COMPILATION() ^^^ / vvv !_CCCL_DEVICE_COMPILATION() vvv +# else // ^^^ _CCCL_DEVICE_COMPILATION() ^^^ / vvv !_CCCL_DEVICE_COMPILATION() vvv return {}; -# endif // ^^^ !_CCCL_DEVICE_COMPILATION() ^^^ +# endif // ^^^ !_CCCL_DEVICE_COMPILATION() ^^^ } _CCCL_END_NAMESPACE_CUDA_DEVICE -# endif // _CCCL_CUDA_COMPILATION - -# include +#endif // _CCCL_CUDA_COMPILATION -#endif // _CCCL_HAS_CTK() +#include #endif // _CUDA___DEVICE_ARCH_TRAITS_H diff --git a/libcudacxx/include/cuda/__event/timed_event.h b/libcudacxx/include/cuda/__event/timed_event.h index 603c140fb7c1..a313f5bacffa 100644 --- a/libcudacxx/include/cuda/__event/timed_event.h +++ b/libcudacxx/include/cuda/__event/timed_event.h @@ -11,9 +11,6 @@ #ifndef _CUDA___EVENT_TIMED_EVENT_H #define _CUDA___EVENT_TIMED_EVENT_H -#include -// cuda_runtime_api needs to come first - #include #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) diff --git a/libcudacxx/include/cuda/__functional/for_each_canceled.h b/libcudacxx/include/cuda/__functional/for_each_canceled.h index 54a97ad0290b..b230911e5e8c 100644 --- a/libcudacxx/include/cuda/__functional/for_each_canceled.h +++ b/libcudacxx/include/cuda/__functional/for_each_canceled.h @@ -21,14 +21,14 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include +#if _CCCL_CUDA_COMPILATION() -#include +# include +# include +# include +# include -#if _CCCL_CUDA_COMPILATION() +# include # include diff --git a/libcudacxx/include/cuda/__fwd/devices.h b/libcudacxx/include/cuda/__fwd/devices.h index 5e77301b2e6a..3158b66ec2e1 100644 --- a/libcudacxx/include/cuda/__fwd/devices.h +++ b/libcudacxx/include/cuda/__fwd/devices.h @@ -27,10 +27,13 @@ _CCCL_BEGIN_NAMESPACE_CUDA +#if _CCCL_HAS_CTK() class __physical_device; class device_ref; template <::cudaDeviceAttr _Attr> struct __dev_attr; +#endif // _CCCL_HAS_CTK() + struct arch_traits_t; class compute_capability; enum class arch_id : int; diff --git a/libcudacxx/include/cuda/__hierarchy/block_level.h b/libcudacxx/include/cuda/__hierarchy/block_level.h index 37257ce9a1eb..dc746527bf6a 100644 --- a/libcudacxx/include/cuda/__hierarchy/block_level.h +++ b/libcudacxx/include/cuda/__hierarchy/block_level.h @@ -21,14 +21,16 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -43,6 +45,8 @@ struct block_level : __native_hierarchy_level_base using __base_type = __native_hierarchy_level_base; using __base_type::count_as; using __base_type::extents_as; + +# if _CCCL_CUDA_COMPILATION() using __base_type::index_as; using __base_type::rank_as; @@ -132,12 +136,15 @@ struct block_level : __native_hierarchy_level_base const auto __idx = index_as<_Tp>(__level, __hier_unpacked); return static_cast<_Tp>((__idx.z * __dims.y + __idx.y) * __dims.x + __idx.x); } +# endif // _CCCL_CUDA_COMPILATION() }; _CCCL_GLOBAL_CONSTANT block_level block; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_BLOCK_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/cluster_level.h b/libcudacxx/include/cuda/__hierarchy/cluster_level.h index 70ef8fdce89c..72a40127fbb4 100644 --- a/libcudacxx/include/cuda/__hierarchy/cluster_level.h +++ b/libcudacxx/include/cuda/__hierarchy/cluster_level.h @@ -21,14 +21,16 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -42,6 +44,8 @@ struct cluster_level : __native_hierarchy_level_base using __base_type = __native_hierarchy_level_base; using __base_type::extents_as; + +# if _CCCL_CUDA_COMPILATION() using __base_type::index_as; // interactions with grid level @@ -63,12 +67,15 @@ struct cluster_level : __native_hierarchy_level_base NV_IF_TARGET(NV_PROVIDES_SM_90, (__idx = ::__clusterIdx();)) return {static_cast<_Tp>(__idx.x), static_cast<_Tp>(__idx.y), static_cast<_Tp>(__idx.z)}; } +# endif // _CCCL_CUDA_COMPILATION() }; _CCCL_GLOBAL_CONSTANT cluster_level cluster; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_CLUSTER_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/dimensions.h b/libcudacxx/include/cuda/__hierarchy/dimensions.h index 395c4c6e996e..97151555b330 100644 --- a/libcudacxx/include/cuda/__hierarchy/dimensions.h +++ b/libcudacxx/include/cuda/__hierarchy/dimensions.h @@ -21,10 +21,12 @@ # pragma system_header #endif // no system header -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -166,6 +168,8 @@ template } // namespace __detail _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_DIMENSIONS_H diff --git a/libcudacxx/include/cuda/__hierarchy/get_launch_dimensions.h b/libcudacxx/include/cuda/__hierarchy/get_launch_dimensions.h index ee3d07ee8ae7..173c07413c15 100644 --- a/libcudacxx/include/cuda/__hierarchy/get_launch_dimensions.h +++ b/libcudacxx/include/cuda/__hierarchy/get_launch_dimensions.h @@ -21,14 +21,16 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -85,6 +87,8 @@ constexpr auto _CCCL_HOST get_launch_dimensions(const hierarchy_dimensions<_Leve _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_GET_LAUNCH_DIMENSIONS_H diff --git a/libcudacxx/include/cuda/__hierarchy/grid_level.h b/libcudacxx/include/cuda/__hierarchy/grid_level.h index 9cd6a111709d..3ab829798028 100644 --- a/libcudacxx/include/cuda/__hierarchy/grid_level.h +++ b/libcudacxx/include/cuda/__hierarchy/grid_level.h @@ -21,10 +21,12 @@ # pragma system_header #endif // no system header -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -39,6 +41,8 @@ _CCCL_GLOBAL_CONSTANT grid_level grid; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_GRID_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h index 816c5553f26c..6bf0f9894936 100644 --- a/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h @@ -21,22 +21,24 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -256,7 +258,7 @@ struct __hierarchy_extents_helper }; template -[[nodiscard]] _CCCL_DEVICE constexpr auto __static_index_hint(const dimensions<_Tp, _Extents...>& __dims, ::dim3 __index) +[[nodiscard]] _CCCL_DEVICE constexpr auto __static_index_hint(const dimensions<_Tp, _Extents...>&, ::dim3 __index) { using _HintedIndexT = dimensions<_Tp, (_Extents == 1 ? 0 : ::cuda::std::dynamic_extent)...>; return _HintedIndexT(__index.x, __index.y, __index.z); @@ -378,10 +380,10 @@ struct hierarchy_dimensions : levels(__ls) {} -#if !defined(_CCCL_NO_THREE_WAY_COMPARISON) && !_CCCL_COMPILER(MSVC, <, 19, 39) && !_CCCL_COMPILER(GCC, <, 12) +# if !defined(_CCCL_NO_THREE_WAY_COMPARISON) && !_CCCL_COMPILER(MSVC, <, 19, 39) && !_CCCL_COMPILER(GCC, <, 12) [[nodiscard]] _CCCL_HIDE_FROM_ABI constexpr bool operator==(const hierarchy_dimensions&) const noexcept = default; -#else // ^^^ !_CCCL_NO_THREE_WAY_COMPARISON ^^^ / vvv - // _CCCL_NO_THREE_WAY_COMPARISON vvv +# else // ^^^ !_CCCL_NO_THREE_WAY_COMPARISON ^^^ / vvv + // _CCCL_NO_THREE_WAY_COMPARISON vvv [[nodiscard]] _CCCL_API friend constexpr bool operator==(const hierarchy_dimensions& __left, const hierarchy_dimensions& __right) noexcept { @@ -393,7 +395,7 @@ struct hierarchy_dimensions { return __left.levels != __right.levels; } -#endif // _CCCL_NO_THREE_WAY_COMPARISON +# endif // _CCCL_NO_THREE_WAY_COMPARISON private: // This being static is a bit of a hack to make extents_type working without @@ -831,12 +833,12 @@ struct hierarchy_dimensions } } -#ifndef _CCCL_DOXYGEN_INVOKED // Do not document +# ifndef _CCCL_DOXYGEN_INVOKED // Do not document constexpr hierarchy_dimensions combine([[maybe_unused]] __empty_hierarchy __empty) const { return *this; } -#endif // _CCCL_DOXYGEN_INVOKED +# endif // _CCCL_DOXYGEN_INVOKED }; // TODO consider having LUnit optional argument for template argument deduction @@ -916,6 +918,8 @@ constexpr auto hierarchy_add_level(const hierarchy_dimensions<_Unit, _Levels...> } _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_HIERARCHY_DIMENSIONS_H diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_level_base.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_level_base.h index 6914f4882dee..b473e7826da6 100644 --- a/libcudacxx/include/cuda/__hierarchy/hierarchy_level_base.h +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_level_base.h @@ -21,19 +21,21 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -171,7 +173,7 @@ struct hierarchy_level_base __level, ::cuda::__unpack_hierarchy_if_needed(__hier)); } -#if _CCCL_CUDA_COMPILATION() +# if _CCCL_CUDA_COMPILATION() _CCCL_TEMPLATE(class _InLevel, class _Hierarchy) _CCCL_REQUIRES(__is_hierarchy_level_v<_InLevel> _CCCL_AND __is_or_has_hierarchy_member_v<_Hierarchy>) [[nodiscard]] _CCCL_DEVICE_API static constexpr auto index(const _InLevel& __level, const _Hierarchy& __hier) noexcept @@ -188,7 +190,7 @@ struct hierarchy_level_base return _Level::template rank_as<__default_1d_query_type<_InLevel>>( __level, ::cuda::__unpack_hierarchy_if_needed(__hier)); } -#endif // _CCCL_CUDA_COMPILATION() +# endif // _CCCL_CUDA_COMPILATION() _CCCL_TEMPLATE(class _Tp, class _InLevel, class _Hierarchy) _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_level_v<_InLevel> _CCCL_AND @@ -238,7 +240,7 @@ struct hierarchy_level_base return __count_as_impl<_Tp>(__level, ::cuda::__unpack_hierarchy_if_needed(__hier)); } -#if _CCCL_CUDA_COMPILATION() +# if _CCCL_CUDA_COMPILATION() _CCCL_TEMPLATE(class _Tp, class _InLevel, class _Hierarchy) _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_level_v<_InLevel> _CCCL_AND __is_or_has_hierarchy_member_v<_Hierarchy>) @@ -310,7 +312,7 @@ struct hierarchy_level_base } return __ret; } -#endif // _CCCL_CUDA_COMPILATION() +# endif // _CCCL_CUDA_COMPILATION() private: template @@ -361,6 +363,8 @@ struct hierarchy_level_base _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_HIERARCHY_LEVEL_BASE_H diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h index fe32c5e69bc8..2569f0403be9 100644 --- a/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h @@ -21,13 +21,15 @@ # pragma system_header #endif // no system header -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include -#include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -119,6 +121,7 @@ namespace __detail template struct __dims_helper; +# if _CCCL_CUDA_COMPILATION() template struct __dims_helper<_Level, _Level> { @@ -188,6 +191,7 @@ struct __dims_helper NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90, (return __clusterIdx();), (return ::dim3(0, 0, 0);)); } }; +# endif // _CCCL_CUDA_COMPILATION() // Seems like a compiler bug, where NODISCARD is marked as ignored due to void // return type, while its not possible to ever have void return type here @@ -414,6 +418,8 @@ _CCCL_DEVICE auto index(const _Unit&, const _Level&) } // namespace hierarchy _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_HIERARCHY_LEVELS_H diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_query_result.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_query_result.h index 0192198ba036..dd73dce727d0 100644 --- a/libcudacxx/include/cuda/__hierarchy/hierarchy_query_result.h +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_query_result.h @@ -21,13 +21,15 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -144,6 +146,8 @@ struct hierarchy_query_result _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_HIERARCHY_QUERY_RESULT_H diff --git a/libcudacxx/include/cuda/__hierarchy/level_dimensions.h b/libcudacxx/include/cuda/__hierarchy/level_dimensions.h index ec4fcd38d9fd..42a76be3a689 100644 --- a/libcudacxx/include/cuda/__hierarchy/level_dimensions.h +++ b/libcudacxx/include/cuda/__hierarchy/level_dimensions.h @@ -21,15 +21,17 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -140,10 +142,10 @@ struct level_dimensions _CCCL_API constexpr level_dimensions() : dims(){}; -#if !defined(_CCCL_NO_THREE_WAY_COMPARISON) && !_CCCL_COMPILER(MSVC, <, 19, 39) && !_CCCL_COMPILER(GCC, <, 12) +# if !defined(_CCCL_NO_THREE_WAY_COMPARISON) && !_CCCL_COMPILER(MSVC, <, 19, 39) && !_CCCL_COMPILER(GCC, <, 12) [[nodiscard]] _CCCL_HIDE_FROM_ABI constexpr bool operator==(const level_dimensions&) const noexcept = default; -#else // ^^^ !_CCCL_NO_THREE_WAY_COMPARISON ^^^ / vvv - // _CCCL_NO_THREE_WAY_COMPARISON vvv +# else // ^^^ !_CCCL_NO_THREE_WAY_COMPARISON ^^^ / vvv + // _CCCL_NO_THREE_WAY_COMPARISON vvv [[nodiscard]] _CCCL_API friend constexpr bool operator==(const level_dimensions& __left, const level_dimensions& __right) noexcept { @@ -155,7 +157,7 @@ struct level_dimensions { return __left.dims != __right.dims; } -#endif // _CCCL_NO_THREE_WAY_COMPARISON +# endif // _CCCL_NO_THREE_WAY_COMPARISON }; /** @@ -234,6 +236,8 @@ _CCCL_API constexpr auto block_dims(_Dims __dims) noexcept } _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_LEVEL_DIMENSIONS_H diff --git a/libcudacxx/include/cuda/__hierarchy/native_hierarchy_level_base.h b/libcudacxx/include/cuda/__hierarchy/native_hierarchy_level_base.h index 2d97796a7a41..fed901e297f4 100644 --- a/libcudacxx/include/cuda/__hierarchy/native_hierarchy_level_base.h +++ b/libcudacxx/include/cuda/__hierarchy/native_hierarchy_level_base.h @@ -21,33 +21,40 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA // cudafe++ makes the queries (that are device only) return void when compiling for host, which causes host compilers // to warn about applying [[nodiscard]] to a function that returns void. _CCCL_DIAG_PUSH -#if _CCCL_CUDA_COMPILER(NVCC) +# if _CCCL_CUDA_COMPILER(NVCC) _CCCL_DIAG_SUPPRESS_GCC("-Wattributes") _CCCL_DIAG_SUPPRESS_CLANG("-Wignored-attributes") _CCCL_DIAG_SUPPRESS_NVHPC(nodiscard_doesnt_apply) -#endif // _CCCL_CUDA_COMPILER(NVCC) +# endif // _CCCL_CUDA_COMPILER(NVCC) template struct __native_hierarchy_level_base : hierarchy_level_base<_Level> { + template + using __default_md_query_type = unsigned; + template + using __default_1d_query_type = ::cuda::std::size_t; + using __base_type = hierarchy_level_base<_Level>; using __base_type::count; using __base_type::count_as; @@ -55,18 +62,13 @@ struct __native_hierarchy_level_base : hierarchy_level_base<_Level> using __base_type::dims_as; using __base_type::extents; using __base_type::extents_as; + using __base_type::static_dims; + +# if _CCCL_CUDA_COMPILATION() using __base_type::index; using __base_type::index_as; using __base_type::rank; using __base_type::rank_as; - using __base_type::static_dims; - - template - using __default_md_query_type = unsigned; - template - using __default_1d_query_type = ::cuda::std::size_t; - -#if _CCCL_CUDA_COMPILATION() _CCCL_TEMPLATE(class _InLevel) _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) @@ -187,7 +189,7 @@ struct __native_hierarchy_level_base : hierarchy_level_base<_Level> return __ret; } -#endif // _CCCL_CUDA_COMPILATION() +# endif // _CCCL_CUDA_COMPILATION() }; _CCCL_DIAG_POP @@ -198,6 +200,8 @@ struct __native_hierarchy_level_base : hierarchy_level_base +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_NATIVE_HIERARCHY_LEVEL_BASE_H diff --git a/libcudacxx/include/cuda/__hierarchy/thread_level.h b/libcudacxx/include/cuda/__hierarchy/thread_level.h index f8beda576d1f..ad117180262b 100644 --- a/libcudacxx/include/cuda/__hierarchy/thread_level.h +++ b/libcudacxx/include/cuda/__hierarchy/thread_level.h @@ -21,14 +21,16 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -42,6 +44,8 @@ struct thread_level : __native_hierarchy_level_base using __base_type = __native_hierarchy_level_base; using __base_type::extents_as; + +# if _CCCL_CUDA_COMPILATION() using __base_type::index_as; using __base_type::rank_as; @@ -85,12 +89,15 @@ struct thread_level : __native_hierarchy_level_base { return static_cast<_Tp>(::cuda::ptx::get_sreg_laneid()); } +# endif // _CCCL_CUDA_COMPILATION() }; _CCCL_GLOBAL_CONSTANT thread_level gpu_thread; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_THREAD_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/traits.h b/libcudacxx/include/cuda/__hierarchy/traits.h index 77e44b8cfd3a..9818d7fc21cf 100644 --- a/libcudacxx/include/cuda/__hierarchy/traits.h +++ b/libcudacxx/include/cuda/__hierarchy/traits.h @@ -21,15 +21,17 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -124,6 +126,8 @@ inline constexpr bool __is_or_has_hierarchy_member_v = __has_hierarchy_member<_T _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_TRAITS_H diff --git a/libcudacxx/include/cuda/__hierarchy/warp_level.h b/libcudacxx/include/cuda/__hierarchy/warp_level.h index 6e1d91703916..bfbfff2af66f 100644 --- a/libcudacxx/include/cuda/__hierarchy/warp_level.h +++ b/libcudacxx/include/cuda/__hierarchy/warp_level.h @@ -21,17 +21,19 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -41,6 +43,8 @@ struct warp_level : __native_hierarchy_level_base using __base_type = __native_hierarchy_level_base; using __base_type::extents_as; + +# if _CCCL_CUDA_COMPILATION() using __base_type::index_as; _CCCL_TEMPLATE(class _Tp) @@ -56,12 +60,15 @@ struct warp_level : __native_hierarchy_level_base { return {static_cast<_Tp>(gpu_thread.rank(block) / 32), 0, 0}; } +# endif // _CCCL_CUDA_COMPILATION() }; _CCCL_GLOBAL_CONSTANT warp_level warp; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___HIERARCHY_WARP_LEVEL_H diff --git a/libcudacxx/include/cuda/__memcpy_async/dispatch_memcpy_async.h b/libcudacxx/include/cuda/__memcpy_async/dispatch_memcpy_async.h index 8fdfae61c6e9..5a281a1f1c20 100644 --- a/libcudacxx/include/cuda/__memcpy_async/dispatch_memcpy_async.h +++ b/libcudacxx/include/cuda/__memcpy_async/dispatch_memcpy_async.h @@ -55,8 +55,8 @@ template <::cuda::std::size_t _Align, typename _Group> char* __dest_char, char const* __src_char, ::cuda::std::size_t __size, - ::cuda::std::uint32_t __allowed_completions, - ::cuda::std::uint64_t* __bar_handle) + [[maybe_unused]] ::cuda::std::uint32_t __allowed_completions, + [[maybe_unused]] ::cuda::std::uint64_t* __bar_handle) { ::cuda::__cp_async_fallback_mechanism<_Align>(__group, __dest_char, __src_char, __size); return __completion_mechanism::__sync; @@ -68,8 +68,8 @@ template <::cuda::std::size_t _Align, typename _Group> char* __dest_char, char const* __src_char, ::cuda::std::size_t __size, - ::cuda::std::uint32_t __allowed_completions, - ::cuda::std::uint64_t* __bar_handle) + [[maybe_unused]] ::cuda::std::uint32_t __allowed_completions, + [[maybe_unused]] ::cuda::std::uint64_t* __bar_handle) { #if __cccl_ptx_isa >= 800 NV_IF_TARGET( @@ -114,8 +114,8 @@ template <::cuda::std::size_t _Align, typename _Group> char* __dest_char, char const* __src_char, ::cuda::std::size_t __size, - ::cuda::std::uint32_t __allowed_completions, - ::cuda::std::uint64_t* __bar_handle) + [[maybe_unused]] ::cuda::std::uint32_t __allowed_completions, + [[maybe_unused]] ::cuda::std::uint64_t* __bar_handle) { NV_IF_ELSE_TARGET( NV_IS_DEVICE, @@ -149,8 +149,8 @@ template <::cuda::std::size_t _Align, typename _Group> _Group const& __group, char* __dest_char, char const* __src_char, - ::cuda::std::size_t __size, - ::cuda::std::uint32_t __allowed_completions) + [[maybe_unused]] ::cuda::std::size_t __size, + [[maybe_unused]] ::cuda::std::uint32_t __allowed_completions) { _CCCL_ASSERT(!(__allowed_completions & uint32_t(__completion_mechanism::__mbarrier_complete_tx)), "Cannot allow mbarrier_complete_tx completion mechanism when not passing a barrier. "); diff --git a/libcudacxx/include/cuda/__memcpy_async/is_local_smem_barrier.h b/libcudacxx/include/cuda/__memcpy_async/is_local_smem_barrier.h index 7c9f69bf8d9d..994d557f4353 100644 --- a/libcudacxx/include/cuda/__memcpy_async/is_local_smem_barrier.h +++ b/libcudacxx/include/cuda/__memcpy_async/is_local_smem_barrier.h @@ -38,7 +38,7 @@ template > -_CCCL_API inline bool __is_local_smem_barrier(barrier<_Sco, _CompF>& __barrier) +_CCCL_API inline bool __is_local_smem_barrier([[maybe_unused]] barrier<_Sco, _CompF>& __barrier) { NV_IF_ELSE_TARGET( NV_IS_DEVICE, diff --git a/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h b/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h index a9441d0d40ee..cd0dcccfc9d6 100644 --- a/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h +++ b/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h @@ -107,8 +107,8 @@ struct __memcpy_completion_impl } template - [[nodiscard]] _CCCL_API inline static async_contract_fulfillment __defer_non_smem_barrier( - __completion_mechanism __cm, _Group const& __group, ::cuda::std::size_t __size, barrier<_Sco, _CompF>& __barrier) + [[nodiscard]] _CCCL_API inline static async_contract_fulfillment + __defer_non_smem_barrier(__completion_mechanism __cm, _Group const&, ::cuda::std::size_t, barrier<_Sco, _CompF>&) { // Overload for non-smem barriers. switch (__cm) diff --git a/libcudacxx/include/cuda/__memcpy_async/try_get_barrier_handle.h b/libcudacxx/include/cuda/__memcpy_async/try_get_barrier_handle.h index f73962782dc3..5817dfae535b 100644 --- a/libcudacxx/include/cuda/__memcpy_async/try_get_barrier_handle.h +++ b/libcudacxx/include/cuda/__memcpy_async/try_get_barrier_handle.h @@ -36,7 +36,7 @@ _CCCL_BEGIN_NAMESPACE_CUDA //! @brief __try_get_barrier_handle returns barrier handle of block-scoped barriers and a nullptr otherwise. template -_CCCL_API inline ::cuda::std::uint64_t* __try_get_barrier_handle(barrier<_Sco, _CompF>& __barrier) +_CCCL_API inline ::cuda::std::uint64_t* __try_get_barrier_handle(barrier<_Sco, _CompF>&) { return nullptr; } diff --git a/libcudacxx/include/cuda/__memory_pool/device_memory_pool.h b/libcudacxx/include/cuda/__memory_pool/device_memory_pool.h index 9cc447088df0..1c8838b0e2b5 100644 --- a/libcudacxx/include/cuda/__memory_pool/device_memory_pool.h +++ b/libcudacxx/include/cuda/__memory_pool/device_memory_pool.h @@ -21,18 +21,15 @@ # pragma system_header #endif // no system header -#if _CCCL_CUDA_COMPILER(CLANG) -# include -# include -#endif // _CCCL_CUDA_COMPILER(CLANG) +#if _CCCL_HAS_CTK() -#include -#include -#include -#include -#include +# include +# include +# include +# include +# include -#include +# include //! @file //! The \c device_memory_pool class provides an asynchronous memory resource @@ -161,6 +158,8 @@ static_assert(::cuda::mr::resource_with +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_DEVICE_MEMORY_POOL_H diff --git a/libcudacxx/include/cuda/__memory_pool/memory_pool_base.h b/libcudacxx/include/cuda/__memory_pool/memory_pool_base.h index d5de79377806..01db9b8fe3be 100644 --- a/libcudacxx/include/cuda/__memory_pool/memory_pool_base.h +++ b/libcudacxx/include/cuda/__memory_pool/memory_pool_base.h @@ -21,23 +21,20 @@ # pragma system_header #endif // no system header -#if _CCCL_CUDA_COMPILER(CLANG) -# include -# include -#endif // _CCCL_CUDA_COMPILER(CLANG) - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -199,9 +196,9 @@ inline void __verify_device_supports_export_handle_type( return; } if (__location.type != ::CU_MEM_LOCATION_TYPE_DEVICE -#if _CCCL_CTK_AT_LEAST(12, 6) +# if _CCCL_CTK_AT_LEAST(12, 6) && __location.type != ::CU_MEM_LOCATION_TYPE_HOST_NUMA -#endif +# endif ) { ::cuda::__throw_cuda_error(::cudaErrorNotSupported, @@ -222,17 +219,17 @@ __get_default_memory_pool(const CUmemLocation __location, [[maybe_unused]] const auto __device = __location.type == ::CU_MEM_LOCATION_TYPE_DEVICE ? __location.id : 0; ::cuda::__verify_device_supports_stream_ordered_allocations(__device); -#if _CCCL_CTK_AT_LEAST(13, 0) +# if _CCCL_CTK_AT_LEAST(13, 0) ::cudaMemPool_t __pool = ::cuda::__driver::__getDefaultMemPool(__location, __allocation_type); if (::cuda::memory_pool_attributes::release_threshold(__pool) == 0) { ::cuda::memory_pool_attributes::release_threshold.set(__pool, ::cuda::std::numeric_limits::max()); } -#else // ^^^ _CCCL_CTK_AT_LEAST(13, 0) ^^^ / vvv _CCCL_CTK_BELOW(13, 0) vvv +# else // ^^^ _CCCL_CTK_AT_LEAST(13, 0) ^^^ / vvv _CCCL_CTK_BELOW(13, 0) vvv _CCCL_ASSERT(__location.type == ::CU_MEM_LOCATION_TYPE_DEVICE, "Before CUDA 13 only device memory pools have a default"); ::cudaMemPool_t __pool = ::cuda::__driver::__deviceGetDefaultMemPool(__device); -#endif // ^^^ _CCCL_CTK_BELOW(13, 0) ^^^ +# endif // ^^^ _CCCL_CTK_BELOW(13, 0) ^^^ return __pool; } @@ -294,27 +291,27 @@ struct memory_pool_properties __pool_properties.handleTypes = ::CUmemAllocationHandleType(__properties.allocation_handle_type); __pool_properties.location = __location; -#if _CCCL_CTK_AT_LEAST(12, 2) +# if _CCCL_CTK_AT_LEAST(12, 2) if (__properties.max_pool_size != 0) { -# if _CCCL_CTK_AT_LEAST(13, 0) +# if _CCCL_CTK_AT_LEAST(13, 0) if (__allocation_type == ::CU_MEM_ALLOCATION_TYPE_MANAGED) { ::cuda::std::__throw_invalid_argument("Max pool size is not supported for managed memory pools"); } -# endif // _CCCL_CTK_AT_LEAST(13, 0) +# endif // _CCCL_CTK_AT_LEAST(13, 0) if (__properties.initial_pool_size > __properties.max_pool_size) { ::cuda::std::__throw_invalid_argument("Initial pool size must be less than the max pool size"); } } __pool_properties.maxSize = __properties.max_pool_size; -#else +# else if (__properties.max_pool_size != 0) { ::cuda::std::__throw_invalid_argument("Max pool size is not supported on this CUDA version"); } -#endif // _CCCL_CTK_AT_LEAST(12, 2) +# endif // _CCCL_CTK_AT_LEAST(12, 2) if (__properties.initial_pool_size > __properties.release_threshold) { @@ -627,18 +624,20 @@ class __memory_pool_base return __pool_ == __rhs.__pool_; } -#if _CCCL_STD_VER <= 2017 +# if _CCCL_STD_VER <= 2017 //! @brief Inequality comparison with another __memory_pool_base. //! @returns true if underlying \c cudaMemPool_t are not equal. [[nodiscard]] _CCCL_HOST_API bool operator!=(__memory_pool_base const& __rhs) const noexcept { return __pool_ != __rhs.__pool_; } -#endif // _CCCL_STD_VER <= 2017 +# endif // _CCCL_STD_VER <= 2017 }; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___MEMORY_RESOURCE_MEMORY_POOL_BASE_H diff --git a/libcudacxx/include/cuda/__memory_pool/pinned_memory_pool.h b/libcudacxx/include/cuda/__memory_pool/pinned_memory_pool.h index 89318fc59209..38b9702e6563 100644 --- a/libcudacxx/include/cuda/__memory_pool/pinned_memory_pool.h +++ b/libcudacxx/include/cuda/__memory_pool/pinned_memory_pool.h @@ -21,25 +21,22 @@ # pragma system_header #endif // no system header -#if _CCCL_CUDA_COMPILER(CLANG) -# include -# include -#endif // _CCCL_CUDA_COMPILER(CLANG) +#if _CCCL_HAS_CTK() -#include -#include -#include -#include -#include +# include +# include +# include +# include +# include -#include +# include //! @file //! The \c pinned_memory_resource class provides a memory resource that //! allocates pinned memory. _CCCL_BEGIN_NAMESPACE_CUDA -#if _CCCL_CTK_AT_LEAST(12, 6) +# if _CCCL_CTK_AT_LEAST(12, 6) static ::cudaMemPool_t __get_default_host_pinned_pool(); @@ -110,7 +107,7 @@ struct pinned_memory_pool : pinned_memory_pool_ref { using reference_type = pinned_memory_pool_ref; -# if _CCCL_CTK_AT_LEAST(13, 0) +# if _CCCL_CTK_AT_LEAST(13, 0) //! @brief Constructs a \c pinned_memory_pool with optional properties. //! Properties include the initial pool size and the release threshold. If the //! pool size grows beyond the release threshold, unused memory held by the @@ -129,7 +126,7 @@ struct pinned_memory_pool : pinned_memory_pool_ref { enable_access_from(cuda::devices); } -# endif // _CCCL_CTK_AT_LEAST(13, 0) +# endif // _CCCL_CTK_AT_LEAST(13, 0) //! @brief Constructs a \c pinned_memory_pool with the specified NUMA node id //! and optional properties. Properties include the initial pool size and the @@ -190,7 +187,7 @@ static_assert(::cuda::mr::resource_with +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_PINNED_MEMORY_POOL_H diff --git a/libcudacxx/include/cuda/__memory_resource/any_resource.h b/libcudacxx/include/cuda/__memory_resource/any_resource.h index 97f4e370034c..949000c91a21 100644 --- a/libcudacxx/include/cuda/__memory_resource/any_resource.h +++ b/libcudacxx/include/cuda/__memory_resource/any_resource.h @@ -21,19 +21,21 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA_MR -#ifndef _CCCL_DOXYGEN_INVOKED // Do not document this +# ifndef _CCCL_DOXYGEN_INVOKED // Do not document this template using __property_result_t _CCCL_NODEBUG_ALIAS = ::cuda::std::__type_call1< // @@ -376,7 +378,7 @@ synchronous_resource_ref<_Properties...> __as_resource_ref(resource_ref<_Propert return __mr; } -#else // ^^^ !_CCCL_DOXYGEN_INVOKED ^^^ / vvv _CCCL_DOXYGEN_INVOKED vvv +# else // ^^^ !_CCCL_DOXYGEN_INVOKED ^^^ / vvv _CCCL_DOXYGEN_INVOKED vvv enum class _ResourceKind { @@ -821,7 +823,7 @@ using synchronous_resource_ref = basic_resource_ref<_ResourceKind::_Synchronous, template using resource_ref = basic_resource_ref<_ResourceKind::_Asynchronous, _Properties...>; -#endif // _CCCL_DOXYGEN_INVOKED +# endif // _CCCL_DOXYGEN_INVOKED //! @rst //! .. _cudax-memory-resource-make-any-resource: @@ -877,6 +879,8 @@ auto make_any_resource(_Args&&... __args) -> any_resource<_Properties...> _CCCL_END_NAMESPACE_CUDA_MR -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___MEMORY_RESOURCE_ANY_RESOURCE_H diff --git a/libcudacxx/include/cuda/__memory_resource/get_memory_resource.h b/libcudacxx/include/cuda/__memory_resource/get_memory_resource.h index f9a2bd5c1ea4..3c984f86df3b 100644 --- a/libcudacxx/include/cuda/__memory_resource/get_memory_resource.h +++ b/libcudacxx/include/cuda/__memory_resource/get_memory_resource.h @@ -21,15 +21,17 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA_MR @@ -77,6 +79,8 @@ _CCCL_GLOBAL_CONSTANT auto get_memory_resource = get_memory_resource_t{}; _CCCL_END_NAMESPACE_CUDA_MR -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA__MEMORY_RESOURCE_GET_MEMORY_RESOURCE_H diff --git a/libcudacxx/include/cuda/__memory_resource/get_property.h b/libcudacxx/include/cuda/__memory_resource/get_property.h index 4de47677cd9a..9623d4f62107 100644 --- a/libcudacxx/include/cuda/__memory_resource/get_property.h +++ b/libcudacxx/include/cuda/__memory_resource/get_property.h @@ -21,12 +21,14 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -148,6 +150,8 @@ using forward_property = __forward_property::__fn<_Derived, _Upstream>; _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_GET_PROPERTY_H diff --git a/libcudacxx/include/cuda/__memory_resource/legacy_managed_memory_resource.h b/libcudacxx/include/cuda/__memory_resource/legacy_managed_memory_resource.h index 8db6cffbf549..fe3c7af16863 100644 --- a/libcudacxx/include/cuda/__memory_resource/legacy_managed_memory_resource.h +++ b/libcudacxx/include/cuda/__memory_resource/legacy_managed_memory_resource.h @@ -21,20 +21,18 @@ # pragma system_header #endif // no system header -#if _CCCL_CUDA_COMPILER(CLANG) -# include -#endif // _CCCL_CUDA_COMPILER(CLANG) +#if _CCCL_HAS_CTK() -#include -#include -#include -#include -#include -#include -#include -#include +# include +# include +# include +# include +# include +# include +# include +# include -#include +# include //! @file //! The \c managed_memory_resource class provides a memory resource that allocates managed memory. @@ -107,7 +105,7 @@ class legacy_managed_memory_resource { return __flags_ == __other.__flags_; } -#if _CCCL_STD_VER <= 2017 +# if _CCCL_STD_VER <= 2017 //! @brief Inequality comparison with another \c managed_memory_resource. //! @param __other The other \c managed_memory_resource. //! @return Whether both \c managed_memory_resource were constructed with different flags. @@ -115,7 +113,7 @@ class legacy_managed_memory_resource { return __flags_ != __other.__flags_; } -#endif // _CCCL_STD_VER <= 2017 +# endif // _CCCL_STD_VER <= 2017 //! @brief Enables the \c device_accessible property _CCCL_HOST_API friend constexpr void @@ -143,6 +141,8 @@ static_assert(::cuda::mr::synchronous_resource_with +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_LEGACY_MANAGED_MEMORY_RESOURCE_H diff --git a/libcudacxx/include/cuda/__memory_resource/legacy_pinned_memory_resource.h b/libcudacxx/include/cuda/__memory_resource/legacy_pinned_memory_resource.h index a43ccb248001..911dafb428e5 100644 --- a/libcudacxx/include/cuda/__memory_resource/legacy_pinned_memory_resource.h +++ b/libcudacxx/include/cuda/__memory_resource/legacy_pinned_memory_resource.h @@ -21,20 +21,17 @@ # pragma system_header #endif // no system header -#if _CCCL_CUDA_COMPILER(CLANG) -# include -# include -#endif // _CCCL_CUDA_COMPILER(CLANG) +#if _CCCL_HAS_CTK() -#include -#include -#include -#include -#include -#include -#include +# include +# include +# include +# include +# include +# include +# include -#include +# include //! @file //! The \c legacy_pinned_memory_resource class provides a memory resource that allocates pinned memory. @@ -99,7 +96,7 @@ class legacy_pinned_memory_resource { return true; } -#if _CCCL_STD_VER <= 2017 +# if _CCCL_STD_VER <= 2017 //! @brief Equality comparison with another \c legacy_pinned_memory_resource. //! @param __other The other \c legacy_pinned_memory_resource. //! @return Whether both \c legacy_pinned_memory_resource were constructed with different flags. @@ -107,7 +104,7 @@ class legacy_pinned_memory_resource { return false; } -#endif // _CCCL_STD_VER <= 2017 +# endif // _CCCL_STD_VER <= 2017 //! @brief Enables the \c device_accessible property _CCCL_HOST_API friend constexpr void @@ -136,6 +133,8 @@ static_assert(::cuda::mr::synchronous_resource_with +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_LEGACY_PINNED_MEMORY_RESOURCE_H diff --git a/libcudacxx/include/cuda/__memory_resource/properties.h b/libcudacxx/include/cuda/__memory_resource/properties.h index c6f8ade529b7..6eac8634755c 100644 --- a/libcudacxx/include/cuda/__memory_resource/properties.h +++ b/libcudacxx/include/cuda/__memory_resource/properties.h @@ -21,11 +21,13 @@ # pragma system_header #endif // no system header -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA_MR @@ -125,6 +127,8 @@ struct __memory_accessability_from_properties _CCCL_END_NAMESPACE_CUDA_MR -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_PROPERTIES_H diff --git a/libcudacxx/include/cuda/__memory_resource/resource.h b/libcudacxx/include/cuda/__memory_resource/resource.h index 3e169651a95b..c9ffeacfc451 100644 --- a/libcudacxx/include/cuda/__memory_resource/resource.h +++ b/libcudacxx/include/cuda/__memory_resource/resource.h @@ -21,19 +21,21 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA_MR @@ -130,6 +132,8 @@ _CCCL_CONCEPT __non_polymorphic_resources = _CCCL_REQUIRES_EXPR((_Resource, _Oth _CCCL_END_NAMESPACE_CUDA_MR -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_RESOURCE_H diff --git a/libcudacxx/include/cuda/__memory_resource/shared_resource.h b/libcudacxx/include/cuda/__memory_resource/shared_resource.h index fe01c2e798f1..3bfc7d590eb8 100644 --- a/libcudacxx/include/cuda/__memory_resource/shared_resource.h +++ b/libcudacxx/include/cuda/__memory_resource/shared_resource.h @@ -21,19 +21,22 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA_MR + //! @rst //! .. _cudax-memory-resource-shared-resource: //! @@ -254,8 +257,11 @@ auto make_shared_resource(_Args&&... __args) -> shared_resource<_Resource> "_Resource does not satisfy the cuda::mr::synchronous_resource concept"); return shared_resource<_Resource>{::cuda::std::in_place_type<_Resource>, ::cuda::std::forward<_Args>(__args)...}; } + _CCCL_END_NAMESPACE_CUDA_MR -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___MEMORY_RESOURCE_SHARED_RESOURCE_H diff --git a/libcudacxx/include/cuda/__memory_resource/synchronous_resource_adapter.h b/libcudacxx/include/cuda/__memory_resource/synchronous_resource_adapter.h index 634ad7f7b336..65eecbcbe0f1 100644 --- a/libcudacxx/include/cuda/__memory_resource/synchronous_resource_adapter.h +++ b/libcudacxx/include/cuda/__memory_resource/synchronous_resource_adapter.h @@ -21,13 +21,15 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include +#if _CCCL_HAS_CTK() -#include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA_MR @@ -103,12 +105,12 @@ struct synchronous_resource_adapter return __resource == __rhs.__resource; } -#if _CCCL_STD_VER <= 2017 +# if _CCCL_STD_VER <= 2017 [[nodiscard]] _CCCL_HOST_API bool operator!=(const synchronous_resource_adapter& __rhs) const noexcept { return __resource != __rhs.__resource; } -#endif // _CCCL_STD_VER <= 2017 +# endif // _CCCL_STD_VER <= 2017 _CCCL_HOST_API _Resource& upstream_resource() noexcept { @@ -138,6 +140,8 @@ _CCCL_HOST_API decltype(auto) __adapt_if_synchronous(_Resource&& __resource) noe } _CCCL_END_NAMESPACE_CUDA_MR -#include +# include + +#endif // _CCCL_HAS_CTK() #endif //_CUDA___MEMORY_RESOURCE_SYNCHRONOUS_RESOURCE_ADAPTER_H diff --git a/libcudacxx/include/cuda/__stream/internal_streams.h b/libcudacxx/include/cuda/__stream/internal_streams.h index e4f74c9b9b29..6cce1db7c2c3 100644 --- a/libcudacxx/include/cuda/__stream/internal_streams.h +++ b/libcudacxx/include/cuda/__stream/internal_streams.h @@ -21,11 +21,11 @@ # pragma system_header #endif // no system header -#include +#if _CCCL_HAS_CTK() -#include +# include -#include +# include _CCCL_BEGIN_NAMESPACE_CUDA @@ -44,6 +44,8 @@ inline ::cuda::stream_ref __cccl_allocation_stream() _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA___STREAM_INTERNAL_STREAMS_H diff --git a/libcudacxx/include/cuda/__stream/launch_transform.h b/libcudacxx/include/cuda/__stream/launch_transform.h index d0dbfb7851fa..cfe876466503 100644 --- a/libcudacxx/include/cuda/__stream/launch_transform.h +++ b/libcudacxx/include/cuda/__stream/launch_transform.h @@ -21,20 +21,22 @@ # pragma system_header #endif // no system header -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#if _CCCL_HAS_CTK() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include _CCCL_BEGIN_NAMESPACE_CUDA namespace __detail @@ -188,6 +190,8 @@ using transformed_device_argument_t _CCCL_NODEBUG_ALIAS = _CCCL_END_NAMESPACE_CUDA -#include +# include + +#endif // _CCCL_HAS_CTK() #endif // _CUDA__STREAM_LAUNCH_TRANSFORM_H diff --git a/libcudacxx/include/cuda/pipeline b/libcudacxx/include/cuda/pipeline index 110fe60e70a9..e100d6aa3b5c 100644 --- a/libcudacxx/include/cuda/pipeline +++ b/libcudacxx/include/cuda/pipeline @@ -473,15 +473,15 @@ _CCCL_API inline pipeline make_pipeline() } template -_CCCL_API inline void pipeline_consumer_wait_prior(pipeline& __pipeline) +_CCCL_API inline void pipeline_consumer_wait_prior([[maybe_unused]] pipeline& __pipeline) { NV_IF_TARGET(NV_PROVIDES_SM_80, ::cuda::device::__pipeline_consumer_wait<_Prior>(__pipeline); __pipeline.__tail = __pipeline.__head - _Prior;) } template -_CCCL_API inline void -pipeline_producer_commit([[maybe_unused]] pipeline& __pipeline, barrier<_Scope>& __barrier) +_CCCL_API inline void pipeline_producer_commit([[maybe_unused]] pipeline& __pipeline, + [[maybe_unused]] barrier<_Scope>& __barrier) { NV_IF_TARGET(NV_PROVIDES_SM_80, ((void) __memcpy_completion_impl::__defer(