diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index 75832c29813..4e33b7937bf 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -37,7 +37,7 @@ enum BlockHistogramMemoryPreference BLEND }; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() inline ::std::ostream& operator<<(::std::ostream& os, BlockHistogramMemoryPreference mempref) { switch (mempref) @@ -52,7 +52,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, BlockHistogramMemoryPrefer return os << "(mempref) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() //! Parameterizable tuning policy type for AgentHistogram //! diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index 15c81f26422..073aa8ea107 100644 --- a/cub/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh @@ -55,7 +55,7 @@ enum RadixSortStoreAlgorithm RADIX_SORT_STORE_ALIGNED }; -#if !_CCCL_COMPILER(NVRTC) && !defined(_CCCL_DOXYGEN_INVOKED) +#if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED) inline ::std::ostream& operator<<(::std::ostream& os, RadixSortStoreAlgorithm algo) { switch (algo) @@ -68,7 +68,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, RadixSortStoreAlgorithm al return os << "(algo) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED template (algo) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED //! @rst //! The BlockLoad class provides :ref:`collective ` data movement methods for loading a linear diff --git a/cub/cub/block/block_radix_rank.cuh b/cub/cub/block/block_radix_rank.cuh index ccec1f4d417..9b571f4227e 100644 --- a/cub/cub/block/block_radix_rank.cuh +++ b/cub/cub/block/block_radix_rank.cuh @@ -71,7 +71,7 @@ enum RadixRankAlgorithm RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR }; -#if !_CCCL_COMPILER(NVRTC) && !defined(_CCCL_DOXYGEN_INVOKED) +#if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED) inline ::std::ostream& operator<<(::std::ostream& os, RadixRankAlgorithm algo) { switch (algo) @@ -90,7 +90,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, RadixRankAlgorithm algo) return os << "(algo) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED /** Empty callback implementation */ template diff --git a/cub/cub/block/block_reduce.cuh b/cub/cub/block/block_reduce.cuh index 30186917294..10526c2cbba 100644 --- a/cub/cub/block/block_reduce.cuh +++ b/cub/cub/block/block_reduce.cuh @@ -149,7 +149,7 @@ enum BlockReduceAlgorithm BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC, }; -#if !_CCCL_COMPILER(NVRTC) && !defined(_CCCL_DOXYGEN_INVOKED) +#if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED) inline ::std::ostream& operator<<(::std::ostream& os, const BlockReduceAlgorithm& alg) { switch (alg) @@ -166,7 +166,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, const BlockReduceAlgorithm return os << "(alg) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED //! @rst //! The BlockReduce class provides :ref:`collective ` methods for computing a diff --git a/cub/cub/block/block_scan.cuh b/cub/cub/block/block_scan.cuh index 59356997522..3adf33110de 100644 --- a/cub/cub/block/block_scan.cuh +++ b/cub/cub/block/block_scan.cuh @@ -100,7 +100,7 @@ enum BlockScanAlgorithm BLOCK_SCAN_WARP_SCANS, }; -#if !_CCCL_COMPILER(NVRTC) && !defined(_CCCL_DOXYGEN_INVOKED) +#if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED) inline ::std::ostream& operator<<(::std::ostream& os, BlockScanAlgorithm algo) { switch (algo) @@ -115,7 +115,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, BlockScanAlgorithm algo) return os << "(algo) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED //! @rst //! The BlockScan class provides :ref:`collective ` methods for computing a parallel prefix diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh index 242a884ba30..cf039baaee3 100644 --- a/cub/cub/block/block_store.cuh +++ b/cub/cub/block/block_store.cuh @@ -542,7 +542,7 @@ enum BlockStoreAlgorithm BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED, }; -#if !_CCCL_COMPILER(NVRTC) && !defined(_CCCL_DOXYGEN_INVOKED) +#if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED) inline ::std::ostream& operator<<(::std::ostream& os, BlockStoreAlgorithm algo) { switch (algo) @@ -563,7 +563,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, BlockStoreAlgorithm algo) return os << "(algo) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED //! @rst //! The BlockStore class provides :ref:`collective ` data movement diff --git a/cub/cub/config.cuh b/cub/cub/config.cuh index a3bc844352b..72f4e423902 100644 --- a/cub/cub/config.cuh +++ b/cub/cub/config.cuh @@ -24,6 +24,6 @@ #include // IWYU pragma: export #include // IWYU pragma: export -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() diff --git a/cub/cub/detail/delay_constructor.cuh b/cub/cub/detail/delay_constructor.cuh index 63099c534bd..fd687ddc573 100644 --- a/cub/cub/detail/delay_constructor.cuh +++ b/cub/cub/detail/delay_constructor.cuh @@ -34,7 +34,7 @@ enum class delay_constructor_kind reduce_by_key }; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() inline ::std::ostream& operator<<(::std::ostream& os, delay_constructor_kind kind) { switch (kind) @@ -61,7 +61,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, delay_constructor_kind kin return os << "(kind) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() struct delay_constructor_policy { @@ -79,13 +79,13 @@ struct delay_constructor_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const delay_constructor_policy& p) { return os << "delay_constructor_policy { .kind = " << p.kind << ", .delay = " << p.delay << ", .l2_write_latency = " << p.l2_write_latency << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; template diff --git a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh index c0af951be57..5da50bc96b0 100644 --- a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh +++ b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh @@ -345,14 +345,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( } const adjacent_difference_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ ::std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceAdjacentDifference to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const int tile_size = active_policy.block_threads * active_policy.items_per_thread; const int num_tiles = static_cast(::cuda::ceil_div(num_items, tile_size)); diff --git a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh index 75166f1c482..ec52663441f 100644 --- a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh +++ b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh @@ -317,7 +317,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( } const batch_memcpy_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ ::std::stringstream ss; ss << active_policy; @@ -325,7 +325,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( static_cast(arch_id), ss.str().c_str()); })) -#endif +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) enum : uint32_t { diff --git a/cub/cub/device/dispatch/dispatch_find.cuh b/cub/cub/device/dispatch/dispatch_find.cuh index 9e96ff14f16..179a4bb0480 100644 --- a/cub/cub/device/dispatch/dispatch_find.cuh +++ b/cub/cub/device/dispatch/dispatch_find.cuh @@ -114,13 +114,13 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( const find_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceFind to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const int tile_size = active_policy.block_threads * active_policy.items_per_thread; const int num_tiles = static_cast(::cuda::ceil_div(num_items, tile_size)); diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index 329aeee7a03..b75bda9e368 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -200,13 +200,13 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( const histogram_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceHistogram to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const auto init_kernel = kernel_source.template HistogramInitKernel(); auto sweep_kernel = [&] { diff --git a/cub/cub/device/dispatch/dispatch_merge.cuh b/cub/cub/device/dispatch/dispatch_merge.cuh index eab8436c847..0bc4cba33a4 100644 --- a/cub/cub/device/dispatch/dispatch_merge.cuh +++ b/cub/cub/device/dispatch/dispatch_merge.cuh @@ -217,14 +217,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( } return dispatch_arch(policy_selector, arch_id, [&](auto policy_getter) { -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ std::stringstream ss; ss << policy_getter(); _CubLog("Dispatching DeviceMerge to arch %d with tuning: %s\n", static_cast(arch_id), ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) static_assert(::cuda::std::is_empty_v); using AgentT = typename choose_merge_agent< diff --git a/cub/cub/device/dispatch/dispatch_merge_sort.cuh b/cub/cub/device/dispatch/dispatch_merge_sort.cuh index 172b38ceacb..fa64da0d11e 100644 --- a/cub/cub/device/dispatch/dispatch_merge_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_merge_sort.cuh @@ -507,13 +507,13 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( constexpr merge_sort_policy active_policy = vsmem_adapted_agents::policy; #endif // CUB_DEFINE_RUNTIME_POLICIES -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceMergeSort to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const auto tile_size = active_policy.items_per_tile(); const auto num_tiles = ::cuda::ceil_div(num_items, tile_size); diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh index 89c4ddea641..725d0ca458b 100644 --- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh @@ -1208,13 +1208,13 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( return error; } -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ std::stringstream ss; ss << policy_selector(arch_id); _CubLog("Dispatching DeviceReduce to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) return dispatch_arch(policy_selector, arch_id, [&](auto policy_getter) { return DispatchRadixSort{ diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh index b1943f4741a..7b681441bd5 100644 --- a/cub/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce.cuh @@ -780,13 +780,13 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( } const reduce_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceReduce to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) // Check for small, single tile size if (num_items diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index 40a11f826fe..26f13c557d6 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -710,7 +710,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch( } return detail::dispatch_arch(policy_selector, arch_id, [&](auto policy_getter) { -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ ::std::stringstream ss; ss << policy_getter(); @@ -718,7 +718,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch( static_cast(arch_id), ss.str().c_str()); })) -#endif +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const auto [block_threads, items_per_thread, vsmem_per_block] = determine_threads_items_vsmem< decltype(policy_getter), diff --git a/cub/cub/device/dispatch/dispatch_reduce_nondeterministic.cuh b/cub/cub/device/dispatch/dispatch_reduce_nondeterministic.cuh index 1898c59f03b..e31aa5a60d5 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_nondeterministic.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_nondeterministic.cuh @@ -184,14 +184,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch_nondeterministic( } const reduce_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceReduceNondeterministic to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) // No temp storage needed but keep API consistent if (d_temp_storage == nullptr) diff --git a/cub/cub/device/dispatch/dispatch_rle.cuh b/cub/cub/device/dispatch/dispatch_rle.cuh index 64b6b74bf5e..19edfa24902 100644 --- a/cub/cub/device/dispatch/dispatch_rle.cuh +++ b/cub/cub/device/dispatch/dispatch_rle.cuh @@ -679,14 +679,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch( } const non_trivial_runs::rle_non_trivial_runs_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ ::std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceRle to arch %d with tuning: %s\n", static_cast(arch_id), ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const int block_threads = active_policy.block_threads; const int items_per_thread = active_policy.items_per_thread; diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh index 5cbe630937c..0cdd2b5ebc8 100644 --- a/cub/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/cub/device/dispatch/dispatch_scan.cuh @@ -961,13 +961,13 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( return error; } -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ std::stringstream ss; ss << policy_selector(arch_id); _CubLog("Dispatching DeviceScan to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) struct fake_policy { diff --git a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh index 844d9eee14c..154f4533f77 100644 --- a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh @@ -642,7 +642,7 @@ struct DispatchScanByKey return error; } -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ ::std::stringstream ss; ss << policy_selector(arch_id); @@ -650,7 +650,7 @@ struct DispatchScanByKey static_cast(arch_id), ss.str().c_str()); })) -#endif +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const detail::scan_by_key::scan_by_key_policy active_policy = policy_selector(arch_id); @@ -737,14 +737,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( return error; } -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ ::std::stringstream ss; ss << policy_selector(arch_id); _CubLog("Dispatching DeviceScanByKey to arch %d with tuning: %s\n", static_cast(arch_id), ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) struct fake_policy { diff --git a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh index fc6461fad2d..06dbb0dbb5a 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh @@ -511,14 +511,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( } const segmented_reduce_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ ::std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceSegmentedReduce to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) // Compute segments_per_block based on max_segment_size hint int segments_per_block = 1; diff --git a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh index a000111cefa..9da47dd6f2f 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh @@ -1287,14 +1287,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( // a function CUB_DETAIL_CONSTEXPR_ISH const segmented_sort_policy active_policy = policy_getter(); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ ::std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceSegmentedSort to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const int radix_bits = active_policy.large_segment.radix_bits; diff --git a/cub/cub/device/dispatch/dispatch_select_if.cuh b/cub/cub/device/dispatch/dispatch_select_if.cuh index f4aa32a6ece..695ef7f4984 100644 --- a/cub/cub/device/dispatch/dispatch_select_if.cuh +++ b/cub/cub/device/dispatch/dispatch_select_if.cuh @@ -1107,14 +1107,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( return error; } -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ ::std::stringstream ss; ss << PolicySelector{}(arch_id); _CubLog("Dispatching DeviceSelectIf to arch %d with tuning: %s\n", static_cast(arch_id), ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) return dispatch_arch(policy_selector, arch_id, [&](auto policy_getter) { return dispatch_policy( diff --git a/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh index 0d8419a8d0c..f34655990f4 100644 --- a/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh @@ -75,7 +75,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch_streaming( const reduce_by_key_policy policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ ::std::stringstream ss; ss << policy; @@ -83,7 +83,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch_streaming( static_cast(arch_id), ss.str().c_str()); })) -#endif +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) using local_offset_t = ::cuda::std::int32_t; using global_offset_t = OffsetT; diff --git a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh index 4941d0505a5..be4c479948c 100644 --- a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh +++ b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh @@ -475,14 +475,14 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( const three_way_partition_policy active_policy = policy_selector(arch_id); -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET( NV_IS_HOST, ({ ::std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceThreeWayPartition to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) struct fake_hub { diff --git a/cub/cub/device/dispatch/dispatch_transform.cuh b/cub/cub/device/dispatch/dispatch_transform.cuh index 8eb89d34591..1635dbb1019 100644 --- a/cub/cub/device/dispatch/dispatch_transform.cuh +++ b/cub/cub/device/dispatch/dispatch_transform.cuh @@ -461,13 +461,13 @@ struct invoke_for_arch<::cuda::std::tuple, CUB_DETAIL_CONSTEXPR_ISH transform_policy active_policy = policy_getter(); const auto seq = ::cuda::std::index_sequence_for{}; -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ ::std::stringstream ss; ss << active_policy; _CubLog("Dispatching DeviceTransform to arch %d with tuning: %s\n", (int) arch_id, ss.str().c_str()); })) -#endif // !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) if CUB_DETAIL_CONSTEXPR_ISH (Algorithm::ublkcp == active_policy.algorithm) { diff --git a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh index 552c2117351..fedfa634f03 100644 --- a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh @@ -30,10 +30,7 @@ #include #include #include - -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) -# include -#endif +#include CUB_NAMESPACE_BEGIN @@ -511,7 +508,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( return error; } -#if !_CCCL_COMPILER(NVRTC) && defined(CUB_DEBUG_LOG) +#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) NV_IF_TARGET(NV_IS_HOST, ({ ::std::stringstream ss; ss << policy_selector(arch_id); @@ -519,7 +516,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch( static_cast(arch_id), ss.str().c_str()); })) -#endif +#endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) struct fake_policy { diff --git a/cub/cub/device/dispatch/kernels/kernel_segmented_sort.cuh b/cub/cub/device/dispatch/kernels/kernel_segmented_sort.cuh index 0b9ec3e5e22..f528d7089ef 100644 --- a/cub/cub/device/dispatch/kernels/kernel_segmented_sort.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_segmented_sort.cuh @@ -45,7 +45,7 @@ struct LargeSegmentsSelectorT , d_offset_begin(d_offset_begin) , d_offset_end(d_offset_end) {} -#endif +#endif // !_CCCL_COMPILER(NVRTC) _CCCL_DEVICE _CCCL_FORCEINLINE bool operator()(local_segment_index_t segment_id) const { @@ -70,7 +70,7 @@ struct SmallSegmentsSelectorT , d_offset_begin(d_offset_begin) , d_offset_end(d_offset_end) {} -#endif +#endif // !_CCCL_COMPILER(NVRTC) _CCCL_DEVICE _CCCL_FORCEINLINE bool operator()(local_segment_index_t segment_id) const { diff --git a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh index ea9e4c235fc..8108a4c38ea 100644 --- a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh @@ -46,14 +46,14 @@ struct adjacent_difference_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const adjacent_difference_policy& p) { return os << "adjacent_difference_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier << ", .store_algorithm = " << p.store_algorithm << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh b/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh index d4c192d4b26..5f5c63329e9 100644 --- a/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_batch_memcpy.cuh @@ -54,7 +54,7 @@ struct small_buffer_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const small_buffer_policy& policy) { return os @@ -65,7 +65,7 @@ struct small_buffer_policy << ", .block_level_threshold = " << policy.block_level_threshold << ", .buff_delay_constructor = " << policy.buff_delay_constructor << ", .block_delay_constructor = " << policy.block_delay_constructor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct large_buffer_policy @@ -85,13 +85,13 @@ struct large_buffer_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const large_buffer_policy& policy) { return os << "large_buffer_policy { .block_threads = " << policy.block_threads << ", .bytes_per_thread = " << policy.bytes_per_thread << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct batch_memcpy_policy @@ -111,13 +111,13 @@ struct batch_memcpy_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const batch_memcpy_policy& policy) { return os << "batch_memcpy_policy { .small_buffer = " << policy.small_buffer << ", .large_buffer = " << policy.large_buffer << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh b/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh index 164c5a08dd1..2c2cc208467 100644 --- a/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_batched_topk.cuh @@ -41,14 +41,14 @@ struct worker_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const worker_policy& p) { return os << "worker_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm << ", .store_algorithm = " << p.store_algorithm << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct batched_topk_policy @@ -67,7 +67,7 @@ struct batched_topk_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const batched_topk_policy& p) { os << "batched_topk_policy { .worker_per_segment_policies = { "; @@ -81,7 +81,7 @@ struct batched_topk_policy } return os << " } }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_find.cuh b/cub/cub/device/dispatch/tuning/tuning_find.cuh index f0651941f19..a93feb9bcee 100644 --- a/cub/cub/device/dispatch/tuning/tuning_find.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_find.cuh @@ -42,13 +42,13 @@ struct find_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const find_policy& p) { return os << "find_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .vector_load_length = " << p.vector_load_length << ", .load_modifier = " << p.load_modifier << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_for.cuh b/cub/cub/device/dispatch/tuning/tuning_for.cuh index c37657fb1c2..9eafcd90eec 100644 --- a/cub/cub/device/dispatch/tuning/tuning_for.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_for.cuh @@ -37,13 +37,13 @@ struct for_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const for_policy& policy) { return os << "for_policy { .block_threads = " << policy.block_threads << ", .items_per_thread = " << policy.items_per_thread << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh index f334fbb84fe..7343e074176 100644 --- a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh @@ -256,7 +256,7 @@ struct histogram_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const histogram_policy& p) { return os @@ -267,7 +267,7 @@ struct histogram_policy << ", .pdl_trigger_next_launch_in_init_kernel_max_bin_count = " << p.pdl_trigger_next_launch_in_init_kernel_max_bin_count << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_merge.cuh b/cub/cub/device/dispatch/tuning/tuning_merge.cuh index 97b825ea2d3..a82c0f20a7d 100644 --- a/cub/cub/device/dispatch/tuning/tuning_merge.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_merge.cuh @@ -46,14 +46,14 @@ struct merge_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const merge_policy& p) { return os << "merge_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .load_modifier = " << p.load_modifier << ", .store_algorithm = " << p.store_algorithm << ", .use_block_load_to_shared = " << p.use_block_load_to_shared << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh index e2b152fe183..2b5ef988e72 100644 --- a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh @@ -112,14 +112,14 @@ struct merge_sort_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const merge_sort_policy& p) { return os << "merge_sort_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier << ", .store_algorithm = " << p.store_algorithm << " }"; } -#endif +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh index 3591f464bf6..27a78576546 100644 --- a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh @@ -56,13 +56,13 @@ struct radix_sort_histogram_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const radix_sort_histogram_policy& p) { return os << "radix_sort_histogram_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .num_parts = " << p.num_parts << ", .radix_bits = " << p.radix_bits << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct radix_sort_exclusive_sum_policy @@ -82,13 +82,13 @@ struct radix_sort_exclusive_sum_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const radix_sort_exclusive_sum_policy& p) { return os << "radix_sort_exclusive_sum_policy { .block_threads = " << p.block_threads << ", .radix_bits = " << p.radix_bits << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct radix_sort_onesweep_policy @@ -116,7 +116,7 @@ struct radix_sort_onesweep_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const radix_sort_onesweep_policy& p) { return os @@ -125,7 +125,7 @@ struct radix_sort_onesweep_policy << ", .radix_bits = " << p.radix_bits << ", .rank_algorith = " << p.rank_algorith << ", .scan_algorithm = " << p.scan_algorithm << ", .store_algorithm = " << p.store_algorithm << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_API constexpr auto make_reg_scaled_radix_sort_onesweep_policy( @@ -174,7 +174,7 @@ struct radix_sort_downsweep_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const radix_sort_downsweep_policy& p) { return os @@ -183,7 +183,7 @@ struct radix_sort_downsweep_policy << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier << ", .rank_algorithm = " << p.rank_algorithm << ", .scan_algorithm = " << p.scan_algorithm << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_API constexpr auto make_reg_scaled_radix_sort_downsweep_policy( @@ -225,14 +225,14 @@ struct radix_sort_upsweep_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const radix_sort_upsweep_policy& p) { return os << "radix_sort_upsweep_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .radix_bits = " << p.radix_bits << ", .load_modifier = " << p.load_modifier << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_API constexpr auto make_reg_scaled_radix_sort_upsweep_policy( @@ -277,7 +277,7 @@ struct radix_sort_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const radix_sort_policy& p) { return os @@ -288,7 +288,7 @@ struct radix_sort_policy << ", .alt_upsweep = " << p.alt_upsweep << ", .single_tile = " << p.single_tile << ", .segmented = " << p.segmented << ", .alt_segmented = " << p.alt_segmented << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; // TODO(bgruber): remove for CCCL 4.0 when we drop the public radix sort dispatcher diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce.cuh index 24d811951fd..911c3e9307c 100644 --- a/cub/cub/device/dispatch/tuning/tuning_reduce.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_reduce.cuh @@ -47,14 +47,14 @@ struct agent_reduce_policy // equivalent of AgentReducePolicy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const agent_reduce_policy& p) { return os << "agent_reduce_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .vector_load_length = " << p.vector_load_length << ", .block_algorithm = " << p.block_algorithm << ", .load_modifier = " << p.load_modifier << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct reduce_policy @@ -74,13 +74,13 @@ struct reduce_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const reduce_policy& p) { return os << "reduce_policy { .reduce = " << p.reduce << ", .single_tile = " << p.single_tile << ", .reduce_nondeterministic = " << p.reduce_nondeterministic << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh index fc4e11193e2..758eebcc858 100644 --- a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh @@ -965,7 +965,7 @@ struct reduce_by_key_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const reduce_by_key_policy& p) { return os @@ -973,7 +973,7 @@ struct reduce_by_key_policy << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier << ", .scan_algorithm = " << p.scan_algorithm << ", .delay_constructor = " << p.delay_constructor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh b/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh index 150ec11024e..4e55538df00 100644 --- a/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh @@ -354,7 +354,7 @@ struct rle_non_trivial_runs_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const rle_non_trivial_runs_policy& p) { return os @@ -363,7 +363,7 @@ struct rle_non_trivial_runs_policy << ", .load_modifier = " << p.load_modifier << ", .store_with_time_slicing = " << p.store_with_time_slicing << ", .scan_algorithm = " << p.scan_algorithm << ", .delay_constructor = " << p.delay_constructor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_scan.cuh b/cub/cub/device/dispatch/tuning/tuning_scan.cuh index 793a6119653..67ae80301b0 100644 --- a/cub/cub/device/dispatch/tuning/tuning_scan.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_scan.cuh @@ -569,7 +569,7 @@ struct scan_lookback_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const scan_lookback_policy& p) { return os @@ -578,7 +578,7 @@ struct scan_lookback_policy << ", .load_modifier = " << p.load_modifier << ", .store_algorithm = " << p.store_algorithm << ", .scan_algorithm = " << p.scan_algorithm << ", .delay_constructor = " << p.delay_constructor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct scan_warpspeed_policy @@ -604,14 +604,14 @@ struct scan_warpspeed_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const scan_warpspeed_policy& p) { return os << "scan_warpspeed_policy { .num_reduce_and_scan_warps = " << p.num_reduce_and_scan_warps << ", .look_ahead_items_per_thread = " << p.look_ahead_items_per_thread << ", .items_per_thread = " << p.items_per_thread << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; enum class scan_algorithm @@ -620,7 +620,7 @@ enum class scan_algorithm warpspeed }; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() inline ::std::ostream& operator<<(::std::ostream& os, scan_algorithm algorithm) { switch (algorithm) @@ -633,7 +633,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, scan_algorithm algorithm) return os << "scan_algorithm::"; } } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() struct scan_policy { @@ -651,13 +651,13 @@ struct scan_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const scan_policy& p) { return os << "scan_policy { .algorithm = " << p.algorithm << ", .lookback = " << p.lookback << ", .warpspeed = " << p.warpspeed << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh index 2f51cc53698..0ec233f5348 100644 --- a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh @@ -57,7 +57,7 @@ struct scan_by_key_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const scan_by_key_policy& p) { return os @@ -66,7 +66,7 @@ struct scan_by_key_policy << ", .store_algorithm = " << p.store_algorithm << ", .scan_algorithm = " << p.scan_algorithm << ", .delay_constructor = " << p.delay_constructor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; enum class primitive_accum { diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh index 83c5700ab08..1f519710161 100644 --- a/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh @@ -51,14 +51,14 @@ struct warp_reduce_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const warp_reduce_policy& p) { return os << "warp_reduce_policy { .block_threads = " << p.block_threads << ", .warp_threads = " << p.warp_threads << ", .items_per_thread = " << p.items_per_thread << ", .vector_load_length = " << p.vector_load_length << ", .load_modifier = " << p.load_modifier << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct segmented_reduce_policy @@ -78,13 +78,13 @@ struct segmented_reduce_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const segmented_reduce_policy& p) { return os << "segmented_reduce_policy { .large_reduce = " << p.large_reduce << ", .small_reduce = " << p.small_reduce << ", .medium_reduce = " << p.medium_reduce << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh index bb81bbc72e6..72414d458ca 100644 --- a/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh @@ -53,7 +53,7 @@ struct segmented_radix_sort_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const segmented_radix_sort_policy& p) { return os @@ -62,7 +62,7 @@ struct segmented_radix_sort_policy << ", .load_modifier = " << p.load_modifier << ", .rank_algorithm = " << p.rank_algorithm << ", .scan_algorithm = " << p.scan_algorithm << ", .radix_bits = " << p.radix_bits << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct sub_warp_merge_sort_policy @@ -98,7 +98,7 @@ struct sub_warp_merge_sort_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const sub_warp_merge_sort_policy& p) { return os @@ -106,7 +106,7 @@ struct sub_warp_merge_sort_policy << ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier << ", .store_algorithm = " << p.store_algorithm << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct segmented_sort_policy @@ -129,14 +129,14 @@ struct segmented_sort_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const segmented_sort_policy& p) { return os << "segmented_sort_policy { .large_segment = " << p.large_segment << ", .small_segment = " << p.small_segment << ", .medium_segment = " << p.medium_segment << ", .partitioning_threshold = " << p.partitioning_threshold << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh index c78630fd899..e1743341b3e 100644 --- a/cub/cub/device/dispatch/tuning/tuning_select_if.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_select_if.cuh @@ -1612,7 +1612,7 @@ struct select_if_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const select_if_policy& p) { return os @@ -1620,7 +1620,7 @@ struct select_if_policy << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier << ", .scan_algorithm = " << p.scan_algorithm << ", .delay_constructor = " << p.delay_constructor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh b/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh index d9113c09009..ae1bbeb2526 100644 --- a/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_three_way_partition.cuh @@ -439,7 +439,7 @@ struct three_way_partition_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const three_way_partition_policy& policy) { return os @@ -448,7 +448,7 @@ struct three_way_partition_policy << ", .load_modifier = " << policy.load_modifier << ", .block_scan_algorithm = " << policy.block_scan_algorithm << ", .delay_constructor = " << policy.delay_constructor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_topk.cuh b/cub/cub/device/dispatch/tuning/tuning_topk.cuh index 13b9ce93fc5..dd5f4ded915 100644 --- a/cub/cub/device/dispatch/tuning/tuning_topk.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_topk.cuh @@ -66,14 +66,14 @@ struct topk_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const topk_policy& p) { return os << "topk_policy { .block_threads = " << p.block_threads << ", .items_per_thread = " << p.items_per_thread << ", .bits_per_pass = " << p.bits_per_pass << ", .load_algorithm = " << p.load_algorithm << ", .scan_algorithm = " << p.scan_algorithm << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_transform.cuh b/cub/cub/device/dispatch/tuning/tuning_transform.cuh index 88c24cb185a..b5902cf4863 100644 --- a/cub/cub/device/dispatch/tuning/tuning_transform.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_transform.cuh @@ -60,7 +60,7 @@ enum class Algorithm ublkcp }; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() inline ::std::ostream& operator<<(::std::ostream& os, const Algorithm& algorithm) { switch (algorithm) @@ -77,7 +77,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, const Algorithm& algorithm return os << "Algorithm::"; } } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() struct prefetch_policy { @@ -103,7 +103,7 @@ struct prefetch_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const prefetch_policy& policy) { return os @@ -112,7 +112,7 @@ struct prefetch_policy << ", .max_items_per_thread = " << policy.max_items_per_thread << ", .prefetch_byte_stride = " << policy.prefetch_byte_stride << ", .unroll_factor = " << policy.unroll_factor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct vectorized_policy @@ -132,13 +132,13 @@ struct vectorized_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const vectorized_policy& policy) { return os << "vectorized_policy { .block_threads = " << policy.block_threads << ", .items_per_thread = " << policy.items_per_thread << ", .vec_size = " << policy.vec_size << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct async_copy_policy @@ -163,7 +163,7 @@ struct async_copy_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const async_copy_policy& policy) { return os @@ -172,7 +172,7 @@ struct async_copy_policy << ", .max_items_per_thread = " << policy.max_items_per_thread << ", .unroll_factor = " << policy.unroll_factor << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; struct transform_policy @@ -194,14 +194,14 @@ struct transform_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const transform_policy& policy) { return os << "transform_policy { .min_bytes_in_flight = " << policy.min_bytes_in_flight << ", .algorithm = " << policy.algorithm << ", .prefetch = " << policy.prefetch << ", .vectorized = " << policy.vectorized << ", .async_copy = " << policy.async_copy << " }"; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; #if _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh index 6d24f010565..49b6639389a 100644 --- a/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh @@ -24,12 +24,9 @@ #include #include +#include #include -#if !_CCCL_COMPILER(NVRTC) -# include -#endif - CUB_NAMESPACE_BEGIN namespace detail::unique_by_key @@ -55,7 +52,7 @@ struct unique_by_key_policy return !(lhs == rhs); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const unique_by_key_policy& p) { return os @@ -63,7 +60,7 @@ struct unique_by_key_policy << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier << ", .scan_algorithm = " << p.scan_algorithm << ", .delay_constructor = " << p.delay_constructor << " }"; } -#endif +#endif // _CCCL_HOSTED() }; enum class primitive_key diff --git a/cub/cub/iterator/arg_index_input_iterator.cuh b/cub/cub/iterator/arg_index_input_iterator.cuh index db2a65fa0ba..800bb9f4ddb 100644 --- a/cub/cub/iterator/arg_index_input_iterator.cuh +++ b/cub/cub/iterator/arg_index_input_iterator.cuh @@ -222,12 +222,12 @@ public: offset = 0; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const self_type& /*itr*/) { return os; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; CUB_NAMESPACE_END diff --git a/cub/cub/iterator/cache_modified_input_iterator.cuh b/cub/cub/iterator/cache_modified_input_iterator.cuh index 62c8c0e1e87..e8055488c1a 100644 --- a/cub/cub/iterator/cache_modified_input_iterator.cuh +++ b/cub/cub/iterator/cache_modified_input_iterator.cuh @@ -104,13 +104,13 @@ public: /// The type of a reference to an element the iterator can point to using reference = ValueType; -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() using iterator_category = ::cuda::std::random_access_iterator_tag; -#else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ // vvv !_CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_FREESTANDING() ^^^ // vvv _CCCL_HOSTED() vvv using iterator_category = THRUST_NS_QUALIFIER::detail::iterator_facade_category_t; -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() public: /// Wrapped native pointer @@ -207,12 +207,12 @@ public: } /// ostream operator -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const self_type& /*itr*/) { return os; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; namespace detail diff --git a/cub/cub/iterator/cache_modified_output_iterator.cuh b/cub/cub/iterator/cache_modified_output_iterator.cuh index 9ed915c29b8..88efa7903f5 100644 --- a/cub/cub/iterator/cache_modified_output_iterator.cuh +++ b/cub/cub/iterator/cache_modified_output_iterator.cuh @@ -216,13 +216,13 @@ public: return (ptr != rhs.ptr); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /// ostream operator friend ::std::ostream& operator<<(::std::ostream& os, const self_type& itr) { return os; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; CUB_NAMESPACE_END diff --git a/cub/cub/iterator/tex_obj_input_iterator.cuh b/cub/cub/iterator/tex_obj_input_iterator.cuh index 80694cda2f9..13a152d52e8 100644 --- a/cub/cub/iterator/tex_obj_input_iterator.cuh +++ b/cub/cub/iterator/tex_obj_input_iterator.cuh @@ -130,7 +130,7 @@ public: /// Constructor _CCCL_FORCEINLINE TexObjInputIterator() = default; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /** * @brief Use this iterator to bind @p ptr with a texture reference * @@ -167,7 +167,7 @@ public: { return CubDebug(cudaDestroyTextureObject(tex_obj)); } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /// Postfix increment _CCCL_HOST_DEVICE _CCCL_FORCEINLINE self_type operator++(int) @@ -260,7 +260,7 @@ public: return ((ptr != rhs.ptr) || (tex_offset != rhs.tex_offset) || (tex_obj != rhs.tex_obj)); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /// ostream operator friend ::std::ostream& operator<<(::std::ostream& os, const self_type& itr) { @@ -268,7 +268,7 @@ public: << " )"; return os; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() private: // This is hoisted out of operator* because #pragma can't be used inside of diff --git a/cub/cub/thread/thread_load.cuh b/cub/cub/thread/thread_load.cuh index 6f635e01b5b..20419de1c74 100644 --- a/cub/cub/thread/thread_load.cuh +++ b/cub/cub/thread/thread_load.cuh @@ -46,7 +46,7 @@ enum CacheLoadModifier LOAD_VOLATILE, ///< Volatile (any memory space) }; -#if !_CCCL_COMPILER(NVRTC) && !defined(_CCCL_DOXYGEN_INVOKED) +#if _CCCL_HOSTED() && !defined(_CCCL_DOXYGEN_INVOKED) inline ::std::ostream& operator<<(::std::ostream& os, CacheLoadModifier modifier) { switch (modifier) @@ -69,7 +69,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, CacheLoadModifier modifier return os << "(modifier) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED //! @name Thread I/O (cache modified) //! @{ diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh index 279265beffb..0ace779912f 100644 --- a/cub/cub/util_device.cuh +++ b/cub/cub/util_device.cuh @@ -35,9 +35,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include // saves 146ms compile-time over (CCCL 3.1) -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh index 152409cbe7c..6026cca08af 100644 --- a/cub/cub/util_type.cuh +++ b/cub/cub/util_type.cuh @@ -725,12 +725,12 @@ struct KeyValuePair return (value != b.value) || (key != b.key); } -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() friend ::std::ostream& operator<<(::std::ostream& os, const KeyValuePair& pair) { return os << '(' << pair.key << ',' << pair.value << ')'; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() }; /** diff --git a/cub/cub/warp/warp_load.cuh b/cub/cub/warp/warp_load.cuh index 4b4c93e4f50..222565a6f3e 100644 --- a/cub/cub/warp/warp_load.cuh +++ b/cub/cub/warp/warp_load.cuh @@ -107,7 +107,7 @@ enum WarpLoadAlgorithm WARP_LOAD_TRANSPOSE }; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() inline ::std::ostream& operator<<(::std::ostream& os, WarpLoadAlgorithm algorithm) { switch (algorithm) @@ -124,7 +124,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, WarpLoadAlgorithm algorith return os << "(algorithm) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED //! @rst //! The WarpLoad class provides :ref:`collective ` data movement methods for diff --git a/cub/cub/warp/warp_store.cuh b/cub/cub/warp/warp_store.cuh index 94f30704a2d..7b008c83c4d 100644 --- a/cub/cub/warp/warp_store.cuh +++ b/cub/cub/warp/warp_store.cuh @@ -111,7 +111,7 @@ enum WarpStoreAlgorithm WARP_STORE_TRANSPOSE }; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() inline ::std::ostream& operator<<(::std::ostream& os, WarpStoreAlgorithm algorithm) { switch (algorithm) @@ -128,7 +128,7 @@ inline ::std::ostream& operator<<(::std::ostream& os, WarpStoreAlgorithm algorit return os << "(algorithm) << ">"; } } -#endif // !_CCCL_COMPILER(NVRTC) && !_CCCL_DOXYGEN_INVOKED +#endif // _CCCL_HOSTED() && !_CCCL_DOXYGEN_INVOKED //! @rst //! The WarpStore class provides :ref:`collective ` diff --git a/cudax/include/cuda/experimental/__execution/exception.cuh b/cudax/include/cuda/experimental/__execution/exception.cuh index 42b3052a024..4c72c36d547 100644 --- a/cudax/include/cuda/experimental/__execution/exception.cuh +++ b/cudax/include/cuda/experimental/__execution/exception.cuh @@ -26,15 +26,15 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include // IWYU pragma: keep -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() namespace cuda::experimental::execution { // Since there are no exceptions in device code, we provide a stub implementation of // std::exception_ptr and related functions. -#if _CCCL_COMPILER(NVRTC) || !_CCCL_HOST_COMPILATION() +#if _CCCL_FREESTANDING() || !_CCCL_HOST_COMPILATION() struct exception_ptr { @@ -100,15 +100,15 @@ public: _CCCL_THROW(::cuda::cuda_error, cudaErrorUnknown, "unknown exception"); } -// ^^^ _CCCL_COMPILER(NVRTC) || !_CCCL_HOST_COMPILATION() ^^^ +// ^^^ _CCCL_FREESTANDING() || !_CCCL_HOST_COMPILATION() ^^^ #else -// vvv !_CCCL_COMPILER(NVRTC) && _CCCL_HOST_COMPILATION() vvv +// vvv _CCCL_HOSTED() && _CCCL_HOST_COMPILATION() vvv using ::std::current_exception; using ::std::exception_ptr; using ::std::rethrow_exception; -#endif // !_CCCL_COMPILER(NVRTC) && _CCCL_HOST_COMPILATION() +#endif // _CCCL_HOSTED() && _CCCL_HOST_COMPILATION() } // namespace cuda::experimental::execution #endif // __CUDAX_EXECUTION_EXCEPTION diff --git a/libcudacxx/include/cuda/__complex/get_real_imag.h b/libcudacxx/include/cuda/__complex/get_real_imag.h index a014c70e123..7c79efc55a4 100644 --- a/libcudacxx/include/cuda/__complex/get_real_imag.h +++ b/libcudacxx/include/cuda/__complex/get_real_imag.h @@ -52,7 +52,7 @@ template return __c.imag(); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() // Unless `--expt-relaxed-constexpr` is specified, obtaining values from std::complex is not constexpr :( # if defined(__CUDACC_RELAXED_CONSTEXPR__) @@ -80,7 +80,7 @@ template return reinterpret_cast(__c)[1]; } # endif // ^^^ !__CUDACC_RELAXED_CONSTEXPR__ ^^^ -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() _CCCL_END_NAMESPACE_CUDA diff --git a/libcudacxx/include/cuda/__random/pcg_engine.h b/libcudacxx/include/cuda/__random/pcg_engine.h index 49991e4d261..d2492d4d3bd 100644 --- a/libcudacxx/include/cuda/__random/pcg_engine.h +++ b/libcudacxx/include/cuda/__random/pcg_engine.h @@ -313,7 +313,7 @@ class pcg64_engine } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_API friend ::std::basic_ostream<_CharT, _Traits>& @@ -364,7 +364,7 @@ class pcg64_engine return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; //! @class pcg64 diff --git a/libcudacxx/include/cuda/__utility/__basic_any/rtti.h b/libcudacxx/include/cuda/__utility/__basic_any/rtti.h index 0efcfaae23c..07f33cb88b0 100644 --- a/libcudacxx/include/cuda/__utility/__basic_any/rtti.h +++ b/libcudacxx/include/cuda/__utility/__basic_any/rtti.h @@ -31,9 +31,9 @@ #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include // IWYU pragma: keep (for std::bad_cast) -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index 3d94a117510..8fd3dc9b953 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -27,9 +27,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__cccl/host_std_lib.h b/libcudacxx/include/cuda/std/__cccl/host_std_lib.h index e27e820207a..3a24eb998f2 100644 --- a/libcudacxx/include/cuda/std/__cccl/host_std_lib.h +++ b/libcudacxx/include/cuda/std/__cccl/host_std_lib.h @@ -34,24 +34,37 @@ # include #endif // ^^^ __has_include() ^^^ +// Freestanding environment detection +// NVRTC is treated as freestanding since it has no access to the host standard library +#if defined(_CCCL_ENABLE_FREESTANDING) || _CCCL_COMPILER(NVRTC) +# define _CCCL_FREESTANDING() 1 +# define _CCCL_NO_TYPEID +#else +# define _CCCL_FREESTANDING() 0 +#endif + +#define _CCCL_HOSTED() (!_CCCL_FREESTANDING()) + #define _CCCL_HOST_STD_LIB_MAKE_VERSION(_MAJOR, _MINOR) ((_MAJOR) * 100 + (_MINOR)) #define _CCCL_HOST_STD_LIB(...) _CCCL_VERSION_COMPARE(_CCCL_HOST_STD_LIB_, _CCCL_HOST_STD_LIB_##__VA_ARGS__) -#if defined(_MSVC_STL_VERSION) -# undef _CCCL_HOST_STD_LIB_STL -# define _CCCL_HOST_STD_LIB_STL() (_MSVC_STL_VERSION, 0) -#elif defined(__GLIBCXX__) -# undef _CCCL_HOST_STD_LIB_LIBSTDCXX -# define _CCCL_HOST_STD_LIB_LIBSTDCXX() (_GLIBCXX_RELEASE, 0) -#elif defined(_LIBCPP_VERSION) -# undef _CCCL_HOST_STD_LIB_LIBCXX +#if _CCCL_HOSTED() +# if defined(_MSVC_STL_VERSION) +# undef _CCCL_HOST_STD_LIB_STL +# define _CCCL_HOST_STD_LIB_STL() (_MSVC_STL_VERSION, 0) +# elif defined(__GLIBCXX__) +# undef _CCCL_HOST_STD_LIB_LIBSTDCXX +# define _CCCL_HOST_STD_LIB_LIBSTDCXX() (_GLIBCXX_RELEASE, 0) +# elif defined(_LIBCPP_VERSION) +# undef _CCCL_HOST_STD_LIB_LIBCXX // since llvm-16, the version scheme has been changed from MMppp to MMmmpp -# if _LIBCPP_VERSION / 10000 < 2 -# define _CCCL_HOST_STD_LIB_LIBCXX() (_LIBCPP_VERSION / 1000, 0) -# else -# define _CCCL_HOST_STD_LIB_LIBCXX() (_LIBCPP_VERSION / 10000, (_LIBCPP_VERSION / 100) % 100) -# endif -#endif // ^^^ _LIBCPP_VERSION ^^^ +# if _LIBCPP_VERSION / 10000 < 2 +# define _CCCL_HOST_STD_LIB_LIBCXX() (_LIBCPP_VERSION / 1000, 0) +# else +# define _CCCL_HOST_STD_LIB_LIBCXX() (_LIBCPP_VERSION / 10000, (_LIBCPP_VERSION / 100) % 100) +# endif +# endif // ^^^ _LIBCPP_VERSION ^^^ +#endif // _CCCL_HOSTED() #define _CCCL_HAS_HOST_STD_LIB() \ (_CCCL_HOST_STD_LIB(LIBSTDCXX) || _CCCL_HOST_STD_LIB(LIBCXX) || _CCCL_HOST_STD_LIB(STL)) diff --git a/libcudacxx/include/cuda/std/__chrono/file_clock.h b/libcudacxx/include/cuda/std/__chrono/file_clock.h index b26a75d8450..52b533d1976 100644 --- a/libcudacxx/include/cuda/std/__chrono/file_clock.h +++ b/libcudacxx/include/cuda/std/__chrono/file_clock.h @@ -25,9 +25,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__chrono/steady_clock.h b/libcudacxx/include/cuda/std/__chrono/steady_clock.h index d22448a5fc4..0bedd4c14ea 100644 --- a/libcudacxx/include/cuda/std/__chrono/steady_clock.h +++ b/libcudacxx/include/cuda/std/__chrono/steady_clock.h @@ -26,9 +26,9 @@ # include # include -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() # include -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() # include diff --git a/libcudacxx/include/cuda/std/__chrono/system_clock.h b/libcudacxx/include/cuda/std/__chrono/system_clock.h index 0d73b9ffb71..14ef84777a4 100644 --- a/libcudacxx/include/cuda/std/__chrono/system_clock.h +++ b/libcudacxx/include/cuda/std/__chrono/system_clock.h @@ -25,9 +25,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__cmath/fpclassify.h b/libcudacxx/include/cuda/std/__cmath/fpclassify.h index 8b4b7881302..b48e44824fd 100644 --- a/libcudacxx/include/cuda/std/__cmath/fpclassify.h +++ b/libcudacxx/include/cuda/std/__cmath/fpclassify.h @@ -32,7 +32,7 @@ #include -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() # ifndef FP_NAN # define FP_NAN 0 # endif // ! FP_NAN @@ -48,7 +48,7 @@ # ifndef FP_NORMAL # define FP_NORMAL 4 # endif // ! FP_NORMAL -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() #ifndef FP_ILOGB0 # define FP_ILOGB0 (-INT_MAX - 1) diff --git a/libcudacxx/include/cuda/std/__complex/complex.h b/libcudacxx/include/cuda/std/__complex/complex.h index 28cea8b858a..05ec7d6493b 100644 --- a/libcudacxx/include/cuda/std/__complex/complex.h +++ b/libcudacxx/include/cuda/std/__complex/complex.h @@ -43,12 +43,12 @@ // Compatibility helpers for thrust to convert between `std::complex` and `cuda::std::complex` // todo: find a way to get rid of this include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include // for std::complex stream operators # define _LIBCUDACXX_ACCESS_STD_COMPLEX_REAL(__c) reinterpret_cast(__c)[0] # define _LIBCUDACXX_ACCESS_STD_COMPLEX_IMAG(__c) reinterpret_cast(__c)[1] -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include @@ -122,7 +122,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _LIBCUDACXX_COMPLEX_ALIGNAS complex return *this; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_API inline complex(const ::std::complex<_Up>& __other) : __re_(_LIBCUDACXX_ACCESS_STD_COMPLEX_REAL(__other)) @@ -141,7 +141,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _LIBCUDACXX_COMPLEX_ALIGNAS complex { return {__re_, __im_}; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() [[nodiscard]] _CCCL_API constexpr value_type real() const { @@ -571,7 +571,7 @@ template } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template [[nodiscard]] _CCCL_API constexpr bool operator==(const complex<_Tp>& __x, const ::std::complex<_Up>& __y) { @@ -599,7 +599,7 @@ template return !(__x == __y); } # endif // _CCCL_STD_VER <= 2017 -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // real @@ -657,7 +657,7 @@ template return 0; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template ::std::basic_istream<_CharT, _Traits>& operator>>(::std::basic_istream<_CharT, _Traits>& __is, complex<_Tp>& __x) { @@ -672,7 +672,7 @@ ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _ { return __os << static_cast<::std::complex<_Tp>>(__x); } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__complex/nvbf16.h b/libcudacxx/include/cuda/std/__complex/nvbf16.h index f1421315d4d..47777ac1add 100644 --- a/libcudacxx/include/cuda/std/__complex/nvbf16.h +++ b/libcudacxx/include/cuda/std/__complex/nvbf16.h @@ -32,9 +32,9 @@ # include // todo: find a way to get rid of this include -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() # include // for std::complex stream operators -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() # include @@ -153,7 +153,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__nv_bfloat162)) compl return *this; } -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() template _CCCL_API inline complex(const ::std::complex<_Up>& __other) noexcept : __repr_(_LIBCUDACXX_ACCESS_STD_COMPLEX_REAL(__other), _LIBCUDACXX_ACCESS_STD_COMPLEX_IMAG(__other)) @@ -171,7 +171,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__nv_bfloat162)) compl { return {__repr_.x, __repr_.y}; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() [[nodiscard]] _CCCL_API inline value_type real() const { @@ -303,7 +303,7 @@ struct __get_complex_impl<__nv_bfloat16> } }; -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() template ::std::basic_istream<_CharT, _Traits>& operator>>(::std::basic_istream<_CharT, _Traits>& __is, complex<__nv_bfloat16>& __x) @@ -320,7 +320,7 @@ operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const complex<__nv_bfloa { return __os << complex{__x}; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__complex/nvfp16.h b/libcudacxx/include/cuda/std/__complex/nvfp16.h index 8113f393332..562d8c00cc2 100644 --- a/libcudacxx/include/cuda/std/__complex/nvfp16.h +++ b/libcudacxx/include/cuda/std/__complex/nvfp16.h @@ -32,9 +32,9 @@ # include // todo: find a way to get rid of this include -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() # include // for std::complex stream operators -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() # include @@ -153,7 +153,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__half2)) complex<__ha return *this; } -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() template _CCCL_API inline complex(const ::std::complex<_Up>& __other) : __repr_(_LIBCUDACXX_ACCESS_STD_COMPLEX_REAL(__other), _LIBCUDACXX_ACCESS_STD_COMPLEX_IMAG(__other)) @@ -171,7 +171,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__half2)) complex<__ha { return {__repr_.x, __repr_.y}; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() [[nodiscard]] _CCCL_API inline value_type real() const { @@ -303,7 +303,7 @@ struct __get_complex_impl<__half> } }; -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() template ::std::basic_istream<_CharT, _Traits>& operator>>(::std::basic_istream<_CharT, _Traits>& __is, complex<__half>& __x) { @@ -319,7 +319,7 @@ operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const complex<__half>& _ { return __os << complex{__x}; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__cstddef/types.h b/libcudacxx/include/cuda/std/__cstddef/types.h index eb89069fe60..57208e293fb 100644 --- a/libcudacxx/include/cuda/std/__cstddef/types.h +++ b/libcudacxx/include/cuda/std/__cstddef/types.h @@ -22,24 +22,24 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#else +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv # if !defined(offsetof) # define offsetof(type, member) (::size_t) ((char*) &(((type*) 0)->member) - (char*) 0) # endif // !offsetof -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() #include _CCCL_BEGIN_NAMESPACE_CUDA_STD -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() using max_align_t = long double; -#else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ / vvv !_CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_FREESTANDING() ^^^ / vvv _CCCL_HOSTED() vvv // Re-use the compiler's max_align_t where possible. using ::max_align_t; -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() using nullptr_t = decltype(nullptr); using ::ptrdiff_t; diff --git a/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h b/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h index 1f156622e11..2affe6a3e8d 100644 --- a/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h +++ b/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h @@ -25,9 +25,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__cstdlib/div.h b/libcudacxx/include/cuda/std/__cstdlib/div.h index e24b980ec0b..36c15559c9d 100644 --- a/libcudacxx/include/cuda/std/__cstdlib/div.h +++ b/libcudacxx/include/cuda/std/__cstdlib/div.h @@ -21,9 +21,9 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include @@ -31,11 +31,11 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD // If available, use the host's div_t, ldiv_t, and lldiv_t types because the struct members order is // implementation-defined. -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() using ::div_t; using ::ldiv_t; using ::lldiv_t; -#else // ^^^ !_CCCL_COMPILER(NVRTC) / _CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv struct _CCCL_TYPE_VISIBILITY_DEFAULT div_t { int quot; @@ -53,7 +53,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT lldiv_t long long quot; long long rem; }; -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() [[nodiscard]] _CCCL_API constexpr div_t div(int __x, int __y) noexcept { diff --git a/libcudacxx/include/cuda/std/__cstdlib/malloc.h b/libcudacxx/include/cuda/std/__cstdlib/malloc.h index 2c1ba11162c..50d5a0b8e4d 100644 --- a/libcudacxx/include/cuda/std/__cstdlib/malloc.h +++ b/libcudacxx/include/cuda/std/__cstdlib/malloc.h @@ -25,9 +25,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__cstring/memcpy.h b/libcudacxx/include/cuda/std/__cstring/memcpy.h index b3ea5f59162..00bcc9e0ebc 100644 --- a/libcudacxx/include/cuda/std/__cstring/memcpy.h +++ b/libcudacxx/include/cuda/std/__cstring/memcpy.h @@ -23,9 +23,9 @@ #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__cstring/memset.h b/libcudacxx/include/cuda/std/__cstring/memset.h index fdbc107cecc..688ceb2f730 100644 --- a/libcudacxx/include/cuda/std/__cstring/memset.h +++ b/libcudacxx/include/cuda/std/__cstring/memset.h @@ -23,9 +23,9 @@ #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__exception/cuda_error.h b/libcudacxx/include/cuda/std/__exception/cuda_error.h index 9bcd47a6b0b..1aef652ad0e 100644 --- a/libcudacxx/include/cuda/std/__exception/cuda_error.h +++ b/libcudacxx/include/cuda/std/__exception/cuda_error.h @@ -26,9 +26,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include @@ -40,7 +40,7 @@ using __cuda_error_t = ::cudaError_t; using __cuda_error_t = int; #endif -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() namespace __detail { static char* __format_cuda_error( @@ -92,7 +92,7 @@ class cuda_error : public ::std::runtime_error private: __cuda_error_t __status_; }; -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() _CCCL_END_NAMESPACE_CUDA diff --git a/libcudacxx/include/cuda/std/__exception/exception_macros.h b/libcudacxx/include/cuda/std/__exception/exception_macros.h index 7959e504c6c..c435f062906 100644 --- a/libcudacxx/include/cuda/std/__exception/exception_macros.h +++ b/libcudacxx/include/cuda/std/__exception/exception_macros.h @@ -23,9 +23,9 @@ #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__exception/terminate.h b/libcudacxx/include/cuda/std/__exception/terminate.h index cde5924c431..32ee1465090 100644 --- a/libcudacxx/include/cuda/std/__exception/terminate.h +++ b/libcudacxx/include/cuda/std/__exception/terminate.h @@ -26,9 +26,9 @@ # include #endif // !_CCCL_TILE_COMPILATION() -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__format/format_error.h b/libcudacxx/include/cuda/std/__format/format_error.h index e4b90b4a237..98d3ed3d43e 100644 --- a/libcudacxx/include/cuda/std/__format/format_error.h +++ b/libcudacxx/include/cuda/std/__format/format_error.h @@ -30,7 +30,7 @@ #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() _CCCL_BEGIN_NAMESPACE_CUDA_STD_NOVERSION @@ -54,7 +54,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT format_error : public ::std::runtime_error _CCCL_END_NAMESPACE_CUDA_STD_NOVERSION -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() _CCCL_BEGIN_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__format/format_spec_parser.h b/libcudacxx/include/cuda/std/__format/format_spec_parser.h index 62bf3a5738a..eda403e128c 100644 --- a/libcudacxx/include/cuda/std/__format/format_spec_parser.h +++ b/libcudacxx/include/cuda/std/__format/format_spec_parser.h @@ -40,9 +40,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // This file contains the std-format-spec parser. // diff --git a/libcudacxx/include/cuda/std/__fwd/complex.h b/libcudacxx/include/cuda/std/__fwd/complex.h index acdd65bb33c..3b1a30881aa 100644 --- a/libcudacxx/include/cuda/std/__fwd/complex.h +++ b/libcudacxx/include/cuda/std/__fwd/complex.h @@ -50,10 +50,10 @@ template inline constexpr bool __is_std_complex_v = __is_std_complex_v<_Tp>; template inline constexpr bool __is_std_complex_v = __is_std_complex_v<_Tp>; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template inline constexpr bool __is_std_complex_v<::std::complex<_Tp>> = true; -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // __is_cuda_std_complex_v diff --git a/libcudacxx/include/cuda/std/__host_stdlib/algorithm b/libcudacxx/include/cuda/std/__host_stdlib/algorithm index 849bc5fc249..b951336fff2 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/algorithm +++ b/libcudacxx/include/cuda/std/__host_stdlib/algorithm @@ -27,10 +27,10 @@ // is defined only when CCCL is including an algorithms-related header, giving // the compiler a chance to detect and break the cycle of includes. -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # define THRUST_INCLUDING_ALGORITHMS_HEADER # include # undef THRUST_INCLUDING_ALGORITHMS_HEADER -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_ALGORITHM diff --git a/libcudacxx/include/cuda/std/__host_stdlib/istream b/libcudacxx/include/cuda/std/__host_stdlib/istream index 15fa36a9a66..e79be0a02be 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/istream +++ b/libcudacxx/include/cuda/std/__host_stdlib/istream @@ -21,8 +21,8 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_ISTREAM diff --git a/libcudacxx/include/cuda/std/__host_stdlib/math.h b/libcudacxx/include/cuda/std/__host_stdlib/math.h index 3e4d19c1149..d6396304e22 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/math.h +++ b/libcudacxx/include/cuda/std/__host_stdlib/math.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include // Standard C++ library comes with it's own C++ compatible header. However, if the include paths are jumbled, @@ -44,6 +44,6 @@ "libcu++ requires the C++ compatibility header, not the C header. Please, check your include paths." # endif // math functions defined as macros -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_MATH_H diff --git a/libcudacxx/include/cuda/std/__host_stdlib/memory b/libcudacxx/include/cuda/std/__host_stdlib/memory index 67eda6d13e4..119de005333 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/memory +++ b/libcudacxx/include/cuda/std/__host_stdlib/memory @@ -27,10 +27,10 @@ // is defined only when CCCL is including an algorithms-related header, giving // the compiler a chance to detect and break the cycle of includes. -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # define THRUST_INCLUDING_ALGORITHMS_HEADER # include # undef THRUST_INCLUDING_ALGORITHMS_HEADER -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_MEMORY diff --git a/libcudacxx/include/cuda/std/__host_stdlib/new b/libcudacxx/include/cuda/std/__host_stdlib/new index 45470d1ab91..fcd65d037ce 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/new +++ b/libcudacxx/include/cuda/std/__host_stdlib/new @@ -21,8 +21,8 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_NEW diff --git a/libcudacxx/include/cuda/std/__host_stdlib/numeric b/libcudacxx/include/cuda/std/__host_stdlib/numeric index e6a5c5cb2a2..5527e64d8da 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/numeric +++ b/libcudacxx/include/cuda/std/__host_stdlib/numeric @@ -27,10 +27,10 @@ // is defined only when CCCL is including an algorithms-related header, giving // the compiler a chance to detect and break the cycle of includes. -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # define THRUST_INCLUDING_ALGORITHMS_HEADER # include # undef THRUST_INCLUDING_ALGORITHMS_HEADER -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_NUMERIC diff --git a/libcudacxx/include/cuda/std/__host_stdlib/ostream b/libcudacxx/include/cuda/std/__host_stdlib/ostream index 6df28052af0..3c6bb37daed 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/ostream +++ b/libcudacxx/include/cuda/std/__host_stdlib/ostream @@ -21,8 +21,8 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_OSTREAM diff --git a/libcudacxx/include/cuda/std/__host_stdlib/sstream b/libcudacxx/include/cuda/std/__host_stdlib/sstream index cc35d28ea6f..d977062875a 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/sstream +++ b/libcudacxx/include/cuda/std/__host_stdlib/sstream @@ -21,8 +21,8 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_SSTREAM diff --git a/libcudacxx/include/cuda/std/__host_stdlib/stdexcept b/libcudacxx/include/cuda/std/__host_stdlib/stdexcept index 3801fc78151..053f822fde8 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/stdexcept +++ b/libcudacxx/include/cuda/std/__host_stdlib/stdexcept @@ -21,8 +21,8 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_STDEXCEPT diff --git a/libcudacxx/include/cuda/std/__host_stdlib/time.h b/libcudacxx/include/cuda/std/__host_stdlib/time.h index 538db257569..da8ef857c9f 100644 --- a/libcudacxx/include/cuda/std/__host_stdlib/time.h +++ b/libcudacxx/include/cuda/std/__host_stdlib/time.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include // Standard C++ library comes with it's own C++ compatible header. However, if the include paths are jumbled, @@ -34,6 +34,6 @@ "libcu++ requires the C++ compatibility header, not the C header. Please, check your include paths." # endif // math functions defined as macros -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___HOST_STDLIB_TIME_H diff --git a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h index 638dd44970a..f6a8c10b392 100644 --- a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h +++ b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h @@ -42,7 +42,7 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # if _CCCL_COMPILER(MSVC) # include // for ::std::input_iterator_tag # else // ^^^ _CCCL_COMPILER(MSVC) ^^^ / vvv !_CCCL_COMPILER(MSVC) vvv @@ -51,7 +51,7 @@ # ifdef _GLIBCXX_DEBUG # include -# endif +# endif // _GLIBCXX_DEBUG # if _CCCL_STD_VER >= 2020 # include @@ -73,7 +73,7 @@ struct __cccl_std_contiguous_iterator_tag_exists : __cccl_type_is_defined # endif // _CCCL_STD_VER >= 2020 -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include @@ -122,7 +122,7 @@ using iter_reference_t = enable_if_t<__dereferenceable<_Tp>, decltype(*declval<_ #endif // _CCCL_HAS_CONCEPTS() -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() struct _CCCL_TYPE_VISIBILITY_DEFAULT input_iterator_tag {}; @@ -137,7 +137,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT random_access_iterator_tag : public bidirec struct _CCCL_TYPE_VISIBILITY_DEFAULT contiguous_iterator_tag : public random_access_iterator_tag {}; -#else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ / vvv !_CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_FREESTANDING() ^^^ / vvv _CCCL_HOSTED() vvv using input_iterator_tag = ::std::input_iterator_tag; using output_iterator_tag = ::std::output_iterator_tag; @@ -157,7 +157,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT contiguous_iterator_tag : public random_acc {}; # endif // _CCCL_STD_VER <= 2017 -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() template struct __iter_traits_cache @@ -440,12 +440,12 @@ template struct __iterator_traits {}; -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() // We need to properly accept specializations of `std::iterator_traits` template <__specialized_from_std _Ip> struct __iterator_traits<_Ip> : public ::std::iterator_traits<_Ip> {}; -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() // [iterator.traits]/3.1 // If `I` has valid ([temp.deduct]) member types `difference-type`, `value-type`, `reference`, and @@ -712,11 +712,11 @@ template struct __iterator_traits {}; -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() template struct __iterator_traits<_Ip, enable_if_t<__specialized_from_std<_Ip>>> : public ::std::iterator_traits<_Ip> {}; -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() // [iterator.traits]/3.1 // If `I` has valid ([temp.deduct]) member types `difference-type`, `value-type`, `reference`, and diff --git a/libcudacxx/include/cuda/std/__memory/addressof.h b/libcudacxx/include/cuda/std/__memory/addressof.h index 5d0fbf68fd1..aa650837d4f 100644 --- a/libcudacxx/include/cuda/std/__memory/addressof.h +++ b/libcudacxx/include/cuda/std/__memory/addressof.h @@ -33,6 +33,12 @@ # define _CCCL_HAS_BUILTIN_STD_ADDRESSOF() 0 #endif // _CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION() +// We cannot use host features if we are building in freestanding +#if _CCCL_FREESTANDING() +# undef _CCCL_HAS_BUILTIN_STD_ADDRESSOF +# define _CCCL_HAS_BUILTIN_STD_ADDRESSOF() 0 +#endif // _CCCL_FREESTANDING() + // include minimal std:: headers #if _CCCL_HAS_BUILTIN_STD_ADDRESSOF() # if _CCCL_HOST_STD_LIB(LIBSTDCXX) && __has_include() diff --git a/libcudacxx/include/cuda/std/__pstl/adjacent_difference.h b/libcudacxx/include/cuda/std/__pstl/adjacent_difference.h index 302cd6f1815..03faf648b20 100644 --- a/libcudacxx/include/cuda/std/__pstl/adjacent_difference.h +++ b/libcudacxx/include/cuda/std/__pstl/adjacent_difference.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -92,6 +92,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_ADJACENT_DIFFERENCE_H diff --git a/libcudacxx/include/cuda/std/__pstl/adjacent_find.h b/libcudacxx/include/cuda/std/__pstl/adjacent_find.h index 87e810da9f7..1a998613e7f 100644 --- a/libcudacxx/include/cuda/std/__pstl/adjacent_find.h +++ b/libcudacxx/include/cuda/std/__pstl/adjacent_find.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -87,6 +87,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_ADJACENT_FIND_H diff --git a/libcudacxx/include/cuda/std/__pstl/all_of.h b/libcudacxx/include/cuda/std/__pstl/all_of.h index be5b5589c97..96bcdfa4c47 100644 --- a/libcudacxx/include/cuda/std/__pstl/all_of.h +++ b/libcudacxx/include/cuda/std/__pstl/all_of.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -82,6 +82,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_ALL_OF_H diff --git a/libcudacxx/include/cuda/std/__pstl/any_of.h b/libcudacxx/include/cuda/std/__pstl/any_of.h index 158370b5458..b7d1a867a94 100644 --- a/libcudacxx/include/cuda/std/__pstl/any_of.h +++ b/libcudacxx/include/cuda/std/__pstl/any_of.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -82,6 +82,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_ANY_OF_H diff --git a/libcudacxx/include/cuda/std/__pstl/copy.h b/libcudacxx/include/cuda/std/__pstl/copy.h index c6fe5e4e2ad..b890579b5c9 100644 --- a/libcudacxx/include/cuda/std/__pstl/copy.h +++ b/libcudacxx/include/cuda/std/__pstl/copy.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -77,6 +77,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_COPY_H diff --git a/libcudacxx/include/cuda/std/__pstl/copy_if.h b/libcudacxx/include/cuda/std/__pstl/copy_if.h index d6a52e0f50d..797c19854c0 100644 --- a/libcudacxx/include/cuda/std/__pstl/copy_if.h +++ b/libcudacxx/include/cuda/std/__pstl/copy_if.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -88,6 +88,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_COPY_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/copy_n.h b/libcudacxx/include/cuda/std/__pstl/copy_n.h index 3ce581130e0..a323043f8b7 100644 --- a/libcudacxx/include/cuda/std/__pstl/copy_n.h +++ b/libcudacxx/include/cuda/std/__pstl/copy_n.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -78,6 +78,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_COPY_N_H diff --git a/libcudacxx/include/cuda/std/__pstl/count.h b/libcudacxx/include/cuda/std/__pstl/count.h index 2687306e673..c0a1c9e2ed7 100644 --- a/libcudacxx/include/cuda/std/__pstl/count.h +++ b/libcudacxx/include/cuda/std/__pstl/count.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -89,6 +89,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_COUNT_H diff --git a/libcudacxx/include/cuda/std/__pstl/count_if.h b/libcudacxx/include/cuda/std/__pstl/count_if.h index d3964e28b13..957903ac897 100644 --- a/libcudacxx/include/cuda/std/__pstl/count_if.h +++ b/libcudacxx/include/cuda/std/__pstl/count_if.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -89,6 +89,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_COUNT_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/equal.h b/libcudacxx/include/cuda/std/__pstl/equal.h index 6c0ddc39f80..d2b1862d9b0 100644 --- a/libcudacxx/include/cuda/std/__pstl/equal.h +++ b/libcudacxx/include/cuda/std/__pstl/equal.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -151,6 +151,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_EQUAL_H diff --git a/libcudacxx/include/cuda/std/__pstl/exclusive_scan.h b/libcudacxx/include/cuda/std/__pstl/exclusive_scan.h index 21e3d0f4d9a..06e26f1b890 100644 --- a/libcudacxx/include/cuda/std/__pstl/exclusive_scan.h +++ b/libcudacxx/include/cuda/std/__pstl/exclusive_scan.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -121,6 +121,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_EXCLUSIVE_SCAN_H diff --git a/libcudacxx/include/cuda/std/__pstl/fill.h b/libcudacxx/include/cuda/std/__pstl/fill.h index 55b67a0f916..6d1f0f22154 100644 --- a/libcudacxx/include/cuda/std/__pstl/fill.h +++ b/libcudacxx/include/cuda/std/__pstl/fill.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -97,6 +97,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_FILL_H diff --git a/libcudacxx/include/cuda/std/__pstl/fill_n.h b/libcudacxx/include/cuda/std/__pstl/fill_n.h index 619c97411fd..e9134b647d0 100644 --- a/libcudacxx/include/cuda/std/__pstl/fill_n.h +++ b/libcudacxx/include/cuda/std/__pstl/fill_n.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -80,6 +80,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_FILL_N_H diff --git a/libcudacxx/include/cuda/std/__pstl/find.h b/libcudacxx/include/cuda/std/__pstl/find.h index 8cf7db2af60..84561b45937 100644 --- a/libcudacxx/include/cuda/std/__pstl/find.h +++ b/libcudacxx/include/cuda/std/__pstl/find.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -82,6 +82,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_FIND_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/find_if.h b/libcudacxx/include/cuda/std/__pstl/find_if.h index 1b1a422f060..16ec610e861 100644 --- a/libcudacxx/include/cuda/std/__pstl/find_if.h +++ b/libcudacxx/include/cuda/std/__pstl/find_if.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -78,6 +78,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_FIND_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/find_if_not.h b/libcudacxx/include/cuda/std/__pstl/find_if_not.h index 823b057053e..b4dc9b0e066 100644 --- a/libcudacxx/include/cuda/std/__pstl/find_if_not.h +++ b/libcudacxx/include/cuda/std/__pstl/find_if_not.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -80,6 +80,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_FIND_IF_NOT_H diff --git a/libcudacxx/include/cuda/std/__pstl/for_each.h b/libcudacxx/include/cuda/std/__pstl/for_each.h index b5e76fa1022..4ef4e70b326 100644 --- a/libcudacxx/include/cuda/std/__pstl/for_each.h +++ b/libcudacxx/include/cuda/std/__pstl/for_each.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -73,6 +73,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_FOR_EACH_H diff --git a/libcudacxx/include/cuda/std/__pstl/for_each_n.h b/libcudacxx/include/cuda/std/__pstl/for_each_n.h index 400d3283346..270259f1695 100644 --- a/libcudacxx/include/cuda/std/__pstl/for_each_n.h +++ b/libcudacxx/include/cuda/std/__pstl/for_each_n.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -72,6 +72,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_FOR_EACH_N_H diff --git a/libcudacxx/include/cuda/std/__pstl/generate.h b/libcudacxx/include/cuda/std/__pstl/generate.h index bcdc90f5687..57a890f260a 100644 --- a/libcudacxx/include/cuda/std/__pstl/generate.h +++ b/libcudacxx/include/cuda/std/__pstl/generate.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -83,6 +83,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_GENERATE_H diff --git a/libcudacxx/include/cuda/std/__pstl/generate_n.h b/libcudacxx/include/cuda/std/__pstl/generate_n.h index cf81ca47900..8b91e4d8cf4 100644 --- a/libcudacxx/include/cuda/std/__pstl/generate_n.h +++ b/libcudacxx/include/cuda/std/__pstl/generate_n.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -82,6 +82,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_GENERATE_N_H diff --git a/libcudacxx/include/cuda/std/__pstl/inclusive_scan.h b/libcudacxx/include/cuda/std/__pstl/inclusive_scan.h index fa1bcd11cbd..3e7bd919451 100644 --- a/libcudacxx/include/cuda/std/__pstl/inclusive_scan.h +++ b/libcudacxx/include/cuda/std/__pstl/inclusive_scan.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -160,6 +160,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_INCLUSIVE_SCAN_H diff --git a/libcudacxx/include/cuda/std/__pstl/is_partitioned.h b/libcudacxx/include/cuda/std/__pstl/is_partitioned.h index 29a19796e4e..0cfa5de2398 100644 --- a/libcudacxx/include/cuda/std/__pstl/is_partitioned.h +++ b/libcudacxx/include/cuda/std/__pstl/is_partitioned.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -101,6 +101,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_IS_PARTITIONED_H diff --git a/libcudacxx/include/cuda/std/__pstl/is_sorted.h b/libcudacxx/include/cuda/std/__pstl/is_sorted.h index 42309644435..c3e6a65e991 100644 --- a/libcudacxx/include/cuda/std/__pstl/is_sorted.h +++ b/libcudacxx/include/cuda/std/__pstl/is_sorted.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -88,6 +88,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_IS_SORTED_H diff --git a/libcudacxx/include/cuda/std/__pstl/is_sorted_until.h b/libcudacxx/include/cuda/std/__pstl/is_sorted_until.h index 84fab890090..228fccfb1a3 100644 --- a/libcudacxx/include/cuda/std/__pstl/is_sorted_until.h +++ b/libcudacxx/include/cuda/std/__pstl/is_sorted_until.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -90,6 +90,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_IS_SORTED_UNTIL_H diff --git a/libcudacxx/include/cuda/std/__pstl/merge.h b/libcudacxx/include/cuda/std/__pstl/merge.h index b38121a94bc..3e9e69e72a6 100644 --- a/libcudacxx/include/cuda/std/__pstl/merge.h +++ b/libcudacxx/include/cuda/std/__pstl/merge.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -124,6 +124,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_MERGE_H diff --git a/libcudacxx/include/cuda/std/__pstl/mismatch.h b/libcudacxx/include/cuda/std/__pstl/mismatch.h index d4991245f85..ccafa49aa72 100644 --- a/libcudacxx/include/cuda/std/__pstl/mismatch.h +++ b/libcudacxx/include/cuda/std/__pstl/mismatch.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -147,6 +147,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_MISMATCH_H diff --git a/libcudacxx/include/cuda/std/__pstl/none_of.h b/libcudacxx/include/cuda/std/__pstl/none_of.h index 9c22d006cd6..ceb1f7f5558 100644 --- a/libcudacxx/include/cuda/std/__pstl/none_of.h +++ b/libcudacxx/include/cuda/std/__pstl/none_of.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -82,6 +82,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_NONE_OF_H diff --git a/libcudacxx/include/cuda/std/__pstl/partition.h b/libcudacxx/include/cuda/std/__pstl/partition.h index 9c6f8263706..0d5d875b7b6 100644 --- a/libcudacxx/include/cuda/std/__pstl/partition.h +++ b/libcudacxx/include/cuda/std/__pstl/partition.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -84,6 +84,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_PARTITION_H diff --git a/libcudacxx/include/cuda/std/__pstl/partition_copy.h b/libcudacxx/include/cuda/std/__pstl/partition_copy.h index 228e11cb0c6..2ac5e214df8 100644 --- a/libcudacxx/include/cuda/std/__pstl/partition_copy.h +++ b/libcudacxx/include/cuda/std/__pstl/partition_copy.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -106,6 +106,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_PARTITION_COPY_H diff --git a/libcudacxx/include/cuda/std/__pstl/reduce.h b/libcudacxx/include/cuda/std/__pstl/reduce.h index e7df736e70d..802e7e1a65e 100644 --- a/libcudacxx/include/cuda/std/__pstl/reduce.h +++ b/libcudacxx/include/cuda/std/__pstl/reduce.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -119,6 +119,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REDUCE_H diff --git a/libcudacxx/include/cuda/std/__pstl/remove.h b/libcudacxx/include/cuda/std/__pstl/remove.h index 9ff73ca6e62..67e863d1f41 100644 --- a/libcudacxx/include/cuda/std/__pstl/remove.h +++ b/libcudacxx/include/cuda/std/__pstl/remove.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -95,6 +95,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REMOVE_H diff --git a/libcudacxx/include/cuda/std/__pstl/remove_copy.h b/libcudacxx/include/cuda/std/__pstl/remove_copy.h index e2940134e9e..a5864e3dcef 100644 --- a/libcudacxx/include/cuda/std/__pstl/remove_copy.h +++ b/libcudacxx/include/cuda/std/__pstl/remove_copy.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -88,6 +88,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REMOVE_COPY_H diff --git a/libcudacxx/include/cuda/std/__pstl/remove_copy_if.h b/libcudacxx/include/cuda/std/__pstl/remove_copy_if.h index a10deaf7adc..896bedc02a4 100644 --- a/libcudacxx/include/cuda/std/__pstl/remove_copy_if.h +++ b/libcudacxx/include/cuda/std/__pstl/remove_copy_if.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -94,6 +94,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REMOVE_COPY_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/remove_if.h b/libcudacxx/include/cuda/std/__pstl/remove_if.h index f5f9b356202..f48ff00476c 100644 --- a/libcudacxx/include/cuda/std/__pstl/remove_if.h +++ b/libcudacxx/include/cuda/std/__pstl/remove_if.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -81,6 +81,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REMOVE_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/replace.h b/libcudacxx/include/cuda/std/__pstl/replace.h index a622800ecf2..3b4f54035e4 100644 --- a/libcudacxx/include/cuda/std/__pstl/replace.h +++ b/libcudacxx/include/cuda/std/__pstl/replace.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -109,6 +109,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REPLACE_H diff --git a/libcudacxx/include/cuda/std/__pstl/replace_copy.h b/libcudacxx/include/cuda/std/__pstl/replace_copy.h index 22253531193..000949a247b 100644 --- a/libcudacxx/include/cuda/std/__pstl/replace_copy.h +++ b/libcudacxx/include/cuda/std/__pstl/replace_copy.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -112,6 +112,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REPLACE_COPY_H diff --git a/libcudacxx/include/cuda/std/__pstl/replace_copy_if.h b/libcudacxx/include/cuda/std/__pstl/replace_copy_if.h index 5aac8359540..60fc0f6f680 100644 --- a/libcudacxx/include/cuda/std/__pstl/replace_copy_if.h +++ b/libcudacxx/include/cuda/std/__pstl/replace_copy_if.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -118,6 +118,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REPLACE_COPY_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/replace_if.h b/libcudacxx/include/cuda/std/__pstl/replace_if.h index bb9eacbf1f8..29114a32165 100644 --- a/libcudacxx/include/cuda/std/__pstl/replace_if.h +++ b/libcudacxx/include/cuda/std/__pstl/replace_if.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -91,6 +91,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REPLACE_IF_H diff --git a/libcudacxx/include/cuda/std/__pstl/reverse.h b/libcudacxx/include/cuda/std/__pstl/reverse.h index 2b068365afa..f30663ac43c 100644 --- a/libcudacxx/include/cuda/std/__pstl/reverse.h +++ b/libcudacxx/include/cuda/std/__pstl/reverse.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -104,6 +104,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REVERSE_H diff --git a/libcudacxx/include/cuda/std/__pstl/reverse_copy.h b/libcudacxx/include/cuda/std/__pstl/reverse_copy.h index efe36e198b5..b2b7be55ea1 100644 --- a/libcudacxx/include/cuda/std/__pstl/reverse_copy.h +++ b/libcudacxx/include/cuda/std/__pstl/reverse_copy.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -80,6 +80,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_REVERSE_COPY_H diff --git a/libcudacxx/include/cuda/std/__pstl/rotate.h b/libcudacxx/include/cuda/std/__pstl/rotate.h index a31b4547847..d11f49a523a 100644 --- a/libcudacxx/include/cuda/std/__pstl/rotate.h +++ b/libcudacxx/include/cuda/std/__pstl/rotate.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -73,6 +73,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_ROTATE_H diff --git a/libcudacxx/include/cuda/std/__pstl/rotate_copy.h b/libcudacxx/include/cuda/std/__pstl/rotate_copy.h index b37cbcca5b0..977f86490a7 100644 --- a/libcudacxx/include/cuda/std/__pstl/rotate_copy.h +++ b/libcudacxx/include/cuda/std/__pstl/rotate_copy.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -94,6 +94,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_ROTATE_COPY_H diff --git a/libcudacxx/include/cuda/std/__pstl/shift_left.h b/libcudacxx/include/cuda/std/__pstl/shift_left.h index 8ca9a2d25b3..e68c03959aa 100644 --- a/libcudacxx/include/cuda/std/__pstl/shift_left.h +++ b/libcudacxx/include/cuda/std/__pstl/shift_left.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -80,6 +80,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_SHIFT_LEFT_H diff --git a/libcudacxx/include/cuda/std/__pstl/shift_right.h b/libcudacxx/include/cuda/std/__pstl/shift_right.h index 7822f352a54..2170fa91b1c 100644 --- a/libcudacxx/include/cuda/std/__pstl/shift_right.h +++ b/libcudacxx/include/cuda/std/__pstl/shift_right.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -80,6 +80,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_SHIFT_RIGHT_H diff --git a/libcudacxx/include/cuda/std/__pstl/swap_ranges.h b/libcudacxx/include/cuda/std/__pstl/swap_ranges.h index 10b9b7a7fd3..cd68015f971 100644 --- a/libcudacxx/include/cuda/std/__pstl/swap_ranges.h +++ b/libcudacxx/include/cuda/std/__pstl/swap_ranges.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -153,6 +153,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_SWAP_RANGES_H diff --git a/libcudacxx/include/cuda/std/__pstl/transform.h b/libcudacxx/include/cuda/std/__pstl/transform.h index dba98ef3dc8..97513f6570f 100644 --- a/libcudacxx/include/cuda/std/__pstl/transform.h +++ b/libcudacxx/include/cuda/std/__pstl/transform.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -145,6 +145,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_TRANSFORM_H diff --git a/libcudacxx/include/cuda/std/__pstl/transform_exclusive_scan.h b/libcudacxx/include/cuda/std/__pstl/transform_exclusive_scan.h index ca5871090f2..26629e11891 100644 --- a/libcudacxx/include/cuda/std/__pstl/transform_exclusive_scan.h +++ b/libcudacxx/include/cuda/std/__pstl/transform_exclusive_scan.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -117,6 +117,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_TRANSFORM_EXCLUSIVE_SCAN_H diff --git a/libcudacxx/include/cuda/std/__pstl/transform_inclusive_scan.h b/libcudacxx/include/cuda/std/__pstl/transform_inclusive_scan.h index 94d046ee72b..46744f635c8 100644 --- a/libcudacxx/include/cuda/std/__pstl/transform_inclusive_scan.h +++ b/libcudacxx/include/cuda/std/__pstl/transform_inclusive_scan.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -178,6 +178,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_TRANSFORM_INCLUSIVE_SCAN_H diff --git a/libcudacxx/include/cuda/std/__pstl/transform_reduce.h b/libcudacxx/include/cuda/std/__pstl/transform_reduce.h index ecbf3744e67..dfdee608d07 100644 --- a/libcudacxx/include/cuda/std/__pstl/transform_reduce.h +++ b/libcudacxx/include/cuda/std/__pstl/transform_reduce.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -165,6 +165,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_TRANSFORM_REDUCE_H diff --git a/libcudacxx/include/cuda/std/__pstl/unique.h b/libcudacxx/include/cuda/std/__pstl/unique.h index 3085b3426b7..dd8d092bbee 100644 --- a/libcudacxx/include/cuda/std/__pstl/unique.h +++ b/libcudacxx/include/cuda/std/__pstl/unique.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -79,6 +79,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_UNIQUE_H diff --git a/libcudacxx/include/cuda/std/__pstl/unique_copy.h b/libcudacxx/include/cuda/std/__pstl/unique_copy.h index 3986157cab7..fef15cd722e 100644 --- a/libcudacxx/include/cuda/std/__pstl/unique_copy.h +++ b/libcudacxx/include/cuda/std/__pstl/unique_copy.h @@ -21,7 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include @@ -99,6 +99,6 @@ _CCCL_END_NAMESPACE_CUDA_STD # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #endif // _CUDA_STD___PSTL_UNIQUE_COPY_H diff --git a/libcudacxx/include/cuda/std/__random/bernoulli_distribution.h b/libcudacxx/include/cuda/std/__random/bernoulli_distribution.h index 5bd4386cb36..a85a11ecdf2 100644 --- a/libcudacxx/include/cuda/std/__random/bernoulli_distribution.h +++ b/libcudacxx/include/cuda/std/__random/bernoulli_distribution.h @@ -128,7 +128,7 @@ class bernoulli_distribution return !(__x == __y); } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const bernoulli_distribution& __x) @@ -162,7 +162,7 @@ class bernoulli_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/binomial_distribution.h b/libcudacxx/include/cuda/std/__random/binomial_distribution.h index e1c61886980..6b254dd4b3d 100644 --- a/libcudacxx/include/cuda/std/__random/binomial_distribution.h +++ b/libcudacxx/include/cuda/std/__random/binomial_distribution.h @@ -212,7 +212,7 @@ class binomial_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const binomial_distribution& __x) @@ -246,7 +246,7 @@ class binomial_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/cauchy_distribution.h b/libcudacxx/include/cuda/std/__random/cauchy_distribution.h index 58dd8422ef5..9f520baad79 100644 --- a/libcudacxx/include/cuda/std/__random/cauchy_distribution.h +++ b/libcudacxx/include/cuda/std/__random/cauchy_distribution.h @@ -142,7 +142,7 @@ class cauchy_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const cauchy_distribution& __x) @@ -182,7 +182,7 @@ class cauchy_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/chi_squared_distribution.h b/libcudacxx/include/cuda/std/__random/chi_squared_distribution.h index 9283cb2f70c..297ece6a984 100644 --- a/libcudacxx/include/cuda/std/__random/chi_squared_distribution.h +++ b/libcudacxx/include/cuda/std/__random/chi_squared_distribution.h @@ -133,7 +133,7 @@ class chi_squared_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const chi_squared_distribution& __x) @@ -169,7 +169,7 @@ class chi_squared_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/exponential_distribution.h b/libcudacxx/include/cuda/std/__random/exponential_distribution.h index 771825fbd71..76b8269829b 100644 --- a/libcudacxx/include/cuda/std/__random/exponential_distribution.h +++ b/libcudacxx/include/cuda/std/__random/exponential_distribution.h @@ -138,7 +138,7 @@ class exponential_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const exponential_distribution& __x) @@ -175,7 +175,7 @@ class exponential_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/extreme_value_distribution.h b/libcudacxx/include/cuda/std/__random/extreme_value_distribution.h index 7d6a0305135..40a56eba571 100644 --- a/libcudacxx/include/cuda/std/__random/extreme_value_distribution.h +++ b/libcudacxx/include/cuda/std/__random/extreme_value_distribution.h @@ -145,7 +145,7 @@ class extreme_value_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const extreme_value_distribution& __x) @@ -184,7 +184,7 @@ class extreme_value_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h b/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h index 3f791f450af..64bc87a65f1 100644 --- a/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h +++ b/libcudacxx/include/cuda/std/__random/fisher_f_distribution.h @@ -145,7 +145,7 @@ class fisher_f_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const fisher_f_distribution& __x) @@ -184,7 +184,7 @@ class fisher_f_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/gamma_distribution.h b/libcudacxx/include/cuda/std/__random/gamma_distribution.h index 3c1073f5fc7..cf0a0845a9f 100644 --- a/libcudacxx/include/cuda/std/__random/gamma_distribution.h +++ b/libcudacxx/include/cuda/std/__random/gamma_distribution.h @@ -207,7 +207,7 @@ class gamma_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const gamma_distribution& __x) @@ -246,7 +246,7 @@ class gamma_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/geometric_distribution.h b/libcudacxx/include/cuda/std/__random/geometric_distribution.h index 414bcd98238..7be05298950 100644 --- a/libcudacxx/include/cuda/std/__random/geometric_distribution.h +++ b/libcudacxx/include/cuda/std/__random/geometric_distribution.h @@ -133,7 +133,7 @@ class geometric_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const geometric_distribution& __x) @@ -167,7 +167,7 @@ class geometric_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/linear_congruential_engine.h b/libcudacxx/include/cuda/std/__random/linear_congruential_engine.h index 80a001df401..4fc1eb82dd8 100644 --- a/libcudacxx/include/cuda/std/__random/linear_congruential_engine.h +++ b/libcudacxx/include/cuda/std/__random/linear_congruential_engine.h @@ -308,7 +308,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT linear_congruential_engine return !(__x == __y); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_API friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const linear_congruential_engine& __e) @@ -336,7 +336,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT linear_congruential_engine __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() private: _CCCL_API constexpr void seed(true_type, true_type, result_type __s) noexcept diff --git a/libcudacxx/include/cuda/std/__random/lognormal_distribution.h b/libcudacxx/include/cuda/std/__random/lognormal_distribution.h index 71c25447997..3e41a80b490 100644 --- a/libcudacxx/include/cuda/std/__random/lognormal_distribution.h +++ b/libcudacxx/include/cuda/std/__random/lognormal_distribution.h @@ -148,7 +148,7 @@ class lognormal_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const lognormal_distribution& __x) @@ -162,7 +162,7 @@ class lognormal_distribution { return __is >> __x.__nd_; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/negative_binomial_distribution.h b/libcudacxx/include/cuda/std/__random/negative_binomial_distribution.h index f635d1bb651..cc71f3f6e49 100644 --- a/libcudacxx/include/cuda/std/__random/negative_binomial_distribution.h +++ b/libcudacxx/include/cuda/std/__random/negative_binomial_distribution.h @@ -166,7 +166,7 @@ class negative_binomial_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const negative_binomial_distribution& __x) @@ -201,7 +201,7 @@ class negative_binomial_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/normal_distribution.h b/libcudacxx/include/cuda/std/__random/normal_distribution.h index e9989f9f193..83d7016b96f 100644 --- a/libcudacxx/include/cuda/std/__random/normal_distribution.h +++ b/libcudacxx/include/cuda/std/__random/normal_distribution.h @@ -173,7 +173,7 @@ class normal_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const normal_distribution& __x) @@ -220,7 +220,7 @@ class normal_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/philox_engine.h b/libcudacxx/include/cuda/std/__random/philox_engine.h index 8b205c52516..2e0a0c75d98 100644 --- a/libcudacxx/include/cuda/std/__random/philox_engine.h +++ b/libcudacxx/include/cuda/std/__random/philox_engine.h @@ -321,7 +321,7 @@ class philox_engine } #endif -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() //! This function streams a philox_engine to a std::basic_ostream. //! @param os The basic_ostream to stream out to. //! @param e The philox_engine to stream out. @@ -425,7 +425,7 @@ class philox_engine return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() private: _CCCL_API constexpr void __increment_counter() noexcept diff --git a/libcudacxx/include/cuda/std/__random/poisson_distribution.h b/libcudacxx/include/cuda/std/__random/poisson_distribution.h index 3b17683e3e4..6e85688eeb5 100644 --- a/libcudacxx/include/cuda/std/__random/poisson_distribution.h +++ b/libcudacxx/include/cuda/std/__random/poisson_distribution.h @@ -300,7 +300,7 @@ class poisson_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const poisson_distribution& __x) @@ -332,7 +332,7 @@ class poisson_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/student_t_distribution.h b/libcudacxx/include/cuda/std/__random/student_t_distribution.h index 56de187533f..00288b9f3d5 100644 --- a/libcudacxx/include/cuda/std/__random/student_t_distribution.h +++ b/libcudacxx/include/cuda/std/__random/student_t_distribution.h @@ -140,7 +140,7 @@ class student_t_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const student_t_distribution& __x) @@ -174,7 +174,7 @@ class student_t_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/uniform_int_distribution.h b/libcudacxx/include/cuda/std/__random/uniform_int_distribution.h index 5c5f85b02be..90733b76196 100644 --- a/libcudacxx/include/cuda/std/__random/uniform_int_distribution.h +++ b/libcudacxx/include/cuda/std/__random/uniform_int_distribution.h @@ -298,7 +298,7 @@ class uniform_int_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const uniform_int_distribution& __x) @@ -331,7 +331,7 @@ class uniform_int_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/uniform_real_distribution.h b/libcudacxx/include/cuda/std/__random/uniform_real_distribution.h index 5eb4393fa32..ea92d587945 100644 --- a/libcudacxx/include/cuda/std/__random/uniform_real_distribution.h +++ b/libcudacxx/include/cuda/std/__random/uniform_real_distribution.h @@ -145,7 +145,7 @@ class uniform_real_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const uniform_real_distribution& __x) @@ -180,7 +180,7 @@ class uniform_real_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__random/weibull_distribution.h b/libcudacxx/include/cuda/std/__random/weibull_distribution.h index d77b0d42b72..38abf896f81 100644 --- a/libcudacxx/include/cuda/std/__random/weibull_distribution.h +++ b/libcudacxx/include/cuda/std/__random/weibull_distribution.h @@ -142,7 +142,7 @@ class weibull_distribution } #endif // _CCCL_STD_VER <= 2017 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template friend ::std::basic_ostream<_CharT, _Traits>& operator<<(::std::basic_ostream<_CharT, _Traits>& __os, const weibull_distribution& __x) @@ -177,7 +177,7 @@ class weibull_distribution __is.flags(__flags); return __is; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__string/constexpr_c_functions.h b/libcudacxx/include/cuda/std/__string/constexpr_c_functions.h index 827103826c5..dc1219cb380 100644 --- a/libcudacxx/include/cuda/std/__string/constexpr_c_functions.h +++ b/libcudacxx/include/cuda/std/__string/constexpr_c_functions.h @@ -26,9 +26,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/__system_error/errc.h b/libcudacxx/include/cuda/std/__system_error/errc.h index 7fdc8cc466e..95a972127ba 100644 --- a/libcudacxx/include/cuda/std/__system_error/errc.h +++ b/libcudacxx/include/cuda/std/__system_error/errc.h @@ -21,17 +21,17 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include _CCCL_BEGIN_NAMESPACE_CUDA_STD -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() using ::std::errc; -#else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv enum class errc { invalid_argument = 22, @@ -42,7 +42,7 @@ enum class errc value_too_large = 75, # endif // ^^^ !_CCCL_OS(WINDOWS) ^^^ }; -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__tuple_dir/structured_bindings.h b/libcudacxx/include/cuda/std/__tuple_dir/structured_bindings.h index fcacc055ad1..86bb146387f 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/structured_bindings.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/structured_bindings.h @@ -47,7 +47,7 @@ _CCCL_BEGIN_NAMESPACE_STD template struct tuple_size; -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() template struct tuple_size< @@ -71,12 +71,12 @@ struct tuple_size< ::cuda::std::integral_constant)>>> : public ::cuda::std::integral_constant::value> {}; -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() template struct tuple_element; -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() template struct tuple_element<_Ip, const _Tp> { @@ -94,7 +94,7 @@ struct tuple_element<_Ip, const volatile _Tp> { using type _CCCL_NODEBUG_ALIAS = const volatile typename tuple_element<_Ip, _Tp>::type; }; -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() template struct tuple_size<::cuda::std::array<_Tp, _Size>> diff --git a/libcudacxx/include/cuda/std/__type_traits/is_primary_template.h b/libcudacxx/include/cuda/std/__type_traits/is_primary_template.h index 64def4d5b51..35945fac824 100644 --- a/libcudacxx/include/cuda/std/__type_traits/is_primary_template.h +++ b/libcudacxx/include/cuda/std/__type_traits/is_primary_template.h @@ -29,9 +29,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include @@ -57,9 +57,9 @@ using __is_primary_cccl_template = _IsValidExpansion<__test_for_primary_template #endif // !_CCCL_COMPILER(MSVC) -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() -// No ::std::traits with NVRTC +// No ::std::traits with freestanding implementations template struct __is_primary_std_template : true_type {}; @@ -67,7 +67,7 @@ struct __is_primary_std_template : true_type template using __select_traits = conditional_t<__is_primary_cccl_template<_Iter>::value, _OtherTraits, iterator_traits<_Iter>>; -#else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ / vvv !_CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_FREESTANDING() ^^^ / vvv _CCCL_HOSTED() vvv // We also need to respect what the user is defining to std::iterator_traits # if _CCCL_HOST_STD_LIB(LIBSTDCXX) @@ -112,7 +112,7 @@ using __select_traits = conditional_t<__is_primary_cccl_template<_Iter>::value, _OtherTraits, iterator_traits<_Iter>>, ::std::iterator_traits<_Iter>>; -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/__utility/as_const.h b/libcudacxx/include/cuda/std/__utility/as_const.h index db108573d9f..f58912717c0 100644 --- a/libcudacxx/include/cuda/std/__utility/as_const.h +++ b/libcudacxx/include/cuda/std/__utility/as_const.h @@ -34,13 +34,19 @@ # define _CCCL_HAS_BUILTIN_STD_AS_CONST() 0 #endif // _CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION() +// We cannot use host features if we are building in freestanding +#if _CCCL_FREESTANDING() +# undef _CCCL_HAS_BUILTIN_STD_AS_CONST +# define _CCCL_HAS_BUILTIN_STD_AS_CONST() 0 +#endif // _CCCL_FREESTANDING() + // include minimal std:: headers #if _CCCL_HAS_BUILTIN_STD_AS_CONST() # if _CCCL_HOST_STD_LIB(LIBCXX) && __has_include(<__utility/as_const.h>) # include <__utility/as_const.h> -# elif !_CCCL_COMPILER(NVRTC) +# elif _CCCL_HOSTED() # include -# endif +# endif // _CCCL_HOSTED() #endif // _CCCL_HAS_BUILTIN_STD_AS_CONST() #include diff --git a/libcudacxx/include/cuda/std/__utility/forward.h b/libcudacxx/include/cuda/std/__utility/forward.h index bc19c770417..1b91ca4ee99 100644 --- a/libcudacxx/include/cuda/std/__utility/forward.h +++ b/libcudacxx/include/cuda/std/__utility/forward.h @@ -38,15 +38,22 @@ # define _CCCL_HAS_BUILTIN_STD_FORWARD() 1 #endif // _CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION() +// We cannot use host features if we are building in freestanding +// However, NVRTC handles it specially +#if _CCCL_FREESTANDING() && !_CCCL_COMPILER(NVRTC) +# undef _CCCL_HAS_BUILTIN_STD_FORWARD +# define _CCCL_HAS_BUILTIN_STD_FORWARD() 0 +#endif // _CCCL_ENABLE_FREESTANDING + // include minimal std:: headers, nvcc in device mode doesn't need the std:: header #if _CCCL_HAS_BUILTIN_STD_FORWARD() && !(_CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION()) # if _CCCL_HOST_STD_LIB(LIBSTDCXX) && __has_include() # include # elif _CCCL_HOST_STD_LIB(LIBCXX) && __has_include(<__utility/forward.h>) # include <__utility/forward.h> -# elif !_CCCL_COMPILER(NVRTC) +# elif _CCCL_HOSTED() # include -# endif +# endif // _CCCL_HOSTED() #endif // _CCCL_HAS_BUILTIN_STD_FORWARD() && !(_CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION()) #include diff --git a/libcudacxx/include/cuda/std/__utility/forward_like.h b/libcudacxx/include/cuda/std/__utility/forward_like.h index f09fd5363c2..f462349207f 100644 --- a/libcudacxx/include/cuda/std/__utility/forward_like.h +++ b/libcudacxx/include/cuda/std/__utility/forward_like.h @@ -37,15 +37,21 @@ # define _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE() 0 #endif // _CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION() +// We cannot use host features if we are building in freestanding +#if _CCCL_FREESTANDING() +# undef _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE +# define _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE() 0 +#endif // _CCCL_FREESTANDING() + // include minimal std:: headers #if _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE() # if _CCCL_HOST_STD_LIB(LIBSTDCXX) && __has_include() # include # elif _CCCL_HOST_STD_LIB(LIBCXX) && __has_include(<__utility/forward_like.h>) # include <__utility/forward_like.h> -# elif !_CCCL_COMPILER(NVRTC) +# elif _CCCL_HOSTED() # include -# endif +# endif // _CCCL_HOSTED() #endif // _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE() #include diff --git a/libcudacxx/include/cuda/std/__utility/move.h b/libcudacxx/include/cuda/std/__utility/move.h index a7da248ff08..defff1bf8c4 100644 --- a/libcudacxx/include/cuda/std/__utility/move.h +++ b/libcudacxx/include/cuda/std/__utility/move.h @@ -39,6 +39,13 @@ # define _CCCL_HAS_BUILTIN_STD_MOVE() 1 #endif // _CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION() +// We cannot use host features if we are building in freestanding +// However, NVRTC handles it specially +#if _CCCL_FREESTANDING() && !_CCCL_COMPILER(NVRTC) +# undef _CCCL_HAS_BUILTIN_STD_MOVE +# define _CCCL_HAS_BUILTIN_STD_MOVE() 0 +#endif // _CCCL_ENABLE_FREESTANDING + #if _CCCL_COMPILER(CLANG, >=, 15) # define _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT() 1 #else // ^^^ has builtin std::move_if_noexcept ^^^ / vvv no builtin std::move_if_noexcept vvv @@ -51,15 +58,21 @@ # define _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT() 0 #endif // _CCCL_CUDA_COMPILER(NVCC) && _CCCL_DEVICE_COMPILATION() +// We cannot use host features if we are building in freestanding +#if _CCCL_FREESTANDING() +# undef _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT +# define _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT() 0 +#endif // _CCCL_FREESTANDING() + // include minimal std:: headers, nvcc in device mode doesn't need the std:: header #if _CCCL_HAS_BUILTIN_STD_MOVE() || _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT() # if _CCCL_HOST_STD_LIB(LIBSTDCXX) && __has_include() # include # elif _CCCL_HOST_STD_LIB(LIBCXX) && __has_include(<__utility/move.h>) # include <__utility/move.h> // includes std::move_if_noexcept, too -# elif !_CCCL_COMPILER(NVRTC) +# elif _CCCL_HOSTED() # include -# endif +# endif // _CCCL_HOSTED() #endif // _CCCL_HAS_BUILTIN_STD_MOVE() || _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT() // _CCCL_MOVE macro always expands to std::move builtin, if available, and fallbacks to cuda::std::move otherwise. diff --git a/libcudacxx/include/cuda/std/__utility/pair.h b/libcudacxx/include/cuda/std/__utility/pair.h index 954b06b8828..bd223fcee7a 100644 --- a/libcudacxx/include/cuda/std/__utility/pair.h +++ b/libcudacxx/include/cuda/std/__utility/pair.h @@ -330,7 +330,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> {} // std compatibility -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template , @@ -366,7 +366,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> is_nothrow_constructible_v<_T1, _U1> && is_nothrow_constructible_v<_T2, _U2>) : __base(::cuda::std::forward<_U1>(__p.first), ::cuda::std::forward<_U2>(__p.second)) {} -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // assignments _CCCL_HIDE_FROM_ABI pair& operator=(const pair&) = default; @@ -397,7 +397,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> } // std assignments -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template && is_copy_assignable_v<_T2>, int> = 0> _CCCL_HOST_API constexpr pair& operator=(::std::pair<_T1, _T2> const& __p) noexcept( is_nothrow_copy_assignable_v<_T1> && is_nothrow_copy_assignable_v<_T2>) @@ -415,7 +415,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> this->second = ::cuda::std::forward<_T2>(__p.second); return *this; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #if _CCCL_STD_VER >= 2023 _CCCL_API constexpr const pair& operator=(pair const& __p) const @@ -427,7 +427,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> return *this; } -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() _CCCL_HOST_API inline constexpr const pair& operator=(::std::pair<_T1, _T2> const& __p) const noexcept(is_nothrow_copy_assignable_v && is_nothrow_copy_assignable_v) requires(is_copy_assignable_v && is_copy_assignable_v) @@ -436,7 +436,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> this->second = __p.second; return *this; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() _CCCL_API constexpr const pair& operator=(pair&& __p) const noexcept(is_nothrow_assignable_v && is_nothrow_assignable_v) @@ -447,7 +447,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> return *this; } -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() _CCCL_HOST_API inline constexpr const pair& operator=(::std::pair<_T1, _T2>&& __p) const noexcept(is_nothrow_assignable_v && is_nothrow_assignable_v) requires(is_assignable_v && is_assignable_v) @@ -456,7 +456,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> this->second = ::cuda::std::forward<_T2>(__p.second); return *this; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() template _CCCL_API constexpr const pair& operator=(const pair<_U1, _U2>& __p) const @@ -467,7 +467,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> return *this; } -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() template _CCCL_HOST_API inline constexpr const pair& operator=(const ::std::pair<_U1, _U2>& __p) const requires(is_assignable_v && is_assignable_v) @@ -476,7 +476,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> this->second = __p.second; return *this; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() template _CCCL_API constexpr const pair& operator=(pair<_U1, _U2>&& __p) const @@ -487,7 +487,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> return *this; } -# if !_CCCL_COMPILER(NVRTC) +# if _CCCL_HOSTED() template _CCCL_HOST_API inline constexpr const pair& operator=(::std::pair<_U1, _U2>&& __p) const requires(is_assignable_v && is_assignable_v) @@ -496,7 +496,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> this->second = ::cuda::std::forward<_U2>(__p.second); return *this; } -# endif // !_CCCL_COMPILER(NVRTC) +# endif // _CCCL_HOSTED() #endif // _CCCL_STD_VER >= 2023 _CCCL_API inline _CCCL_CONSTEXPR_CXX20 void @@ -517,12 +517,12 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT pair : public __pair_base<_T1, _T2> } #endif // _CCCL_STD_VER >= 2023 -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() _CCCL_HOST_API constexpr operator ::std::pair<_T1, _T2>() const { return {this->first, this->second}; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() }; template diff --git a/libcudacxx/include/cuda/std/cassert b/libcudacxx/include/cuda/std/cassert index b8816743844..ace1d2a8091 100644 --- a/libcudacxx/include/cuda/std/cassert +++ b/libcudacxx/include/cuda/std/cassert @@ -21,8 +21,7 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif //_CCCL_COMPILER(NVRTC) - +#endif // _CCCL_HOSTED() #endif // _CUDA_STD_CASSERT diff --git a/libcudacxx/include/cuda/std/cfloat b/libcudacxx/include/cuda/std/cfloat index bf88acb5b27..2bd0431a6ff 100644 --- a/libcudacxx/include/cuda/std/cfloat +++ b/libcudacxx/include/cuda/std/cfloat @@ -21,9 +21,9 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv # define FLT_RADIX 2 # define DECIMAL_DIG DBL_DECIMAL_DIG # define FLT_ROUNDS 1 @@ -54,6 +54,6 @@ # define DBL_MAX_EXP 1024 # define DBL_MAX_10_EXP 308 # define DBL_HAS_SUBNORM 1 -#endif // ^^^ _CCCL_COMPILER(NVRTC) ^^^ +#endif // ^^^ _CCCL_FREESTANDING() ^^^ #endif // _CUDA_STD_CFLOAT diff --git a/libcudacxx/include/cuda/std/climits b/libcudacxx/include/cuda/std/climits index dc16a3a3229..c41351180f1 100644 --- a/libcudacxx/include/cuda/std/climits +++ b/libcudacxx/include/cuda/std/climits @@ -21,9 +21,9 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv # define _CCCL_CHAR_IS_UNSIGNED() ('\xff' > 0) // CURSED # define CHAR_BIT 8 @@ -56,6 +56,6 @@ # define LLONG_MIN (-LLONG_MAX - 1) # define LLONG_MAX 0x7fffffffffffffffLL # define ULLONG_MAX 0xffffffffffffffffUL -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() #endif // _CUDA_STD_CLIMITS diff --git a/libcudacxx/include/cuda/std/cmath b/libcudacxx/include/cuda/std/cmath index 1caa25f92dd..86504a6e97a 100644 --- a/libcudacxx/include/cuda/std/cmath +++ b/libcudacxx/include/cuda/std/cmath @@ -57,26 +57,26 @@ #include #include -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() # define INFINITY ::cuda::std::numeric_limits::infinity() # define NAN ::cuda::std::numeric_limits::quiet_NaN() -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() #include _CCCL_BEGIN_NAMESPACE_CUDA_STD -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() using double_t = double; using float_t = float; -#else // ^^^ _CCCL_COMPILER(NVRTC) ^^^ / vvv !_CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_FREESTANDING() ^^^ / vvv _CCCL_HOSTED() vvv using ::double_t; using ::float_t; -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() _CCCL_END_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/cstdint b/libcudacxx/include/cuda/std/cstdint index 51d63427d85..f027de19efd 100644 --- a/libcudacxx/include/cuda/std/cstdint +++ b/libcudacxx/include/cuda/std/cstdint @@ -23,9 +23,9 @@ #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv # include using int8_t = signed char; @@ -124,7 +124,7 @@ using uintmax_t = uint64_t; # define INTMAX_C(X) ((::intmax_t) (X)) # define UINTMAX_C(X) ((::uintmax_t) (X)) -#endif // ^^^ _CCCL_COMPILER(NVRTC) +#endif // ^^^ _CCCL_FREESTANDING() #include diff --git a/libcudacxx/include/cuda/std/cstring b/libcudacxx/include/cuda/std/cstring index 80995c01458..1aadb32c26c 100644 --- a/libcudacxx/include/cuda/std/cstring +++ b/libcudacxx/include/cuda/std/cstring @@ -25,9 +25,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include diff --git a/libcudacxx/include/cuda/std/ctime b/libcudacxx/include/cuda/std/ctime index 935b980c979..b5d09ed3023 100644 --- a/libcudacxx/include/cuda/std/ctime +++ b/libcudacxx/include/cuda/std/ctime @@ -29,7 +29,7 @@ #include -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() # define TIME_UTC 1 using time_t = long long int; @@ -39,7 +39,7 @@ struct timespec ::time_t tv_sec; long tv_nsec; }; -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() _CCCL_BEGIN_NAMESPACE_CUDA_STD diff --git a/libcudacxx/include/cuda/std/initializer_list b/libcudacxx/include/cuda/std/initializer_list index 29eca098795..47feb01c884 100644 --- a/libcudacxx/include/cuda/std/initializer_list +++ b/libcudacxx/include/cuda/std/initializer_list @@ -21,13 +21,13 @@ # pragma system_header #endif // no system header -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv # if !defined(__NV_BUILTIN_INITIALIZER_LIST) # error "libcu++ requires -builtin-initializer-list=true when compiling with NVRTC" # endif // !__NV_BUILTIN_INITIALIZER_LIST -#endif // ^^^ _CCCL_COMPILER(NVRTC) ^^^ +#endif // ^^^ _CCCL_FREESTANDING() ^^^ #include diff --git a/libcudacxx/include/cuda/std/string_view b/libcudacxx/include/cuda/std/string_view index 2b8348b341f..a320302bc07 100644 --- a/libcudacxx/include/cuda/std/string_view +++ b/libcudacxx/include/cuda/std/string_view @@ -61,12 +61,12 @@ #include // todo: find a way to get rid of these includes -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include // for std::basic_string operator<< # if __cpp_lib_string_view >= 201606L # include // for std::basic_string_view operator<< # endif // __cpp_lib_string_view >= 201606L -#endif // __cpp_lib_string_view >= 201606L +#endif // _CCCL_HOSTED() #include @@ -86,16 +86,16 @@ _CCCL_CONCEPT __cccl_basic_sv_compatible_range_check_traits = _CCCL_REQUIRES_EXP template _CCCL_CONCEPT __cccl_basic_sv_compatible_range_check_traits_std_ext = -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() _CCCL_REQUIRES_EXPR((_Range, _CharT, _Traits)) // ( // typename(typename _Range::traits_type), // requires(is_same_v>), // requires(is_same_v<_Traits, ::cuda::std::char_traits<_CharT>>) // ); -#else // ^^^ !_CCCL_COMPILER(NVRTC) ^^^ / vvv _CCCL_COMPILER(NVRTC) vvv +#else // ^^^ _CCCL_HOSTED() ^^^ / vvv _CCCL_FREESTANDING() vvv true; -#endif // ^^^ _CCCL_COMPILER(NVRTC) ^^^ +#endif // ^^^ _CCCL_FREESTANDING() ^^^ template _CCCL_CONCEPT __cccl_basic_sv_compatible_range = _CCCL_REQUIRES_EXPR((_Range, _CharT, _Traits)) // @@ -188,7 +188,7 @@ public: _CCCL_ASSERT((__end - __begin) >= 0, "string_view::string_view(iterator, sentinel) received invalid range"); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() _CCCL_TEMPLATE(class _Traits2, class _Alloc) _CCCL_REQUIRES(__cccl_basic_sv_is_std_to_cuda_std_char_traits<_Traits, _Traits2>) _CCCL_HOST_API constexpr basic_string_view(const ::std::basic_string<_CharT, _Traits2, _Alloc>& __sv) noexcept @@ -202,7 +202,7 @@ public: : __data_{__sv.data()} , __size_{__sv.size()} {} -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #if __cpp_lib_string_view >= 201606L _CCCL_TEMPLATE(class _Traits2) @@ -820,13 +820,13 @@ _CCCL_TEMPLATE(class _Range) _CCCL_REQUIRES(::cuda::std::ranges::contiguous_range<_Range>) _CCCL_HOST_DEVICE basic_string_view(_Range&&) -> basic_string_view<::cuda::std::ranges::range_value_t<_Range>>; -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template basic_string_view(::std::basic_string<_CharT, ::std::char_traits<_CharT>, _Alloc>) -> basic_string_view<_CharT>; template basic_string_view(::std::basic_string<_CharT, _Traits, _Alloc>) -> basic_string_view<_CharT, _Traits>; -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #if __cpp_lib_string_view >= 201606L template @@ -838,7 +838,7 @@ basic_string_view(::std::basic_string_view<_CharT, _Traits>) -> basic_string_vie // operator << -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_HOST_API ::std::basic_ostream<_CharT>& operator<<(::std::basic_ostream<_CharT>& __os, basic_string_view<_CharT> __str) @@ -861,7 +861,7 @@ operator<<(::std::basic_ostream<_CharT, _Traits>& __os, basic_string_view<_CharT return __os << ::std::basic_string<_CharT, _Traits, ::std::allocator<_CharT>>{__str.data(), __str.size()}; # endif // ^^^ __cpp_lib_string_view < 201606L ^^^ } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // literals diff --git a/libcudacxx/include/cuda/std/version b/libcudacxx/include/cuda/std/version index 02c5103b716..79e840ae028 100644 --- a/libcudacxx/include/cuda/std/version +++ b/libcudacxx/include/cuda/std/version @@ -25,9 +25,9 @@ // At the same time we want bring in all feature test macros from host #if __has_include() // should be the smallest include possible # include -#elif !_CCCL_COMPILER(NVRTC) +#elif _CCCL_HOSTED() # include // otherwise go for the smallest possible header -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #define __cccl_lib_bitops 201907L #define __cccl_lib_bool_constant 201505L diff --git a/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_complex.pass.cpp b/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_complex.pass.cpp index b0b4fa5d86c..3bf83bcc08a 100644 --- a/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_complex.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_complex.pass.cpp @@ -19,9 +19,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include "test_macros.h" @@ -75,13 +75,13 @@ TEST_FUNC void test_types() test_cccl_types(); static_assert(test_cccl_types()); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() // std::complex is not required to support other than standard floating-point types if constexpr (cuda::std::__is_std_fp_v) { NV_IF_TARGET(NV_IS_HOST, (test_constructor_from_complex(std::complex{U(1), U(2)});)) } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() } template diff --git a/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_tuple_like.pass.cpp b/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_tuple_like.pass.cpp index ecc4969bf1d..ec347581b94 100644 --- a/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_tuple_like.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/complex/complex.number/complex.members/constructors/from_tuple_like.pass.cpp @@ -21,9 +21,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include "test_macros.h" // vvv MyPair & tuple protocol for MyPair vvv @@ -181,7 +181,7 @@ TEST_FUNC constexpr bool test() return true; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template struct cuda::std::tuple_size> : cuda::std::integral_constant @@ -246,14 +246,14 @@ void test_host_types() # endif // _CCCL_HAS_INT128() } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() int main(int, char**) { test(); static_assert(test()); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() NV_IF_TARGET(NV_IS_HOST, (test_host_types();)) -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() return 0; } diff --git a/libcudacxx/test/libcudacxx/cuda/complex/complex.number/deduction.pass.cpp b/libcudacxx/test/libcudacxx/cuda/complex/complex.number/deduction.pass.cpp index f083005da3b..e44e749a4bd 100644 --- a/libcudacxx/test/libcudacxx/cuda/complex/complex.number/deduction.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/complex/complex.number/deduction.pass.cpp @@ -18,9 +18,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include "test_macros.h" @@ -52,7 +52,7 @@ TEST_FUNC void test_deduction() } // 5. Test cuda::complex(std::complex) -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() // std::complex is not required to support other than standard floating-point types if constexpr (cuda::std::__is_std_fp_v) { @@ -60,7 +60,7 @@ TEST_FUNC void test_deduction() ([[maybe_unused]] cuda::complex c{std::complex{}}; // assert((cuda::std::is_same_v) );)) } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // 6. Test cuda::complex(tuple-like) #if _CCCL_STD_VER >= 2020 diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/architecture.compile.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/architecture.compile.pass.cpp index 59ed65e3f25..9922d08e087 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/architecture.compile.pass.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/architecture.compile.pass.cpp @@ -24,7 +24,7 @@ # include # endif // _CCCL_ARCH(ARM64) && defined(__ARM_ACLE) # endif // !_CCCL_COMPILER(NVHPC) -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() int main(int, char**) { diff --git a/libcudacxx/test/libcudacxx/std/random/seed_seq/generate.pass.cpp b/libcudacxx/test/libcudacxx/std/random/seed_seq/generate.pass.cpp index d5e3b3b6ddb..2b4a7394adb 100644 --- a/libcudacxx/test/libcudacxx/std/random/seed_seq/generate.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/random/seed_seq/generate.pass.cpp @@ -12,10 +12,10 @@ // error: dynamic memory allocation is unsupported in tile code #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include "test_macros.h" @@ -29,7 +29,7 @@ TEST_FUNC TEST_CONSTEXPR_CXX20 bool test() return true; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() void test_against_std() { cuda::std::size_t n = 100; @@ -52,7 +52,7 @@ void test_against_std() assert(cuda_output == std_output); } } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() int main(int, char**) { diff --git a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.helper/tuple_size_structured_bindings.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.helper/tuple_size_structured_bindings.pass.cpp index 8e5dae73cea..3c7c5137710 100644 --- a/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.helper/tuple_size_structured_bindings.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/utilities/tuple/tuple.tuple/tuple.helper/tuple_size_structured_bindings.pass.cpp @@ -151,7 +151,7 @@ TEST_FUNC void test_after_tuple_size_specialization() auto& [p] = t; #if !_CCCL_COMPILER(NVRTC) // nvbug4053842 assert(p == -1); -#endif +#endif // !_CCCL_COMPILER(NVRTC) } int main(int, char**) diff --git a/libcudacxx/test/support/count_new.h b/libcudacxx/test/support/count_new.h index 9ab56e732a0..6aa4a955e7e 100644 --- a/libcudacxx/test/support/count_new.h +++ b/libcudacxx/test/support/count_new.h @@ -13,9 +13,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include "test_macros.h" diff --git a/libcudacxx/test/support/random_utilities/test_distribution.h b/libcudacxx/test/support/random_utilities/test_distribution.h index fd3f32868bc..e72684a2595 100644 --- a/libcudacxx/test/support/random_utilities/test_distribution.h +++ b/libcudacxx/test/support/random_utilities/test_distribution.h @@ -15,9 +15,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include "test_macros.h" @@ -65,7 +65,7 @@ TEST_FUNC constexpr bool test_get_param(Param param) return true; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template bool test_io(Param param) { diff --git a/libcudacxx/test/support/random_utilities/test_engine.h b/libcudacxx/test/support/random_utilities/test_engine.h index b89d00ee626..e3355ea04b6 100644 --- a/libcudacxx/test/support/random_utilities/test_engine.h +++ b/libcudacxx/test/support/random_utilities/test_engine.h @@ -7,14 +7,11 @@ // SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// -#if !_CCCL_COMPILER(NVRTC) -# include -#endif // !_CCCL_COMPILER(NVRTC) #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #include "test_macros.h" @@ -152,8 +149,7 @@ TEST_FUNC TEST_CONSTEXPR_CXX20 bool test_min_max() return true; } -#if !_CCCL_COMPILER(NVRTC) -# include +#if _CCCL_HOSTED() template void test_save_restore() { @@ -168,7 +164,7 @@ void test_save_restore() e1.discard(10000); assert(e0() == e1()); } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() template TEST_FUNC TEST_CONSTEXPR_CXX20 bool test_engine() diff --git a/thrust/thrust/complex.h b/thrust/thrust/complex.h index a88c244ba08..2a410b589ca 100644 --- a/thrust/thrust/complex.h +++ b/thrust/thrust/complex.h @@ -26,9 +26,9 @@ #include #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() #define THRUST_STD_COMPLEX_REAL(z) \ reinterpret_cast::value_type(&)[2]>(z)[0] @@ -119,7 +119,7 @@ struct complex template _CCCL_HOST_DEVICE complex(const complex& z); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /*! This converting copy constructor copies from a std::complex with * a type that is convertible to this \p complex's \c value_type. * @@ -144,7 +144,7 @@ struct complex */ template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex(const ::std::complex& z); -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /* --- Assignment Operators --- */ @@ -180,7 +180,7 @@ struct complex template _CCCL_HOST_DEVICE complex& operator=(const complex& z); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /*! Assign `z.real()` and `z.imag()` to the real and imaginary parts of this * \p complex respectively. * @@ -205,7 +205,7 @@ struct complex */ template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex& operator=(const ::std::complex& z); -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /* --- Compound Assignment Operators --- */ @@ -431,7 +431,7 @@ struct complex data.y = im; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /* --- Casting functions --- */ /*! Casts this \p complex to a std::complex of the same type. @@ -444,7 +444,7 @@ struct complex { return ::std::complex(real(), imag()); } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() private: struct alignas(sizeof(T) * 2) storage @@ -982,7 +982,7 @@ _CCCL_HOST_DEVICE complex asinh(const complex& z); template _CCCL_HOST_DEVICE complex atanh(const complex& z); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /* --- Stream Operators --- */ /*! Writes to an output stream a \p complex number in the form (real, imaginary). @@ -1015,7 +1015,7 @@ std::basic_ostream& operator<<(std::basic_ostream& */ template _CCCL_HOST ::std::basic_istream& operator>>(std::basic_istream& is, complex& z); -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /* --- Equality Operators --- */ @@ -1031,7 +1031,7 @@ _CCCL_HOST ::std::basic_istream& operator>>(std::basic_istream _CCCL_HOST_DEVICE bool operator==(const complex& x, const complex& y); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /*! Returns true if two \p complex numbers are equal and false otherwise. * * \param x The first \p complex. @@ -1055,7 +1055,7 @@ _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const complex& x, const */ template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const ::std::complex& x, const complex& y); -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /*! Returns true if the imaginary part of the \p complex number is zero and * the real part is equal to the scalar. Returns false otherwise. @@ -1095,7 +1095,7 @@ _CCCL_HOST_DEVICE bool operator==(const complex& x, const T1& y); template _CCCL_HOST_DEVICE bool operator!=(const complex& x, const complex& y); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /*! Returns true if two \p complex numbers are different and false otherwise. * * \param x The first \p complex. @@ -1119,7 +1119,7 @@ _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const complex& x, const */ template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const ::std::complex& x, const complex& y); -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /*! Returns true if the imaginary part of the \p complex number is not zero or * the real part is different from the scalar. Returns false otherwise. diff --git a/thrust/thrust/detail/complex/complex.inl b/thrust/thrust/detail/complex/complex.inl index ac028a10d7c..f74cfda1c54 100644 --- a/thrust/thrust/detail/complex/complex.inl +++ b/thrust/thrust/detail/complex/complex.inl @@ -36,7 +36,7 @@ _CCCL_HOST_DEVICE complex::complex(const complex& z) : data{T(z.real()), T(z.imag())} {} -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const ::std::complex& z) // Initialize the storage in the member initializer list using C++ unicorn @@ -52,7 +52,7 @@ _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex::complex(const ::std::complex // We do a functional-style cast here to suppress conversion warnings. : data{T(THRUST_STD_COMPLEX_REAL(z)), T(THRUST_STD_COMPLEX_IMAG(z))} {} -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /* --- Assignment Operators --- */ @@ -73,7 +73,7 @@ _CCCL_HOST_DEVICE complex& complex::operator=(const complex& z) return *this; } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex& complex::operator=(const ::std::complex& z) { @@ -90,7 +90,7 @@ _CCCL_HOST THRUST_STD_COMPLEX_DEVICE complex& complex::operator=(const ::s imag(T(THRUST_STD_COMPLEX_IMAG(z))); return *this; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() /* --- Compound Assignment Operators --- */ @@ -166,7 +166,7 @@ _CCCL_HOST_DEVICE bool operator==(const complex& x, const complex& y) return x.real() == y.real() && x.imag() == y.imag(); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const complex& x, const ::std::complex& y) { @@ -178,7 +178,7 @@ _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator==(const ::std::complex& x { return THRUST_STD_COMPLEX_REAL(x) == y.real() && THRUST_STD_COMPLEX_IMAG(x) == y.imag(); } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() template _CCCL_HOST_DEVICE bool operator==(const T0& x, const complex& y) @@ -198,7 +198,7 @@ _CCCL_HOST_DEVICE bool operator!=(const complex& x, const complex& y) return !(x == y); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const complex& x, const ::std::complex& y) { @@ -210,7 +210,7 @@ _CCCL_HOST THRUST_STD_COMPLEX_DEVICE bool operator!=(const ::std::complex& x { return !(x == y); } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() template _CCCL_HOST_DEVICE bool operator!=(const T0& x, const complex& y) diff --git a/thrust/thrust/detail/config.h b/thrust/thrust/detail/config.h index 640220eee28..dc2a3e603d1 100644 --- a/thrust/thrust/detail/config.h +++ b/thrust/thrust/detail/config.h @@ -19,6 +19,6 @@ #include // IWYU pragma: export -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() diff --git a/thrust/thrust/detail/pointer.h b/thrust/thrust/detail/pointer.h index 41dafe9a293..1543e2f1f61 100644 --- a/thrust/thrust/detail/pointer.h +++ b/thrust/thrust/detail/pointer.h @@ -265,14 +265,14 @@ class pointer : public detail::pointer_base::t return static_cast(::cuda::std::addressof(r)); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_HOST friend std::basic_ostream& operator<<(std::basic_ostream& os, const pointer& p) { return os << p.get(); } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // NOTE: This is needed so that Thrust smart pointers can be used in `std::unique_ptr`. _CCCL_HOST_DEVICE friend bool operator==(::cuda::std::nullptr_t, pointer p) diff --git a/thrust/thrust/iterator/detail/tuple_of_iterator_references.h b/thrust/thrust/iterator/detail/tuple_of_iterator_references.h index 07c164b4906..3904b10da8f 100644 --- a/thrust/thrust/iterator/detail/tuple_of_iterator_references.h +++ b/thrust/thrust/iterator/detail/tuple_of_iterator_references.h @@ -16,6 +16,7 @@ #include #include +#include #include #include #include @@ -206,18 +207,17 @@ struct tuple_size struct tuple_element> - : ::cuda::std::tuple_element> + : tuple_element> {}; _CCCL_END_NAMESPACE_CUDA_STD // structured bindings support -#if !_CCCL_COMPILER(NVRTC) namespace std { template struct tuple_size> - : integral_constant + : ::cuda::std::integral_constant {}; template @@ -225,4 +225,3 @@ struct tuple_element> {}; } // namespace std -#endif // !_CCCL_COMPILER(NVRTC) diff --git a/thrust/thrust/iterator/iterator_traits.h b/thrust/thrust/iterator/iterator_traits.h index 972496f69c7..1d9d411eab3 100644 --- a/thrust/thrust/iterator/iterator_traits.h +++ b/thrust/thrust/iterator/iterator_traits.h @@ -78,9 +78,9 @@ struct lazy_trait template using iterator_traits CCCL_DEPRECATED_BECAUSE("Use cuda::std::iterator_traits instead") = // FIXME(bgruber): switching to ::cuda::std::iterator_traits breaks some tests, e.g. cub.test.device_merge_sort.lid_1 -#if _CCCL_COMPILER(NVRTC) +#if _CCCL_FREESTANDING() ::cuda -#endif // _CCCL_COMPILER(NVRTC) +#endif // _CCCL_FREESTANDING() ::std::iterator_traits; _CCCL_SUPPRESS_DEPRECATED_PUSH diff --git a/thrust/thrust/system/cuda/detail/util.h b/thrust/thrust/system/cuda/detail/util.h index 24f6ce32821..2ddc702cf89 100644 --- a/thrust/thrust/system/cuda/detail/util.h +++ b/thrust/thrust/system/cuda/detail/util.h @@ -19,12 +19,12 @@ #include -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() # include # include # include -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() THRUST_NAMESPACE_BEGIN namespace cuda_cub @@ -111,7 +111,7 @@ _CCCL_HOST_DEVICE cudaError_t synchronize_optional(Policy& policy) return synchronize_stream_optional(derived_cast(policy)); } -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() template _CCCL_HOST_API _CCCL_FORCEINLINE cudaError_t trivial_copy_from_device(Type* dst, Type const* src, size_t count, cudaStream_t stream) @@ -158,7 +158,7 @@ trivial_copy_device_to_device(Policy& policy, Type* dst, Type const* src, size_t cuda_cub::synchronize_optional(policy); return status; } -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() _CCCL_HOST_DEVICE inline void throw_on_error(cudaError_t status) { diff --git a/thrust/thrust/system/error_code.h b/thrust/thrust/system/error_code.h index d5a75dd9b61..10c4c08293b 100644 --- a/thrust/thrust/system/error_code.h +++ b/thrust/thrust/system/error_code.h @@ -314,12 +314,12 @@ inline error_code make_error_code(errc::errc_t e); */ inline bool operator<(const error_code& lhs, const error_code& rhs); -#if !_CCCL_COMPILER(NVRTC) +#if _CCCL_HOSTED() /*! Effects: os << ec.category().name() << ':' << ec.value(). */ template std::basic_ostream& operator<<(std::basic_ostream& os, const error_code& ec); -#endif // !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HOSTED() // [19.5.3] class error_condition