diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh index 2776c601766..53cabe2c031 100644 --- a/cub/cub/util_device.cuh +++ b/cub/cub/util_device.cuh @@ -688,6 +688,6 @@ private: CUB_NAMESPACE_END -#if _CCCL_HAS_CUDA_COMPILER && !_CCCL_COMPILER(NVRTC) +#if _CCCL_HAS_CUDA_COMPILER() && !_CCCL_COMPILER(NVRTC) # include // to complete the definition of TripleChevronFactory -#endif // _CCCL_HAS_CUDA_COMPILER && !_CCCL_COMPILER(NVRTC) +#endif // _CCCL_HAS_CUDA_COMPILER() && !_CCCL_COMPILER(NVRTC) diff --git a/cudax/include/cuda/experimental/__detail/config.cuh b/cudax/include/cuda/experimental/__detail/config.cuh index 704b2e720a3..1ea10bea367 100644 --- a/cudax/include/cuda/experimental/__detail/config.cuh +++ b/cudax/include/cuda/experimental/__detail/config.cuh @@ -24,7 +24,7 @@ // Debuggers do not step into functions marked with __attribute__((__artificial__)). // This is useful for small wrapper functions that just dispatch to other functions and // that are inlined into the caller. -#if _CCCL_HAS_ATTRIBUTE(__artificial__) && !_CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_ATTRIBUTE(__artificial__) && !_CCCL_HAS_CUDA_COMPILER() # define _CUDAX_ARTIFICIAL __attribute__((__artificial__)) #else // ^^^ _CCCL_HAS_ATTRIBUTE(__artificial__) ^^^ / vvv !_CCCL_HAS_ATTRIBUTE(__artificial__) vvv # define _CUDAX_ARTIFICIAL diff --git a/docs/cccl/development/macro.rst b/docs/cccl/development/macro.rst index eea2c289d46..3404c5a3a12 100644 --- a/docs/cccl/development/macro.rst +++ b/docs/cccl/development/macro.rst @@ -59,7 +59,7 @@ The ``_CCCL_CUDA_COMPILER`` function-like macro can also be used to check the ve **CUDA identification/version macros**: +----------------------------------+-----------------------------+ -| ``_CCCL_HAS_CUDA_COMPILER`` | CUDA compiler is available | +| ``_CCCL_HAS_CUDA_COMPILER()`` | CUDA compiler is available | +----------------------------------+-----------------------------+ | ``_CCCL_CUDACC_BELOW(12, 7)`` | CUDA version below 12.7 | +----------------------------------+-----------------------------+ diff --git a/libcudacxx/codegen/generators/header.h b/libcudacxx/codegen/generators/header.h index 437a83261f6..f44444af8f6 100644 --- a/libcudacxx/codegen/generators/header.h +++ b/libcudacxx/codegen/generators/header.h @@ -57,7 +57,7 @@ inline void FormatHeader(std::ostream& out) _LIBCUDACXX_BEGIN_NAMESPACE_STD -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() )XXX"; out << header; @@ -66,7 +66,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD inline void FormatTail(std::ostream& out) { std::string tail = R"XXX( -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/__barrier/barrier_arrive_tx.h b/libcudacxx/include/cuda/__barrier/barrier_arrive_tx.h index e1ba83350ac..708658097cd 100644 --- a/libcudacxx/include/cuda/__barrier/barrier_arrive_tx.h +++ b/libcudacxx/include/cuda/__barrier/barrier_arrive_tx.h @@ -22,7 +22,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # if __cccl_ptx_isa >= 800 # include diff --git a/libcudacxx/include/cuda/__barrier/barrier_block_scope.h b/libcudacxx/include/cuda/__barrier/barrier_block_scope.h index af38ab97c8b..27845ca30e9 100644 --- a/libcudacxx/include/cuda/__barrier/barrier_block_scope.h +++ b/libcudacxx/include/cuda/__barrier/barrier_block_scope.h @@ -23,7 +23,7 @@ #include #include -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include # include diff --git a/libcudacxx/include/cuda/__barrier/barrier_expect_tx.h b/libcudacxx/include/cuda/__barrier/barrier_expect_tx.h index 1b8fc49d400..19ac3d9b3fd 100644 --- a/libcudacxx/include/cuda/__barrier/barrier_expect_tx.h +++ b/libcudacxx/include/cuda/__barrier/barrier_expect_tx.h @@ -22,7 +22,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # if __cccl_ptx_isa >= 800 # include diff --git a/libcudacxx/include/cuda/__barrier/barrier_native_handle.h b/libcudacxx/include/cuda/__barrier/barrier_native_handle.h index a685c832723..a12c5d68793 100644 --- a/libcudacxx/include/cuda/__barrier/barrier_native_handle.h +++ b/libcudacxx/include/cuda/__barrier/barrier_native_handle.h @@ -25,7 +25,7 @@ #include #include -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_BEGIN_NAMESPACE_CUDA_DEVICE @@ -36,6 +36,6 @@ _CCCL_DEVICE inline _CUDA_VSTD::uint64_t* barrier_native_handle(barrier _CCCL_NODISCARD _CCCL_HIDE_FROM_ABI _CCCL_DEVICE constexpr _Tp __bit_reverse_device(_Tp __value) noexcept @@ -92,7 +92,7 @@ _CCCL_NODISCARD _CCCL_HIDE_FROM_ABI _CCCL_DEVICE constexpr _Tp __bit_reverse_dev _CCCL_UNREACHABLE(); } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() template _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __bit_reverse_generic(_Tp __value) noexcept diff --git a/libcudacxx/include/cuda/__functional/for_each_canceled.h b/libcudacxx/include/cuda/__functional/for_each_canceled.h index 54e195c55e7..20762549e0e 100644 --- a/libcudacxx/include/cuda/__functional/for_each_canceled.h +++ b/libcudacxx/include/cuda/__functional/for_each_canceled.h @@ -28,7 +28,7 @@ #include -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_BEGIN_NAMESPACE_CUDA @@ -270,6 +270,6 @@ _CCCL_DEVICE _CCCL_HIDE_FROM_ABI void for_each_canceled_block(__UnaryFunction __ _LIBCUDACXX_END_NAMESPACE_CUDA -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() #endif // _CUDA__FUNCTIONAL_FOR_EACH_CANCELED_H diff --git a/libcudacxx/include/cuda/__functional/get_device_address.h b/libcudacxx/include/cuda/__functional/get_device_address.h index 77cdf4e296f..e10d76a62cb 100644 --- a/libcudacxx/include/cuda/__functional/get_device_address.h +++ b/libcudacxx/include/cuda/__functional/get_device_address.h @@ -38,7 +38,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CUDA template _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _Tp* get_device_address(_Tp& __device_object) { -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() NV_IF_ELSE_TARGET( NV_IS_DEVICE, (return _CUDA_VSTD::addressof(__device_object);), @@ -48,9 +48,9 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _Tp* get_device_address(_Tp& __device_ &__device_ptr, __device_object); return static_cast<_Tp*>(__device_ptr);)) -#else // ^^^ _CCCL_HAS_CUDA_COMPILER ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER vvv +#else // ^^^ _CCCL_HAS_CUDA_COMPILER() ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER() vvv return _CUDA_VSTD::addressof(__device_object); -#endif // !_CCCL_HAS_CUDA_COMPILER +#endif // !_CCCL_HAS_CUDA_COMPILER() } _LIBCUDACXX_END_NAMESPACE_CUDA diff --git a/libcudacxx/include/cuda/__memcpy_async/cp_async_bulk_shared_global.h b/libcudacxx/include/cuda/__memcpy_async/cp_async_bulk_shared_global.h index 513bfd387e2..c5900b66ed3 100644 --- a/libcudacxx/include/cuda/__memcpy_async/cp_async_bulk_shared_global.h +++ b/libcudacxx/include/cuda/__memcpy_async/cp_async_bulk_shared_global.h @@ -22,7 +22,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # if __cccl_ptx_isa >= 800 # include diff --git a/libcudacxx/include/cuda/__memcpy_async/cp_async_shared_global.h b/libcudacxx/include/cuda/__memcpy_async/cp_async_shared_global.h index 583f26dcd72..5ef8864d41e 100644 --- a/libcudacxx/include/cuda/__memcpy_async/cp_async_shared_global.h +++ b/libcudacxx/include/cuda/__memcpy_async/cp_async_shared_global.h @@ -22,7 +22,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/libcudacxx/include/cuda/__memcpy_async/memcpy_async.h b/libcudacxx/include/cuda/__memcpy_async/memcpy_async.h index 8afd043027f..2ff8468ec39 100644 --- a/libcudacxx/include/cuda/__memcpy_async/memcpy_async.h +++ b/libcudacxx/include/cuda/__memcpy_async/memcpy_async.h @@ -22,7 +22,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/libcudacxx/include/cuda/__memcpy_async/memcpy_async_tx.h b/libcudacxx/include/cuda/__memcpy_async/memcpy_async_tx.h index b7924a7721b..df4614f407e 100644 --- a/libcudacxx/include/cuda/__memcpy_async/memcpy_async_tx.h +++ b/libcudacxx/include/cuda/__memcpy_async/memcpy_async_tx.h @@ -22,7 +22,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # if __cccl_ptx_isa >= 800 # include diff --git a/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h b/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h index 3dd3e91d125..dde1ecd2d5f 100644 --- a/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h +++ b/libcudacxx/include/cuda/__memcpy_async/memcpy_completion.h @@ -31,7 +31,7 @@ #include #include -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include #endif // _CCCL_CUDA_COMPILER diff --git a/libcudacxx/include/cuda/__ptx/ptx_helper_functions.h b/libcudacxx/include/cuda/__ptx/ptx_helper_functions.h index 1b0dc58cd4f..06d28c60897 100644 --- a/libcudacxx/include/cuda/__ptx/ptx_helper_functions.h +++ b/libcudacxx/include/cuda/__ptx/ptx_helper_functions.h @@ -27,7 +27,7 @@ #include #include -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_BEGIN_NAMESPACE_CUDA_PTX @@ -142,6 +142,6 @@ inline _CCCL_DEVICE _B8 __u32_as_b8(uint32_t __u32) _LIBCUDACXX_END_NAMESPACE_CUDA_PTX -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() #endif // _CUDA_PTX_HELPER_FUNCTIONS_H_ diff --git a/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h b/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h index be06fbd34d1..7aa84b72889 100644 --- a/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h +++ b/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h @@ -33,7 +33,7 @@ // uses inline PTX to bypass __isLocal. _LIBCUDACXX_BEGIN_NAMESPACE_STD -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _CCCL_DEVICE inline bool __cuda_is_local(const volatile void* __ptr) { @@ -203,7 +203,7 @@ _CCCL_DEVICE bool __cuda_fetch_min_weak_if_local(volatile _Type* __ptr, _Type __ return __cuda_fetch_weak_if_local(__ptr, __val, __ret, __cuda_fetch_local_bop_min<_Type>); } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_derived.h b/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_derived.h index 19e89cad11f..c8a3563be86 100644 --- a/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_derived.h +++ b/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_derived.h @@ -29,7 +29,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() template using __cuda_atomic_enable_non_native_arithmetic = @@ -390,7 +390,7 @@ _CCCL_DEVICE static inline void __atomic_signal_fence_cuda(int) asm volatile("" ::: "memory"); } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated.h b/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated.h index 1211688a2dd..5ab554aff26 100644 --- a/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated.h +++ b/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated.h @@ -40,7 +40,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() static inline _CCCL_DEVICE void __cuda_atomic_membar(__thread_scope_block_tag) { asm volatile("membar.cta;" ::: "memory"); } @@ -3958,7 +3958,7 @@ static inline _CCCL_DEVICE _Type __atomic_fetch_sub_cuda(_Type volatile* __ptr, return __atomic_fetch_add_cuda(__ptr, -__op, __memorder, _Sco{}); } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__atomic/wait/notify_wait.h b/libcudacxx/include/cuda/std/__atomic/wait/notify_wait.h index ba9c4104d8b..acc1a942195 100644 --- a/libcudacxx/include/cuda/std/__atomic/wait/notify_wait.h +++ b/libcudacxx/include/cuda/std/__atomic/wait/notify_wait.h @@ -54,7 +54,7 @@ _LIBCUDACXX_HIDE_FROM_ABI void __atomic_notify_all(_Tp const volatile*, _Sco) template _LIBCUDACXX_HIDE_FROM_ABI bool __nonatomic_compare_equal(_Tp const& __lhs, _Tp const& __rhs) { -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() return __lhs == __rhs; #else return _CUDA_VSTD::memcmp(&__lhs, &__rhs, sizeof(_Tp)) == 0; diff --git a/libcudacxx/include/cuda/std/__cccl/assert.h b/libcudacxx/include/cuda/std/__cccl/assert.h index b8acc604a65..38eb06b582f 100644 --- a/libcudacxx/include/cuda/std/__cccl/assert.h +++ b/libcudacxx/include/cuda/std/__cccl/assert.h @@ -100,11 +100,11 @@ _CCCL_HOST_DEVICE _CCCL_BUILTIN_EXPECT(static_cast(expression), 1) \ ? (void) 0 : __assert_fail(message, __FILE__, __LINE__, __func__) # endif // !_CCCL_COMPILER(MSVC) -#elif _CCCL_HAS_CUDA_COMPILER +#elif _CCCL_HAS_CUDA_COMPILER() # define _CCCL_ASSERT_IMPL_DEVICE(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) -#else // ^^^ _CCCL_HAS_CUDA_COMPILER ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER vvv +#else // ^^^ _CCCL_HAS_CUDA_COMPILER() ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER() vvv # define _CCCL_ASSERT_IMPL_DEVICE(expression, message) ((void) 0) -#endif // !_CCCL_HAS_CUDA_COMPILER +#endif // !_CCCL_HAS_CUDA_COMPILER() //! _CCCL_ASSERT_HOST is enabled conditionally depending on CCCL_ENABLE_HOST_ASSERTIONS #ifdef CCCL_ENABLE_HOST_ASSERTIONS @@ -130,7 +130,7 @@ _CCCL_HOST_DEVICE # else # define _CCCL_ASSERT(expression, message) ((void) 0) # endif -#elif _CCCL_HAS_CUDA_COMPILER +#elif _CCCL_HAS_CUDA_COMPILER() # ifdef __CUDA_ARCH__ # define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL_DEVICE(expression, message) # define _CCCL_ASSERT(expression, message) _CCCL_ASSERT_DEVICE(expression, message) @@ -138,9 +138,9 @@ _CCCL_HOST_DEVICE # define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) # define _CCCL_ASSERT(expression, message) _CCCL_ASSERT_HOST(expression, message) # endif // !__CUDA_ARCH__ -#else // ^^^ _CCCL_HAS_CUDA_COMPILER ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER vvv +#else // ^^^ _CCCL_HAS_CUDA_COMPILER() ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER() vvv # define _CCCL_VERIFY(expression, message) _CCCL_ASSERT_IMPL_HOST(expression, message) # define _CCCL_ASSERT(expression, message) _CCCL_ASSERT_HOST(expression, message) -#endif // !_CCCL_HAS_CUDA_COMPILER +#endif // !_CCCL_HAS_CUDA_COMPILER() #endif // __CCCL_ASSERT_H diff --git a/libcudacxx/include/cuda/std/__cccl/builtin.h b/libcudacxx/include/cuda/std/__cccl/builtin.h index 59c17dddef6..cd45dd17479 100644 --- a/libcudacxx/include/cuda/std/__cccl/builtin.h +++ b/libcudacxx/include/cuda/std/__cccl/builtin.h @@ -239,19 +239,19 @@ # undef _CCCL_BUILTIN_BSWAP128 #endif // _CCCL_CUDA_COMPILER(NVCC) -#if _CCCL_CHECK_BUILTIN(builtin_bitreverse8) && !_CCCL_HAS_CUDA_COMPILER +#if _CCCL_CHECK_BUILTIN(builtin_bitreverse8) && !_CCCL_HAS_CUDA_COMPILER() # define _CCCL_BUILTIN_BITREVERSE8(...) __builtin_bitreverse8(__VA_ARGS__) #endif -#if _CCCL_CHECK_BUILTIN(builtin_bitreverse16) && !_CCCL_HAS_CUDA_COMPILER +#if _CCCL_CHECK_BUILTIN(builtin_bitreverse16) && !_CCCL_HAS_CUDA_COMPILER() # define _CCCL_BUILTIN_BITREVERSE16(...) __builtin_bitreverse16(__VA_ARGS__) #endif -#if _CCCL_CHECK_BUILTIN(builtin_bitreverse32) && !_CCCL_HAS_CUDA_COMPILER +#if _CCCL_CHECK_BUILTIN(builtin_bitreverse32) && !_CCCL_HAS_CUDA_COMPILER() # define _CCCL_BUILTIN_BITREVERSE32(...) __builtin_bitreverse32(__VA_ARGS__) #endif -#if _CCCL_CHECK_BUILTIN(builtin_bitreverse64) && !_CCCL_HAS_CUDA_COMPILER +#if _CCCL_CHECK_BUILTIN(builtin_bitreverse64) && !_CCCL_HAS_CUDA_COMPILER() # define _CCCL_BUILTIN_BITREVERSE64(...) __builtin_bitreverse64(__VA_ARGS__) #endif diff --git a/libcudacxx/include/cuda/std/__cccl/compiler.h b/libcudacxx/include/cuda/std/__cccl/compiler.h index db031419e69..e5ceea1ab63 100644 --- a/libcudacxx/include/cuda/std/__cccl/compiler.h +++ b/libcudacxx/include/cuda/std/__cccl/compiler.h @@ -139,12 +139,12 @@ #define _CCCL_CUDACC_AT_LEAST(...) _CCCL_VERSION_COMPARE(_CCCL_CUDACC_, _CCCL_CUDACC, >=, __VA_ARGS__) #if _CCCL_VERSION_IS_INVALID(_CCCL_CUDACC()) -# define _CCCL_HAS_CUDA_COMPILER 0 +# define _CCCL_HAS_CUDA_COMPILER() 0 #else // ^^^ has cuda compiler ^^^ / vvv no cuda compiler vvv -# define _CCCL_HAS_CUDA_COMPILER 1 +# define _CCCL_HAS_CUDA_COMPILER() 1 #endif // ^^^ no cuda compiler ^^^ -#if defined(_CCCL_HAS_CUDA_COMPILER) && _CCCL_CUDACC_BELOW(12) && !defined(CCCL_IGNORE_DEPRECATED_CUDA_BELOW_12) +#if _CCCL_HAS_CUDA_COMPILER() && _CCCL_CUDACC_BELOW(12) && !defined(CCCL_IGNORE_DEPRECATED_CUDA_BELOW_12) # error "CUDA versions below 12 are not supported." \ "Define CCCL_IGNORE_DEPRECATED_CUDA_BELOW_12 to suppress this message." #endif diff --git a/libcudacxx/include/cuda/std/__cccl/diagnostic.h b/libcudacxx/include/cuda/std/__cccl/diagnostic.h index 4248fb8c541..d8de8971d80 100644 --- a/libcudacxx/include/cuda/std/__cccl/diagnostic.h +++ b/libcudacxx/include/cuda/std/__cccl/diagnostic.h @@ -90,7 +90,7 @@ #endif // !_CCCL_COMPILER(CLANG) && !_CCCL_COMPILER(GCC) && !_CCCL_COMPILER(NVHPC) && !_CCCL_COMPILER(MSVC) // Enable us to selectively silence cuda compiler warnings -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # if _CCCL_CUDA_COMPILER(CLANG) # define _CCCL_NV_DIAG_SUPPRESS(_WARNING) # define _CCCL_NV_DIAG_DEFAULT(_WARNING) @@ -114,7 +114,7 @@ # define _CCCL_NV_DIAG_DEFAULT(_WARNING) _CCCL_PRAGMA(diag_default _WARNING) # endif // !_CCCL_COMPILER(GCC) # endif // !__NVCC_DIAG_PRAGMA_SUPPORT__ -#else // ^^^ _CCCL_HAS_CUDA_COMPILER ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER vvv +#else // ^^^ _CCCL_HAS_CUDA_COMPILER() ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER() vvv # define _CCCL_NV_DIAG_SUPPRESS(_WARNING) # define _CCCL_NV_DIAG_DEFAULT(_WARNING) #endif // other compilers diff --git a/libcudacxx/include/cuda/std/__cccl/execution_space.h b/libcudacxx/include/cuda/std/__cccl/execution_space.h index a62cfc31d75..1ed43f3012a 100644 --- a/libcudacxx/include/cuda/std/__cccl/execution_space.h +++ b/libcudacxx/include/cuda/std/__cccl/execution_space.h @@ -23,7 +23,7 @@ #endif // no system header // We need to ensure that we not only compile with a cuda compiler but also compile cuda source files -#if _CCCL_HAS_CUDA_COMPILER && (defined(__CUDACC__) || defined(_NVHPC_CUDA)) +#if _CCCL_HAS_CUDA_COMPILER() && (defined(__CUDACC__) || defined(_NVHPC_CUDA)) # define _CCCL_HOST __host__ # define _CCCL_DEVICE __device__ # define _CCCL_HOST_DEVICE __host__ __device__ diff --git a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h index 5ddf2e22a3c..cfa6cac7b66 100644 --- a/libcudacxx/include/cuda/std/__cccl/extended_data_types.h +++ b/libcudacxx/include/cuda/std/__cccl/extended_data_types.h @@ -51,12 +51,12 @@ #endif // FIXME: Enable this for clang-cuda in a followup -#if !_CCCL_HAS_CUDA_COMPILER +#if !_CCCL_HAS_CUDA_COMPILER() # undef _CCCL_HAS_LONG_DOUBLE # define _CCCL_HAS_LONG_DOUBLE() 1 -#endif // !_CCCL_HAS_CUDA_COMPILER +#endif // !_CCCL_HAS_CUDA_COMPILER() -#if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ +#if _CCCL_HAS_INCLUDE() && (_CCCL_HAS_CUDA_COMPILER() || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) \ && !defined(CCCL_DISABLE_FP16_SUPPORT) # undef _CCCL_HAS_NVFP16 # define _CCCL_HAS_NVFP16() 1 diff --git a/libcudacxx/include/cuda/std/__cccl/is_non_narrowing_convertible.h b/libcudacxx/include/cuda/std/__cccl/is_non_narrowing_convertible.h index 68994a362a7..bb525aa75c2 100644 --- a/libcudacxx/include/cuda/std/__cccl/is_non_narrowing_convertible.h +++ b/libcudacxx/include/cuda/std/__cccl/is_non_narrowing_convertible.h @@ -24,7 +24,7 @@ namespace __cccl_internal { -#if _CCCL_HAS_CUDA_COMPILER && (defined(__CUDACC__) || defined(_NVHPC_CUDA) || _CCCL_COMPILER(NVRTC)) +#if _CCCL_HAS_CUDA_COMPILER() && (defined(__CUDACC__) || defined(_NVHPC_CUDA) || _CCCL_COMPILER(NVRTC)) template __host__ __device__ _Tp&& __cccl_declval(int); template diff --git a/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h b/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h index ce871498fa1..81e3f0c4a3f 100644 --- a/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h +++ b/libcudacxx/include/cuda/std/__cstdlib/aligned_alloc.h @@ -31,9 +31,9 @@ #include -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() extern "C" _CCCL_DEVICE void* __cuda_syscall_aligned_malloc(size_t, size_t); -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_BEGIN_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__cstdlib/malloc.h b/libcudacxx/include/cuda/std/__cstdlib/malloc.h index ffdf0ec73a4..72521d893a6 100644 --- a/libcudacxx/include/cuda/std/__cstdlib/malloc.h +++ b/libcudacxx/include/cuda/std/__cstdlib/malloc.h @@ -36,7 +36,7 @@ using ::free; using ::malloc; // We need to ensure that we not only compile with a cuda compiler but also compile cuda source files -#if _CCCL_HAS_CUDA_COMPILER && (defined(__CUDACC__) || defined(_NVHPC_CUDA)) +#if _CCCL_HAS_CUDA_COMPILER() && (defined(__CUDACC__) || defined(_NVHPC_CUDA)) _CCCL_NODISCARD _CCCL_HIDE_FROM_ABI _CCCL_DEVICE void* __calloc_device(size_t __n, size_t __size) noexcept { void* __ptr{}; @@ -54,7 +54,7 @@ _CCCL_NODISCARD _CCCL_HIDE_FROM_ABI _CCCL_DEVICE void* __calloc_device(size_t __ return __ptr; } -#endif // _CCCL_HAS_CUDA_COMPILER && (defined(__CUDACC__) || defined(_NVHPC_CUDA)) +#endif // _CCCL_HAS_CUDA_COMPILER() && (defined(__CUDACC__) || defined(_NVHPC_CUDA)) _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI void* calloc(size_t __n, size_t __size) noexcept { diff --git a/libcudacxx/include/cuda/std/__cuda/api_wrapper.h b/libcudacxx/include/cuda/std/__cuda/api_wrapper.h index 61dc291b1ce..8c3e0774691 100644 --- a/libcudacxx/include/cuda/std/__cuda/api_wrapper.h +++ b/libcudacxx/include/cuda/std/__cuda/api_wrapper.h @@ -27,7 +27,7 @@ #include -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # define _CCCL_TRY_CUDA_API(_NAME, _MSG, ...) \ { \ const ::cudaError_t __status = _NAME(__VA_ARGS__); \ @@ -47,9 +47,9 @@ _CCCL_ASSERT(__status == cudaSuccess, _MSG); \ (void) __status; \ } -#else // ^^^ _CCCL_HAS_CUDA_COMPILER ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER vvv +#else // ^^^ _CCCL_HAS_CUDA_COMPILER() ^^^ / vvv !_CCCL_HAS_CUDA_COMPILER() vvv # define _CCCL_TRY_CUDA_API(_NAME, _MSG, ...) # define _CCCL_ASSERT_CUDA_API(_NAME, _MSG, ...) -#endif // !_CCCL_HAS_CUDA_COMPILER +#endif // !_CCCL_HAS_CUDA_COMPILER() #endif //_CUDA__STD__CUDA_API_WRAPPER_H diff --git a/libcudacxx/include/cuda/std/__internal/features.h b/libcudacxx/include/cuda/std/__internal/features.h index 35b1426d866..12f2afa0413 100644 --- a/libcudacxx/include/cuda/std/__internal/features.h +++ b/libcudacxx/include/cuda/std/__internal/features.h @@ -26,11 +26,11 @@ #define _LIBCUDACXX_HAS_SPACESHIP_OPERATOR() 0 #define _LIBCUDACXX_HAS_WCHAR_H() 0 -#if _CCCL_HAS_CUDA_COMPILER || __cpp_aligned_new < 201606 +#if _CCCL_HAS_CUDA_COMPILER() || __cpp_aligned_new < 201606 # define _LIBCUDACXX_HAS_ALIGNED_ALLOCATION() 0 #else # define _LIBCUDACXX_HAS_ALIGNED_ALLOCATION() 1 -#endif // !_CCCL_HAS_CUDA_COMPILER && __cpp_aligned_new >= 201606 +#endif // !_CCCL_HAS_CUDA_COMPILER() && __cpp_aligned_new >= 201606 #if _CCCL_STD_VER <= 2017 || !defined(__cpp_char8_t) # define _LIBCUDACXX_HAS_CHAR8_T() 0 @@ -57,7 +57,7 @@ // libcu++ requires host device support for its tests. Until then restrict usage to at least 12.2 #if _CCCL_HAS_NVFP16() && _CCCL_CUDACC_AT_LEAST(12, 2) \ - && (_CCCL_HAS_CUDA_COMPILER || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) + && (_CCCL_HAS_CUDA_COMPILER() || defined(LIBCUDACXX_ENABLE_HOST_NVFP16)) # define _LIBCUDACXX_HAS_NVFP16() 1 #else # define _LIBCUDACXX_HAS_NVFP16() 0 diff --git a/libcudacxx/include/cuda/std/__tuple_dir/vector_types.h b/libcudacxx/include/cuda/std/__tuple_dir/vector_types.h index ed4081fa1b8..3cfb091aaa0 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/vector_types.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/vector_types.h @@ -20,7 +20,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_CLANG("-Wmismatched-tags") diff --git a/libcudacxx/include/cuda/std/cstring b/libcudacxx/include/cuda/std/cstring index 93e46a15c71..df114b327c9 100644 --- a/libcudacxx/include/cuda/std/cstring +++ b/libcudacxx/include/cuda/std/cstring @@ -33,7 +33,7 @@ using ::memcpy; using ::memset; using ::size_t; -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _CCCL_HIDE_FROM_ABI _CCCL_DEVICE const void* __memchr_device(const void* __ptr, int __c, size_t __n) noexcept { auto __p = static_cast(__ptr); @@ -49,7 +49,7 @@ _CCCL_HIDE_FROM_ABI _CCCL_DEVICE const void* __memchr_device(const void* __ptr, return nullptr; } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_HIDE_FROM_ABI const void* memchr(const void* __ptr, int __c, size_t __n) noexcept { @@ -64,7 +64,7 @@ _LIBCUDACXX_HIDE_FROM_ABI void* memchr(void* __ptr, int __c, size_t __n) noexcep (return const_cast(_CUDA_VSTD::memchr(const_cast(__ptr), __c, __n));)) } -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _CCCL_HIDE_FROM_ABI _CCCL_DEVICE void* __memmove_device(void* __dst, const void* __src, size_t __n) noexcept { auto __d = (__dst <= __src) ? static_cast(__dst) : (static_cast(__dst) + __n - 1); @@ -81,7 +81,7 @@ _CCCL_HIDE_FROM_ABI _CCCL_DEVICE void* __memmove_device(void* __dst, const void* return __dst; } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_HIDE_FROM_ABI void* memmove(void* __dst, const void* __src, size_t __n) noexcept { @@ -93,7 +93,7 @@ _LIBCUDACXX_HIDE_FROM_ABI void* memmove(void* __dst, const void* __src, size_t _ #endif // ^^^ !_CCCL_BUILTIN_MEMMOVE ^^^ } -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() _CCCL_HIDE_FROM_ABI _CCCL_DEVICE int __memcmp_device(const void* __lhs, const void* __rhs, size_t __n) noexcept { auto __l = static_cast(__lhs); @@ -110,7 +110,7 @@ _CCCL_HIDE_FROM_ABI _CCCL_DEVICE int __memcmp_device(const void* __lhs, const vo } return 0; } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() _LIBCUDACXX_HIDE_FROM_ABI int memcmp(const void* __lhs, const void* __rhs, size_t __n) noexcept { diff --git a/libcudacxx/include/cuda/std/detail/__annotated_ptr b/libcudacxx/include/cuda/std/detail/__annotated_ptr index 6417942bf1e..49fb9f68106 100644 --- a/libcudacxx/include/cuda/std/detail/__annotated_ptr +++ b/libcudacxx/include/cuda/std/detail/__annotated_ptr @@ -137,7 +137,7 @@ namespace __detail_ap template _CCCL_DEVICE void* __associate_address_space(void* __ptr, _Property __prop) { -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() if (std::is_same<_Property, access_property::shared>::value == true) { bool __b = __isShared(__ptr); @@ -162,7 +162,7 @@ _CCCL_DEVICE void* __associate_address_space(void* __ptr, _Property __prop) (void) __b; # endif // !_CCCL_BUILTIN_ASSUME } -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() return __ptr; } diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__config b/libcudacxx/include/cuda/std/detail/libcxx/include/__config index 6c1c8aad91b..c0220891c07 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__config @@ -118,7 +118,7 @@ extern "C++" { # define _CCCL_ALIGNAS(x) __attribute__((__aligned__(x))) # endif // !_CCCL_COMPILER(MSVC) && !_CCCL_HAS_FEATURE(cxx_alignas) -# if _CCCL_HAS_CUDA_COMPILER +# if _CCCL_HAS_CUDA_COMPILER() # define _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(size, ptr) (size <= 8) # elif _CCCL_COMPILER(CLANG) || _CCCL_COMPILER(GCC) # define _LIBCUDACXX_ATOMIC_ALWAYS_LOCK_FREE(...) __atomic_always_lock_free(__VA_ARGS__) @@ -263,7 +263,7 @@ using char32_t = unsigned int; # endif // _CCCL_CUDACC_BELOW(12, 2) // CUDA Atomics supersede host atomics in order to insert the host/device dispatch layer -# if _CCCL_CUDA_COMPILER(NVCC) || _CCCL_COMPILER(NVRTC) || _CCCL_COMPILER(NVHPC) || _CCCL_CUDACC +# if _CCCL_CUDA_COMPILER(NVCC) || _CCCL_COMPILER(NVRTC) || _CCCL_COMPILER(NVHPC) || _CCCL_HAS_CUDA_COMPILER() # define _LIBCUDACXX_HAS_CUDA_ATOMIC_IMPL # endif diff --git a/libcudacxx/test/support/host_device_types.h b/libcudacxx/test/support/host_device_types.h index e8fa21b85b9..f5ef2873b10 100644 --- a/libcudacxx/test/support/host_device_types.h +++ b/libcudacxx/test/support/host_device_types.h @@ -79,7 +79,7 @@ struct host_only_type }; #endif // !_CCCL_COMPILER(NVRTC) -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() struct device_only_type { int val_; @@ -143,6 +143,6 @@ struct device_only_type cuda::std::swap(val_, other.val_); } }; -#endif // _CCCL_HAS_CUDA_COMPILER +#endif // _CCCL_HAS_CUDA_COMPILER() #endif // TEST_SUPPORT_HOST_DEVICE_TYPES diff --git a/libcudacxx/test/support/test_iterators.h b/libcudacxx/test/support/test_iterators.h index d7fe898ff4f..9c223768bfc 100644 --- a/libcudacxx/test/support/test_iterators.h +++ b/libcudacxx/test/support/test_iterators.h @@ -1129,6 +1129,9 @@ class cpp20_output_iterator : it_(it) {} + cpp20_output_iterator(cpp20_output_iterator&&) = default; + cpp20_output_iterator& operator=(cpp20_output_iterator&&) = default; + __host__ __device__ constexpr decltype(auto) operator*() const { return *it_; diff --git a/libcudacxx/test/support/test_macros.h b/libcudacxx/test/support/test_macros.h index 77b674372b3..b7942ad6ae0 100644 --- a/libcudacxx/test/support/test_macros.h +++ b/libcudacxx/test/support/test_macros.h @@ -16,7 +16,7 @@ // Use the CCCL compiler detection #define TEST_COMPILER(...) _CCCL_COMPILER(__VA_ARGS__) #define TEST_CUDA_COMPILER(...) _CCCL_CUDA_COMPILER(__VA_ARGS__) -#define TEST_HAS_CUDA_COMPILER _CCCL_HAS_CUDA_COMPILER +#define TEST_HAS_CUDA_COMPILER _CCCL_HAS_CUDA_COMPILER() // Use the CCCL diagnostic suppression #define TEST_DIAG_SUPPRESS_CLANG(...) _CCCL_DIAG_SUPPRESS_CLANG(__VA_ARGS__) diff --git a/thrust/examples/include/host_device.h b/thrust/examples/include/host_device.h index a533e898e47..40b65ace0ea 100644 --- a/thrust/examples/include/host_device.h +++ b/thrust/examples/include/host_device.h @@ -16,7 +16,7 @@ #pragma once -#if !_CCCL_HAS_CUDA_COMPILER +#if !_CCCL_HAS_CUDA_COMPILER() # ifndef __host__ # define __host__ diff --git a/thrust/thrust/detail/malloc_and_free.h b/thrust/thrust/detail/malloc_and_free.h index b30eaa0f5f2..1c8db05f5eb 100644 --- a/thrust/thrust/detail/malloc_and_free.h +++ b/thrust/thrust/detail/malloc_and_free.h @@ -63,7 +63,7 @@ malloc(const thrust::detail::execution_policy_base& exec, std::si } // XXX WAR nvbug 992955 -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # if CUDART_VERSION < 5000 // cudafe generates unqualified calls to free(int *volatile) diff --git a/thrust/thrust/random/detail/normal_distribution_base.h b/thrust/thrust/random/detail/normal_distribution_base.h index 07fb48d8f48..97ccfa73b91 100644 --- a/thrust/thrust/random/detail/normal_distribution_base.h +++ b/thrust/thrust/random/detail/normal_distribution_base.h @@ -148,7 +148,7 @@ class normal_distribution_portable template struct normal_distribution_base { -#if _CCCL_HAS_CUDA_COMPILER && !_CCCL_CUDA_COMPILER(NVHPC) +#if _CCCL_HAS_CUDA_COMPILER() && !_CCCL_CUDA_COMPILER(NVHPC) using type = normal_distribution_nvcc; #else using type = normal_distribution_portable; diff --git a/thrust/thrust/system/cuda/detail/adjacent_difference.h b/thrust/thrust/system/cuda/detail/adjacent_difference.h index 13061c5b86a..4830f7fd3f3 100644 --- a/thrust/thrust/system/cuda/detail/adjacent_difference.h +++ b/thrust/thrust/system/cuda/detail/adjacent_difference.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/assign_value.h b/thrust/thrust/system/cuda/detail/assign_value.h index a5f5312cad2..0e3c6c9e558 100644 --- a/thrust/thrust/system/cuda/detail/assign_value.h +++ b/thrust/thrust/system/cuda/detail/assign_value.h @@ -26,7 +26,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/copy.h b/thrust/thrust/system/cuda/detail/copy.h index 857af1eceec..39b8d2ca881 100644 --- a/thrust/thrust/system/cuda/detail/copy.h +++ b/thrust/thrust/system/cuda/detail/copy.h @@ -84,7 +84,7 @@ THRUST_NAMESPACE_BEGIN namespace cuda_cub { -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() // D->D copy requires NVCC compiler _CCCL_EXEC_CHECK_DISABLE diff --git a/thrust/thrust/system/cuda/detail/copy_if.h b/thrust/thrust/system/cuda/detail/copy_if.h index ae049c53e5e..4c90828954a 100644 --- a/thrust/thrust/system/cuda/detail/copy_if.h +++ b/thrust/thrust/system/cuda/detail/copy_if.h @@ -35,7 +35,8 @@ #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER + +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/core/agent_launcher.h b/thrust/thrust/system/cuda/detail/core/agent_launcher.h index 4a7f3a1ba9a..914a6caca2c 100644 --- a/thrust/thrust/system/cuda/detail/core/agent_launcher.h +++ b/thrust/thrust/system/cuda/detail/core/agent_launcher.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/count.h b/thrust/thrust/system/cuda/detail/count.h index 22b373e18c0..20e556a3b9d 100644 --- a/thrust/thrust/system/cuda/detail/count.h +++ b/thrust/thrust/system/cuda/detail/count.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/equal.h b/thrust/thrust/system/cuda/detail/equal.h index a60f21dd158..1b4d9cfa79b 100644 --- a/thrust/thrust/system/cuda/detail/equal.h +++ b/thrust/thrust/system/cuda/detail/equal.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/extrema.h b/thrust/thrust/system/cuda/detail/extrema.h index a39118c99c6..5065ad52f92 100644 --- a/thrust/thrust/system/cuda/detail/extrema.h +++ b/thrust/thrust/system/cuda/detail/extrema.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/fill.h b/thrust/thrust/system/cuda/detail/fill.h index a31b48674d0..e3bb37e14df 100644 --- a/thrust/thrust/system/cuda/detail/fill.h +++ b/thrust/thrust/system/cuda/detail/fill.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include # include diff --git a/thrust/thrust/system/cuda/detail/find.h b/thrust/thrust/system/cuda/detail/find.h index da2a2ffdce2..a2aa0ceb123 100644 --- a/thrust/thrust/system/cuda/detail/find.h +++ b/thrust/thrust/system/cuda/detail/find.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/for_each.h b/thrust/thrust/system/cuda/detail/for_each.h index ef84f31c99e..03518ef4472 100644 --- a/thrust/thrust/system/cuda/detail/for_each.h +++ b/thrust/thrust/system/cuda/detail/for_each.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/gather.h b/thrust/thrust/system/cuda/detail/gather.h index 8a7f31bbac8..08244d115db 100644 --- a/thrust/thrust/system/cuda/detail/gather.h +++ b/thrust/thrust/system/cuda/detail/gather.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/generate.h b/thrust/thrust/system/cuda/detail/generate.h index be891374ab7..d2261c1d832 100644 --- a/thrust/thrust/system/cuda/detail/generate.h +++ b/thrust/thrust/system/cuda/detail/generate.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/get_value.h b/thrust/thrust/system/cuda/detail/get_value.h index 8e7b46dd005..140305c6e75 100644 --- a/thrust/thrust/system/cuda/detail/get_value.h +++ b/thrust/thrust/system/cuda/detail/get_value.h @@ -26,7 +26,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/inner_product.h b/thrust/thrust/system/cuda/detail/inner_product.h index 07ba9c902b1..c5176b2e853 100644 --- a/thrust/thrust/system/cuda/detail/inner_product.h +++ b/thrust/thrust/system/cuda/detail/inner_product.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include # include diff --git a/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h b/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h index c2249bcc85e..614ed1fd562 100644 --- a/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h +++ b/thrust/thrust/system/cuda/detail/internal/copy_cross_system.h @@ -137,7 +137,7 @@ OutputIt _CCCL_HOST cross_system_copy_n( return ret; } -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() // non-trivial copy D->H, only supported with NVCC compiler // because copy ctor must have __device__ annotations, which is nvcc-only // feature diff --git a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h index 9b1e1cd7457..ba8e9bba6fd 100644 --- a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h +++ b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h @@ -37,7 +37,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/iter_swap.h b/thrust/thrust/system/cuda/detail/iter_swap.h index 2740177ca5e..97512182b28 100644 --- a/thrust/thrust/system/cuda/detail/iter_swap.h +++ b/thrust/thrust/system/cuda/detail/iter_swap.h @@ -26,7 +26,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/merge.h b/thrust/thrust/system/cuda/detail/merge.h index c6b480e6c88..f0b2e21852d 100644 --- a/thrust/thrust/system/cuda/detail/merge.h +++ b/thrust/thrust/system/cuda/detail/merge.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/mismatch.h b/thrust/thrust/system/cuda/detail/mismatch.h index 6e723bf4d75..d026b973c1c 100644 --- a/thrust/thrust/system/cuda/detail/mismatch.h +++ b/thrust/thrust/system/cuda/detail/mismatch.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/parallel_for.h b/thrust/thrust/system/cuda/detail/parallel_for.h index cb4e9aac3e8..fdac924843f 100644 --- a/thrust/thrust/system/cuda/detail/parallel_for.h +++ b/thrust/thrust/system/cuda/detail/parallel_for.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/partition.h b/thrust/thrust/system/cuda/detail/partition.h index 474836f2b57..f9117978f27 100644 --- a/thrust/thrust/system/cuda/detail/partition.h +++ b/thrust/thrust/system/cuda/detail/partition.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/per_device_resource.h b/thrust/thrust/system/cuda/detail/per_device_resource.h index 8cc91a4793a..1e1078e4874 100644 --- a/thrust/thrust/system/cuda/detail/per_device_resource.h +++ b/thrust/thrust/system/cuda/detail/per_device_resource.h @@ -37,7 +37,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/reduce.h b/thrust/thrust/system/cuda/detail/reduce.h index 2b232e31be9..552e1c2b7a8 100644 --- a/thrust/thrust/system/cuda/detail/reduce.h +++ b/thrust/thrust/system/cuda/detail/reduce.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/reduce_by_key.h b/thrust/thrust/system/cuda/detail/reduce_by_key.h index a9cee37215f..671c42a9d59 100644 --- a/thrust/thrust/system/cuda/detail/reduce_by_key.h +++ b/thrust/thrust/system/cuda/detail/reduce_by_key.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/remove.h b/thrust/thrust/system/cuda/detail/remove.h index 76d23eeb81d..f53e2263265 100644 --- a/thrust/thrust/system/cuda/detail/remove.h +++ b/thrust/thrust/system/cuda/detail/remove.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include THRUST_NAMESPACE_BEGIN diff --git a/thrust/thrust/system/cuda/detail/replace.h b/thrust/thrust/system/cuda/detail/replace.h index 12257f2f224..ae538b53ce5 100644 --- a/thrust/thrust/system/cuda/detail/replace.h +++ b/thrust/thrust/system/cuda/detail/replace.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/reverse.h b/thrust/thrust/system/cuda/detail/reverse.h index 682b0efcf5f..46fca87ad41 100644 --- a/thrust/thrust/system/cuda/detail/reverse.h +++ b/thrust/thrust/system/cuda/detail/reverse.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include THRUST_NAMESPACE_BEGIN diff --git a/thrust/thrust/system/cuda/detail/scan.h b/thrust/thrust/system/cuda/detail/scan.h index 6f83aafcf21..15c69536686 100644 --- a/thrust/thrust/system/cuda/detail/scan.h +++ b/thrust/thrust/system/cuda/detail/scan.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/scan_by_key.h b/thrust/thrust/system/cuda/detail/scan_by_key.h index 2edb0ba7e41..042a32ef489 100644 --- a/thrust/thrust/system/cuda/detail/scan_by_key.h +++ b/thrust/thrust/system/cuda/detail/scan_by_key.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/scatter.h b/thrust/thrust/system/cuda/detail/scatter.h index 17a61c70f77..80f2ea80eaa 100644 --- a/thrust/thrust/system/cuda/detail/scatter.h +++ b/thrust/thrust/system/cuda/detail/scatter.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index 3d4b7eff4f0..deb58d488f3 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/sort.h b/thrust/thrust/system/cuda/detail/sort.h index a32e5bd0b75..15c673e850a 100644 --- a/thrust/thrust/system/cuda/detail/sort.h +++ b/thrust/thrust/system/cuda/detail/sort.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/swap_ranges.h b/thrust/thrust/system/cuda/detail/swap_ranges.h index 23cd499ad2a..fbd2547714f 100644 --- a/thrust/thrust/system/cuda/detail/swap_ranges.h +++ b/thrust/thrust/system/cuda/detail/swap_ranges.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include # include diff --git a/thrust/thrust/system/cuda/detail/tabulate.h b/thrust/thrust/system/cuda/detail/tabulate.h index dee20887a77..2e8583f8f7a 100644 --- a/thrust/thrust/system/cuda/detail/tabulate.h +++ b/thrust/thrust/system/cuda/detail/tabulate.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/transform.h b/thrust/thrust/system/cuda/detail/transform.h index 6863fba4490..a04b468dab0 100644 --- a/thrust/thrust/system/cuda/detail/transform.h +++ b/thrust/thrust/system/cuda/detail/transform.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/transform_reduce.h b/thrust/thrust/system/cuda/detail/transform_reduce.h index ed7d26942f7..16e8c935421 100644 --- a/thrust/thrust/system/cuda/detail/transform_reduce.h +++ b/thrust/thrust/system/cuda/detail/transform_reduce.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include diff --git a/thrust/thrust/system/cuda/detail/transform_scan.h b/thrust/thrust/system/cuda/detail/transform_scan.h index 335e4ea97c3..e3230d9c77d 100644 --- a/thrust/thrust/system/cuda/detail/transform_scan.h +++ b/thrust/thrust/system/cuda/detail/transform_scan.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include # include diff --git a/thrust/thrust/system/cuda/detail/uninitialized_copy.h b/thrust/thrust/system/cuda/detail/uninitialized_copy.h index 5be6813ebf7..4b78d56832c 100644 --- a/thrust/thrust/system/cuda/detail/uninitialized_copy.h +++ b/thrust/thrust/system/cuda/detail/uninitialized_copy.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include # include diff --git a/thrust/thrust/system/cuda/detail/uninitialized_fill.h b/thrust/thrust/system/cuda/detail/uninitialized_fill.h index 601e0b9f0b1..23dc99628cb 100644 --- a/thrust/thrust/system/cuda/detail/uninitialized_fill.h +++ b/thrust/thrust/system/cuda/detail/uninitialized_fill.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include # include # include diff --git a/thrust/thrust/system/cuda/detail/unique.h b/thrust/thrust/system/cuda/detail/unique.h index a3eb9462cef..985143cf4be 100644 --- a/thrust/thrust/system/cuda/detail/unique.h +++ b/thrust/thrust/system/cuda/detail/unique.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include diff --git a/thrust/thrust/system/cuda/detail/unique_by_key.h b/thrust/thrust/system/cuda/detail/unique_by_key.h index cfe8dc134ee..c3bb8203e20 100644 --- a/thrust/thrust/system/cuda/detail/unique_by_key.h +++ b/thrust/thrust/system/cuda/detail/unique_by_key.h @@ -36,7 +36,7 @@ # pragma system_header #endif // no system header -#if _CCCL_HAS_CUDA_COMPILER +#if _CCCL_HAS_CUDA_COMPILER() # include