Skip to content

Commit

Permalink
Properly test internal headers (#1258)
Browse files Browse the repository at this point in the history
* Add tests for internal headers, so that we know that they do not rely on transitive includes
* Add tests for fallback definitions of traits, to ensure hat all fallbacks are functional
* Add tests that all public headers can be included
* Add tests that all public headers can be included by just a host compiler

Co-authored-by: Wesley Maxey <[email protected]>
  • Loading branch information
miscco and wmaxey authored Jan 18, 2024
1 parent 962ace6 commit f0f8808
Show file tree
Hide file tree
Showing 50 changed files with 598 additions and 250 deletions.
7 changes: 6 additions & 1 deletion CMakePresets.json
Original file line number Diff line number Diff line change
Expand Up @@ -273,7 +273,12 @@
{
"name": "libcudacxx-base",
"hidden": true,
"targets": ["libcudacxx.test.lit.precompile"]
"targets": [
"libcudacxx.test.internal_headers",
"libcudacxx.test.public_headers",
"libcudacxx.test.public_headers_host_only",
"libcudacxx.test.lit.precompile"
]
},
{
"name": "libcudacxx-nvrtc-cpp11",
Expand Down
8 changes: 4 additions & 4 deletions libcudacxx/include/cuda/discard_memory
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@

_LIBCUDACXX_BEGIN_NAMESPACE_CUDA

inline _LIBCUDACXX_HOST_DEVICE void discard_memory(volatile void* __ptr, std::size_t __nbytes) noexcept
inline _LIBCUDACXX_HOST_DEVICE void discard_memory(volatile void* __ptr, size_t __nbytes) noexcept
{
// The discard PTX instruction is only available with PTX ISA 7.4 and later
#if __cccl_ptx_isa < 740ULL
Expand All @@ -30,12 +30,12 @@ inline _LIBCUDACXX_HOST_DEVICE void discard_memory(volatile void* __ptr, std::si

char* __p = reinterpret_cast<char*>(const_cast<void*>(__ptr));
char* const __end_p = __p + __nbytes;
static constexpr std::size_t _LINE_SIZE = 128;
static constexpr size_t _LINE_SIZE = 128;

// Trim the first block and last block if they're not 128 bytes aligned
std::size_t __misalignment = reinterpret_cast<std::uintptr_t>(__p) % _LINE_SIZE;
size_t __misalignment = reinterpret_cast<uintptr_t>(__p) % _LINE_SIZE;
char* __start_aligned = __misalignment == 0 ? __p : __p + (_LINE_SIZE - __misalignment);
char* const __end_aligned = __end_p - (reinterpret_cast<std::uintptr_t>(__end_p) % _LINE_SIZE);
char* const __end_aligned = __end_p - (reinterpret_cast<uintptr_t>(__end_p) % _LINE_SIZE);

while (__start_aligned < __end_aligned) {
asm volatile("discard.global.L2 [%0], 128;" ::"l"(__start_aligned) :);
Expand Down
8 changes: 6 additions & 2 deletions libcudacxx/include/cuda/std/detail/__annotated_ptr
Original file line number Diff line number Diff line change
Expand Up @@ -63,7 +63,9 @@ namespace __detail_ap {
_LIBCUDACXX_ASSERT(__b, "");
#if !defined(_LIBCUDACXX_CUDACC_BELOW_11_2)
__builtin_assume(__b);
#endif // !defined(_LIBCUDACXX_CUDACC_BELOW_11_2)
#else // ^^^ !_LIBCUDACXX_CUDACC_BELOW_11_2 ^^^ / vvv _LIBCUDACXX_CUDACC_BELOW_11_2 vvv
(void)__b;
#endif // _LIBCUDACXX_CUDACC_BELOW_11_2
} else if (std::is_same<_Property, access_property::global>::value == true ||
std::is_same<_Property, access_property::normal>::value == true ||
std::is_same<_Property, access_property::persisting>::value == true ||
Expand All @@ -73,7 +75,9 @@ namespace __detail_ap {
_LIBCUDACXX_ASSERT(__b, "");
#if !defined(_LIBCUDACXX_CUDACC_BELOW_11_2)
__builtin_assume(__b);
#endif // !defined(_LIBCUDACXX_CUDACC_BELOW_11_2)
#else // ^^^ !_LIBCUDACXX_CUDACC_BELOW_11_2 ^^^ / vvv _LIBCUDACXX_CUDACC_BELOW_11_2 vvv
(void)__b;
#endif // _LIBCUDACXX_CUDACC_BELOW_11_2
}

return __ptr;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,7 @@
# include <__config>
#endif // __cuda_std__

#include "../__algorithm/comp.h"
#include "../__iterator/iterator_traits.h"
#include "../__type_traits/add_lvalue_reference.h"

Expand Down
12 changes: 6 additions & 6 deletions libcudacxx/include/cuda/std/detail/libcxx/include/__assert
Original file line number Diff line number Diff line change
Expand Up @@ -51,18 +51,18 @@

#if _LIBCUDACXX_ENABLE_ASSERTIONS
# define _LIBCUDACXX_ASSERT(expression, message) \
(_LIBCUDACXX_DIAGNOSTIC_PUSH \
_LIBCUDACXX_CLANG_DIAGNOSTIC_IGNORED("-Wassume") \
(_CCCL_DIAG_PUSH \
_CCCL_DIAG_SUPPRESS_CLANG("-Wassume") \
__builtin_expect(static_cast<bool>(expression), 1) ? \
(void)0 : \
::_CUDA_VSTD::__libcpp_verbose_abort("%s:%d: assertion %s failed: %s", __FILE__, __LINE__, #expression, message)
_LIBCUDACXX_DIAGNOSTIC_POP)
_CCCL_DIAG_POP)
#elif 0 // !defined(_LIBCUDACXX_ASSERTIONS_DISABLE_ASSUME) && __has_builtin(__builtin_assume)
# define _LIBCUDACXX_ASSERT(expression, message) \
(_LIBCUDACXX_DIAGNOSTIC_PUSH \
_LIBCUDACXX_CLANG_DIAGNOSTIC_IGNORED("-Wassume") \
(_CCCL_DIAG_PUSH \
_CCCL_DIAG_SUPPRESS_CLANG("-Wassume") \
__builtin_assume(static_cast<bool>(expression)) \
_LIBCUDACXX_DIAGNOSTIC_POP)
_CCCL_DIAG_POP)
#else
# define _LIBCUDACXX_ASSERT(expression, message) ((void)0)
#endif
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -26,12 +26,14 @@
#if !defined(_CCCL_EXEC_CHECK_DISABLE)
# if defined(_CCCL_COMPILER_NVRTC) || defined(_CCCL_CUDA_COMPILER_NVHPC) || defined(_CCCL_CUDA_COMPILER_CLANG)
# define _CCCL_EXEC_CHECK_DISABLE
# else // defined(_CCCL_CUDA_COMPILER_NVCC)
# elif defined(_CCCL_CUDA_COMPILER_NVCC) // defined(_CCCL_CUDA_COMPILER_NVCC)
# if defined(_CCCL_COMPILER_MSVC)
# define _CCCL_EXEC_CHECK_DISABLE __pragma("nv_exec_check_disable")
# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv
# define _CCCL_EXEC_CHECK_DISABLE _Pragma("nv_exec_check_disable")
# endif // !_CCCL_COMPILER_MSVC
# else
# define _CCCL_EXEC_CHECK_DISABLE
# endif // _CCCL_CUDA_COMPILER_NVCC
#endif // !_CCCL_EXEC_CHECK_DISABLE

Expand Down
24 changes: 23 additions & 1 deletion libcudacxx/include/cuda/std/detail/libcxx/include/__cccl_config
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,12 @@
# define _CCCL_CUDA_COMPILER_CLANG
#endif

// Shorthand to check whether there is a cuda compiler available
#if defined(_CCCL_CUDA_COMPILER_NVCC) || defined(_CCCL_CUDA_COMPILER_NVHPC) || defined(_CCCL_CUDA_COMPILER_CLANG) \
|| defined(_CCCL_COMPILER_NVRTC)
# define _CCCL_CUDA_COMPILER
#endif // cuda compiler available

// Enforce that cccl headers are treated as system headers
#if defined(_CCCL_COMPILER_GCC) || defined(_CCCL_COMPILER_NVHPC) || defined(_CCCL_COMPILER_ICC) \
|| defined(_CCCL_COMPILER_ICC_LLVM)
Expand Down Expand Up @@ -92,24 +98,35 @@
# define _CCCL_DIAG_SUPPRESS_CLANG(str) _Pragma(_CCCL_TOSTRING(clang diagnostic ignored str))
# define _CCCL_DIAG_SUPPRESS_GCC(str)
# define _CCCL_DIAG_SUPPRESS_NVHPC(str)
# define _CCCL_DIAG_SUPPRESS_MSVC(str)
#elif defined(_CCCL_COMPILER_GCC)
# define _CCCL_DIAG_PUSH _Pragma("GCC diagnostic push")
# define _CCCL_DIAG_POP _Pragma("GCC diagnostic pop")
# define _CCCL_DIAG_SUPPRESS_CLANG(str)
# define _CCCL_DIAG_SUPPRESS_GCC(str) _Pragma(_CCCL_TOSTRING(GCC diagnostic ignored str))
# define _CCCL_DIAG_SUPPRESS_NVHPC(str)
# define _CCCL_DIAG_SUPPRESS_MSVC(str)
#elif defined(_CCCL_COMPILER_NVHPC)
# define _CCCL_DIAG_PUSH _Pragma("diagnostic push")
# define _CCCL_DIAG_POP _Pragma("diagnostic pop")
# define _CCCL_DIAG_SUPPRESS_CLANG(str)
# define _CCCL_DIAG_SUPPRESS_GCC(str)
# define _CCCL_DIAG_SUPPRESS_NVHPC(str) _Pragma(_LIBCUDACXX_TOSTRING(diag_suppress str))
# define _CCCL_DIAG_SUPPRESS_NVHPC(str) _Pragma(_CCCL_TOSTRING(diag_suppress str))
# define _CCCL_DIAG_SUPPRESS_MSVC(str)
#elif defined(_CCCL_COMPILER_MSVC)
# define _CCCL_DIAG_PUSH __pragma(warning(push))
# define _CCCL_DIAG_POP __pragma(warning(pop))
# define _CCCL_DIAG_SUPPRESS_CLANG(str)
# define _CCCL_DIAG_SUPPRESS_GCC(str)
# define _CCCL_DIAG_SUPPRESS_NVHPC(str)
# define _CCCL_DIAG_SUPPRESS_MSVC(str) __pragma(warning(disable : str))
#else
# define _CCCL_DIAG_PUSH
# define _CCCL_DIAG_POP
# define _CCCL_DIAG_SUPPRESS_CLANG(str)
# define _CCCL_DIAG_SUPPRESS_GCC(str)
# define _CCCL_DIAG_SUPPRESS_NVHPC(str)
# define _CCCL_DIAG_SUPPRESS_MSVC(str)
#endif

// Convenient shortcuts to silence common warnings
Expand All @@ -125,6 +142,11 @@
_CCCL_DIAG_SUPPRESS_GCC("-Wdeprecated") \
_CCCL_DIAG_SUPPRESS_GCC("-Wdeprecated-declarations")
# define _CCCL_DIAG_SUPPRESS_DEPRECATED_POP _CCCL_DIAG_POP
#elif defined(_CCCL_COMPILER_MSVC)
# define _CCCL_DIAG_SUPPRESS_DEPRECATED_PUSH \
_CCCL_DIAG_PUSH \
_CCCL_DIAG_SUPPRESS_MSVC(4996)
# define _CCCL_DIAG_SUPPRESS_DEPRECATED_POP _CCCL_DIAG_POP
#else // !_CCCL_COMPILER_CLANG && !_CCCL_COMPILER_GCC
# define _CCCL_DIAG_SUPPRESS_DEPRECATED_PUSH
# define _CCCL_DIAG_SUPPRESS_DEPRECATED_POP
Expand Down
39 changes: 11 additions & 28 deletions libcudacxx/include/cuda/std/detail/libcxx/include/__config
Original file line number Diff line number Diff line change
Expand Up @@ -389,7 +389,7 @@ extern "C++" {
# define _LIBCUDACXX_NO_CFI
#endif

#if __ISO_C_VISIBLE >= 2011 || __cplusplus >= 201103L
#if (defined(__ISO_C_VISIBLE) && __ISO_C_VISIBLE >= 2011) || __cplusplus >= 201103L
# if defined(__FreeBSD__)
# define _LIBCUDACXX_HAS_QUICK_EXIT
# define _LIBCUDACXX_HAS_C11_FEATURES
Expand Down Expand Up @@ -572,9 +572,11 @@ extern "C++" {

// TODO: Clang incorrectly reports that __is_array is true for T[0].
// Re-enable the branch once https://llvm.org/PR54705 is fixed.
#ifndef _LIBCUDACXX_USE_IS_ARRAY_FALLBACK
#if defined(_LIBCUDACXX_COMPILER_CLANG)
#define _LIBCUDACXX_USE_IS_ARRAY_FALLBACK
#endif // _LIBCUDACXX_COMPILER_CLANG
#endif // !_LIBCUDACXX_USE_IS_ARRAY_FALLBACK

#if __check_builtin(is_assignable) \
|| defined(_LIBCUDACXX_COMPILER_MSVC)
Expand Down Expand Up @@ -650,9 +652,11 @@ extern "C++" {
#define _LIBCUDACXX_IS_LVALUE_REFERENCE(...) __is_lvalue_reference(__VA_ARGS__)
#endif // __check_builtin(is_lvalue_reference)

#ifndef _LIBCUDACXX_USE_IS_LVALUE_REFERENCE_FALLBACK
#if defined(_LIBCUDACXX_CUDACC_BELOW_11_3)
#define _LIBCUDACXX_USE_IS_LVALUE_REFERENCE_FALLBACK
#endif // nvcc < 11.3
#endif // !_LIBCUDACXX_USE_IS_LVALUE_REFERENCE_FALLBACK

#if __check_builtin(is_nothrow_assignable) \
|| defined(_LIBCUDACXX_COMPILER_MSVC) \
Expand All @@ -676,9 +680,11 @@ extern "C++" {
#define _LIBCUDACXX_IS_OBJECT(...) __is_object(__VA_ARGS__)
#endif // __check_builtin(is_object)

#ifndef _LIBCUDACXX_USE_IS_OBJECT_FALLBACK
#if defined(_LIBCUDACXX_CUDACC_BELOW_11_3)
#define _LIBCUDACXX_USE_IS_OBJECT_FALLBACK
#endif // nvcc < 11.3
#endif // !_LIBCUDACXX_USE_IS_OBJECT_FALLBACK

#if __check_builtin(is_pod) \
|| (defined(_LIBCUDACXX_COMPILER_GCC) && _GNUC_VER >= 403) \
Expand Down Expand Up @@ -778,9 +784,11 @@ extern "C++" {
#define _LIBCUDACXX_IS_UNSIGNED(...) __is_unsigned(__VA_ARGS__)
#endif // __check_builtin(is_unsigned)

#ifndef _LIBCUDACXX_USE_IS_UNSIGNED_FALLBACK
#if defined(_LIBCUDACXX_CUDACC_BELOW_11_3)
#define _LIBCUDACXX_USE_IS_UNSIGNED_FALLBACK
#endif // nvcc < 11.3
#endif // !_LIBCUDACXX_USE_IS_UNSIGNED_FALLBACK

// libstdc++ defines this as a function, breaking functionality
#if 0 // __check_builtin(is_void)
Expand Down Expand Up @@ -954,11 +962,12 @@ typedef __char32_t char32_t;

#elif defined(_LIBCUDACXX_COMPILER_GCC)

#ifndef _LIBCUDACXX_USE_IS_ASSIGNABLE_FALLBACK
// FIXME: GCC 8.0 supports this trait, but it has a bug.
// https://gcc.gnu.org/bugzilla/show_bug.cgi?id=91592
// https://godbolt.org/z/IljfIw
#define _LIBCUDACXX_USE_IS_ASSIGNABLE_FALLBACK

#endif // _LIBCUDACXX_USE_IS_ASSIGNABLE_FALLBACK

// Determine if GCC supports relaxed constexpr
#if !defined(__cpp_constexpr) || __cpp_constexpr < 201304L
Expand Down Expand Up @@ -1742,32 +1751,6 @@ typedef unsigned int char32_t;
# define _LIBCUDACXX_NODISCARD_FRIEND _LIBCUDACXX_NODISCARD_ATTRIBUTE friend
#endif // !_LIBCUDACXX_CUDACC_BELOW_11_3 && !_LIBCUDACXX_COMPILER_CLANG

# ifdef _LIBCUDACXX_COMPILER_CLANG
# define _LIBCUDACXX_DIAGNOSTIC_PUSH _Pragma("clang diagnostic push")
# define _LIBCUDACXX_DIAGNOSTIC_POP _Pragma("clang diagnostic pop")
# define _LIBCUDACXX_CLANG_DIAGNOSTIC_IGNORED(str) _Pragma(_LIBCUDACXX_TOSTRING(clang diagnostic ignored str))
# define _LIBCUDACXX_GCC_DIAGNOSTIC_IGNORED(str)
# define _LIBCUDACXX_NVHPC_DIAGNOSTIC_IGNORED(str)
# elif defined(_LIBCUDACXX_COMPILER_GCC)
# define _LIBCUDACXX_DIAGNOSTIC_PUSH _Pragma("GCC diagnostic push")
# define _LIBCUDACXX_DIAGNOSTIC_POP _Pragma("GCC diagnostic pop")
# define _LIBCUDACXX_CLANG_DIAGNOSTIC_IGNORED(str)
# define _LIBCUDACXX_GCC_DIAGNOSTIC_IGNORED(str) _Pragma(_LIBCUDACXX_TOSTRING(GCC diagnostic ignored str))
# define _LIBCUDACXX_NVHPC_DIAGNOSTIC_IGNORED(str)
# elif defined(_LIBCUDACXX_COMPILER_NVHPC)
# define _LIBCUDACXX_DIAGNOSTIC_PUSH _Pragma("diagnostic push")
# define _LIBCUDACXX_DIAGNOSTIC_POP _Pragma("diagnostic pop")
# define _LIBCUDACXX_CLANG_DIAGNOSTIC_IGNORED(str)
# define _LIBCUDACXX_GCC_DIAGNOSTIC_IGNORED(str)
# define _LIBCUDACXX_NVHPC_DIAGNOSTIC_IGNORED(str) _Pragma(_LIBCUDACXX_TOSTRING(diag_suppress str))
# else
# define _LIBCUDACXX_DIAGNOSTIC_PUSH
# define _LIBCUDACXX_DIAGNOSTIC_POP
# define _LIBCUDACXX_CLANG_DIAGNOSTIC_IGNORED(str)
# define _LIBCUDACXX_GCC_DIAGNOSTIC_IGNORED(str)
# define _LIBCUDACXX_NVHPC_DIAGNOSTIC_IGNORED(str)
# endif

// _LIBCUDACXX_NODISCARD_EXT may be used to apply [[nodiscard]] to entities not
// specified as such as an extension.
#if defined(_LIBCUDACXX_ENABLE_NODISCARD) && !defined(_LIBCUDACXX_DISABLE_NODISCARD_EXT)
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,10 @@

#include "../cstdlib" // _LIBCUDACXX_UNREACHABLE
#include "../__type_traits/void_t.h" // _CUDA_VSTD::__void_t

#if defined(_CCCL_CUDA_COMPILER)
#include "../__cuda/ptx.h" // cuda::ptx::*
#endif // _CCCL_CUDA_COMPILER

#if defined(_LIBCUDACXX_COMPILER_NVRTC)
#define _LIBCUDACXX_OFFSET_IS_ZERO(type, member) !(&(((type *)0)->member))
Expand Down Expand Up @@ -567,6 +570,7 @@ inline _CUDA_VSTD::uint64_t * barrier_native_handle(barrier<thread_scope_block>
return reinterpret_cast<_CUDA_VSTD::uint64_t *>(&b.__barrier);
}

#if defined(_CCCL_CUDA_COMPILER)

// Hide arrive_tx when CUDA architecture is insufficient. Note the
// (!defined(__CUDA_MINIMUM_ARCH__)). This is required to make sure the function
Expand Down Expand Up @@ -686,9 +690,12 @@ void barrier_expect_tx(
: "memory");
}
#endif // __CUDA_MINIMUM_ARCH__
#endif // _CCCL_CUDA_COMPILER

_LIBCUDACXX_END_NAMESPACE_CUDA_DEVICE

#if defined(_CCCL_CUDA_COMPILER)

_LIBCUDACXX_BEGIN_NAMESPACE_CUDA

template<>
Expand Down Expand Up @@ -1218,4 +1225,6 @@ async_contract_fulfillment memcpy_async(void * __destination, void const * __sou

_LIBCUDACXX_END_NAMESPACE_CUDA

#endif // _CCCL_CUDA_COMPILER

#endif // _LIBCUDACXX___CUDA_BARRIER_H
Original file line number Diff line number Diff line change
Expand Up @@ -13,9 +13,6 @@
#include <__config>
#endif // __cuda_std__

#include "../__utility/move.h"
#include "../exception"

#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
# pragma GCC system_header
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
Expand All @@ -24,6 +21,9 @@
# pragma system_header
#endif // no system header

#include "../__utility/move.h"
#include "../exception"

#if _CCCL_STD_VER > 2011

_LIBCUDACXX_BEGIN_NAMESPACE_STD
Expand All @@ -39,15 +39,15 @@ class bad_expected_access<void> : public exception {
bad_expected_access(bad_expected_access&&) = default;
bad_expected_access& operator=(const bad_expected_access&) = default;
bad_expected_access& operator=(bad_expected_access&&) = default;
~bad_expected_access() override = default;
~bad_expected_access() noexcept override = default;

public:
// The way this has been designed (by using a class template below) means that we'll already
// have a profusion of these vtables in TUs, and the dynamic linker will already have a bunch
// of work to do. So it is not worth hiding the <void> specialization in the dylib, given that
// it adds deployment target restrictions.
_LIBCUDACXX_INLINE_VISIBILITY
const char* what() const noexcept override { return "bad access to _CUDA_VSTD::expected"; }
const char* what() const noexcept override { return "bad access to std::expected"; }
};

template <class _Err>
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -45,12 +45,10 @@ template <class _Arg1, class _Arg2, class _Result> struct __binary_function_keep
};

#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_UNARY_BINARY_FUNCTION)
_LIBCUDACXX_DIAGNOSTIC_PUSH
_LIBCUDACXX_CLANG_DIAGNOSTIC_IGNORED("-Wdeprecated-declarations")
_LIBCUDACXX_GCC_DIAGNOSTIC_IGNORED("-Wdeprecated-declarations")
_CCCL_DIAG_SUPPRESS_DEPRECATED_PUSH
template <class _Arg1, class _Arg2, class _Result>
using __binary_function = binary_function<_Arg1, _Arg2, _Result>;
_LIBCUDACXX_DIAGNOSTIC_POP
_CCCL_DIAG_SUPPRESS_DEPRECATED_POP
#else
template <class _Arg1, class _Arg2, class _Result>
using __binary_function = __binary_function_keep_layout_base<_Arg1, _Arg2, _Result>;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,8 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD

#if _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_BINDERS)

_CCCL_DIAG_SUPPRESS_DEPRECATED_PUSH

template <class __Operation>
class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_DEPRECATED_IN_CXX11 binder1st
: public __unary_function<typename __Operation::second_argument_type, typename __Operation::result_type>
Expand Down Expand Up @@ -56,6 +58,8 @@ binder1st<__Operation>
bind1st(const __Operation& __op, const _Tp& __x)
{return binder1st<__Operation>(__op, __x);}

_CCCL_DIAG_SUPPRESS_DEPRECATED_POP

#endif // _CCCL_STD_VER <= 2014 || defined(_LIBCUDACXX_ENABLE_CXX17_REMOVED_BINDERS)

_LIBCUDACXX_END_NAMESPACE_STD
Expand Down
Loading

0 comments on commit f0f8808

Please sign in to comment.