diff --git a/ci/matrix.yaml b/ci/matrix.yaml index 9d41224c0450..d9122a539d6e 100644 --- a/ci/matrix.yaml +++ b/ci/matrix.yaml @@ -105,6 +105,8 @@ workflows: - {jobs: ['build'], project: 'cudax', std: 'max', cxx: ['gcc12']} - {jobs: ['build'], project: 'cudax', std: 'max', cxx: ['clang16', 'clang17', 'clang18', 'clang19', 'clang20']} - {jobs: ['build'], project: 'cudax', std: 'all', cxx: ['gcc', 'clang', 'msvc']} # Newest + # CTK '13.X' build with tile support: default projects + - {jobs: ['build'], ctk: '13.X', cxx: ['gcc', 'clang'], args: '--enable-tile'} # Current CTK testing: - {jobs: ['test'], project: 'thrust', std: 'max', cxx: ['gcc', 'clang', 'msvc'], gpu: 'rtx4090'} - {jobs: ['test'], project: ['libcudacxx', 'cudax'], std: 'max', cxx: ['gcc', 'clang', 'msvc'], gpu: 't4'} @@ -234,6 +236,8 @@ workflows: - {project: 'cudax', jobs: ['test'], std: 'max', cxx: ['gcc', 'msvc'], gpu: 'rtx2080', sm: 'gpu'} - {project: 'cudax', jobs: ['build'], std: 'max', cxx: 'clang', sm: '75;120'} - {project: 'cudax', jobs: ['build'], std: 'max', ctk: 'nvhpc', cxx: 'nvhpc', sm: '75;120'} + # CTK '13.X' build with tile support: default projects + - {jobs: ['build'], ctk: '13.X', cxx: ['gcc', 'clang'], args: '--enable-tile'} # stdpar - {project: 'stdpar', jobs: ['build'], std: 'max', ctk: 'nvhpc', cxx: 'nvhpc'} # Python + support diff --git a/cub/test/catch2_test_device_find.cu b/cub/test/catch2_test_device_find.cu index 85913025f79d..d68d08978c3b 100644 --- a/cub/test/catch2_test_device_find.cu +++ b/cub/test/catch2_test_device_find.cu @@ -10,6 +10,7 @@ #include #include +#include #include "catch2_test_device_reduce.cuh" #include "catch2_test_launch_helper.h" @@ -43,8 +44,8 @@ enum class gen_data_t template auto compute_find_if_reference(InputIt first, InputIt last, Predicate predicate) -> OffsetT { - const auto it = std::find_if(first, last, predicate); // not thrust::find_if because it will rely on cub::FindIf - return static_cast(std::distance(first, it)); + const auto it = cuda::std::find_if(first, last, predicate); // not thrust::find_if because it will rely on cub::FindIf + return static_cast(cuda::std::distance(first, it)); } CUB_TEST("Device find_if works", "[device][find_if]", CUB_SMALL, value_types, offset_types) @@ -87,7 +88,7 @@ CUB_TEST("Device find_if works", "[device][find_if]", CUB_SMALL, value_types, of else { // omit the largest value from the random values so we have a value to that does not occur - c2h::gen(C2H_SEED(1), in_items, input_t{0}, static_cast(::cuda::std::numeric_limits::max() - 1)); + c2h::gen(C2H_SEED(1), in_items, input_t{0}, static_cast(cuda::std::numeric_limits::max() - 1)); } } else @@ -111,7 +112,7 @@ CUB_TEST("Device find_if works", "[device][find_if]", CUB_SMALL, value_types, of else { // max value is neither in the random input and nor in the constant - val_to_find = ::cuda::std::numeric_limits::max(); + val_to_find = cuda::std::numeric_limits::max(); } auto predicate = predice_t{val_to_find}; @@ -191,7 +192,7 @@ CUB_TEST("Device find_if works with non primitive iterator", "[device][find_if]" } { // transform_iterator of counting_iterator input and thrust device_ptr output - auto t_it = cuda::make_transform_iterator(c_it, ::cuda::std::negate{}); + auto t_it = cuda::make_transform_iterator(c_it, cuda::std::negate{}); c2h::device_vector out_result(1, thrust::no_init); auto predicate = cuda::equal_to_value{-val_to_find}; find_if(t_it, out_result.data(), predicate, num_items); @@ -201,7 +202,7 @@ CUB_TEST("Device find_if works with non primitive iterator", "[device][find_if]" { // counting_iterator input and transform_output_iterator output c2h::device_vector out_result(1, thrust::no_init); auto predicate = cuda::equal_to_value{val_to_find}; - auto out_it = cuda::make_transform_output_iterator(out_result.begin(), ::cuda::std::negate{}); + auto out_it = cuda::make_transform_output_iterator(out_result.begin(), cuda::std::negate{}); find_if(c_it, out_it, predicate, num_items); REQUIRE(-expected_if_found == out_result[0]); } @@ -284,7 +285,7 @@ struct std_lower_bound_t template RangeIteratorT operator()(RangeIteratorT first, RangeIteratorT last, const T& value, CompareOpT comp) const { - return std::lower_bound(first, last, value, comp); + return cuda::std::lower_bound(first, last, value, comp); } } std_lower_bound; @@ -293,7 +294,7 @@ struct std_upper_bound_t template RangeIteratorT operator()(RangeIteratorT first, RangeIteratorT last, const T& value, CompareOpT comp) const { - return std::upper_bound(first, last, value, comp); + return cuda::std::upper_bound(first, last, value, comp); } } std_upper_bound; diff --git a/cub/test/catch2_test_device_reduce_deterministic.cu b/cub/test/catch2_test_device_reduce_deterministic.cu index 434b880b5264..4e2733877aad 100644 --- a/cub/test/catch2_test_device_reduce_deterministic.cu +++ b/cub/test/catch2_test_device_reduce_deterministic.cu @@ -11,8 +11,7 @@ #include #include #include - -#include +#include #include "catch2_test_device_reduce.cuh" #include "cub_test_macros.h" @@ -53,10 +52,9 @@ CUB_TEST("Deterministic Device reduce works with float and double on gpu", c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - // Requires `std::accumulate` to produce deterministic result which is required for comparison + // Requires `cuda::std::accumulate` to produce deterministic result which is required for comparison // with the device RFA result. - // NOTE: `std::reduce` is not equivalent - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus()); + h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus()); REQUIRE_APPROX_EQ_EPSILON(h_expected, d_output, type{0.02}); } @@ -84,9 +82,9 @@ CUB_TEST("Deterministic Device reduce works with float and double on gpu with la CUB_SMALL, large_offset_type_list) { - using type = typename c2h::get<0, TestType>; - const size_t random_num_items = static_cast(cuda::std::numeric_limits<::cuda::std::int32_t>::max()) - + GENERATE_COPY(take(1, random(1, 1000))); + using type = typename c2h::get<0, TestType>; + const size_t random_num_items = + static_cast(cuda::std::numeric_limits::max()) + GENERATE_COPY(take(1, random(1, 1000))); const size_t half_chunk_size = GENERATE_COPY(take(1, random(1, 128))); @@ -105,7 +103,7 @@ CUB_TEST("Deterministic Device reduce works with float and double on gpu with la d_chunk.begin(), d_chunk.begin() + half_chunk_size, d_chunk.begin() + half_chunk_size, - ::cuda::std::negate{}); + cuda::std::negate{}); cyclic_chunk_accessor wrapper{d_chunk.data(), chunk_size}; auto d_input = cuda::transform_iterator(cuda::counting_iterator{}, wrapper); @@ -167,7 +165,7 @@ CUB_TEST("Deterministic Device reduce works with float and double and is determi c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus()); + h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus()); // device RFA result should be approximately equal to host result REQUIRE_APPROX_EQ_EPSILON(h_expected, d_output_p1, type{0.05}); @@ -200,10 +198,10 @@ CUB_TEST("Deterministic Device reduce works with float and double on gpu with di c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - // Requires `std::accumulate` to produce deterministic result which is required for comparison + // Requires `cuda::std::accumulate` to produce deterministic result which is required for comparison // with the device RFA result. - // NOTE: `std::reduce` is not equivalent - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus()); + // NOTE: `cuda::std::reduce` is not equivalent + h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus()); c2h::host_vector h_output = d_output; REQUIRE_APPROX_EQ_EPSILON(h_expected, h_output, type{0.01}); @@ -218,10 +216,10 @@ CUB_TEST("Deterministic Device reduce works with float and double on gpu with di REQUIRE(error == cudaSuccess); c2h::host_vector h_expected(1); - // Requires `std::accumulate` to produce deterministic result which is required for comparison + // Requires `cuda::std::accumulate` to produce deterministic result which is required for comparison // with the device RFA result. - // NOTE: `std::reduce` is not equivalent - h_expected[0] = std::accumulate(input, input + num_items, type{}, cuda::std::plus()); + // NOTE: `cuda::std::reduce` is not equivalent + h_expected[0] = cuda::std::accumulate(input, input + num_items, type{}, cuda::std::plus()); c2h::host_vector h_output = d_output; REQUIRE_APPROX_EQ_EPSILON(h_expected, h_output, type{0.01}); @@ -259,8 +257,8 @@ CUB_TEST("Deterministic Device reduce works with float and double on gpu with di auto h_input = cuda::transform_iterator(input, transform_t{}); c2h::host_vector h_expected(1); - // Requires `std::accumulate` to produce deterministic result which is required for comparison - h_expected[0] = std::accumulate(h_input, h_input + num_items, type{}, cuda::std::plus()); + // Requires `cuda::std::accumulate` to produce deterministic result which is required for comparison + h_expected[0] = cuda::std::accumulate(h_input, h_input + num_items, type{}, cuda::std::plus()); // device RFA result should be approximately equal to host result REQUIRE_APPROX_EQ_EPSILON(h_expected, d_output, type{0.01}); @@ -290,10 +288,10 @@ CUB_TEST("Deterministic Device reduce works with float and double on gpu with di c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - // Requires `std::accumulate` to produce deterministic result which is required for comparison + // Requires `cuda::std::accumulate` to produce deterministic result which is required for comparison // with the device RFA result. - // NOTE: `std::reduce` is not equivalent - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), init_value, cuda::std::plus()); + // NOTE: `cuda::std::reduce` is not equivalent + h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), init_value, cuda::std::plus()); REQUIRE_APPROX_EQ_EPSILON(h_expected, d_output, type{0.01}); } @@ -329,7 +327,7 @@ CUB_TEST("Deterministic Device reduce works with integral types on gpu with diff c2h::device_vector d_input(num_items); c2h::gen(C2H_SEED(2), d_input, min_value, max_value); - if constexpr (::cuda::std::is_integral_v) + if constexpr (cuda::std::is_integral_v) { SECTION("plus") { @@ -342,10 +340,10 @@ CUB_TEST("Deterministic Device reduce works with integral types on gpu with diff c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - // Requires `std::accumulate` to produce deterministic result which is required for comparison + // Requires `cuda::std::accumulate` to produce deterministic result which is required for comparison // with the device RFA result. - // NOTE: `std::reduce` is not equivalent - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), init_value_t{}, cuda::std::plus{}); + // NOTE: `cuda::std::reduce` is not equivalent + h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), init_value_t{}, cuda::std::plus{}); c2h::host_vector h_output = d_output; REQUIRE(h_expected == h_output); @@ -363,7 +361,8 @@ CUB_TEST("Deterministic Device reduce works with integral types on gpu with diff c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::std::bit_xor{}); + h_expected[0] = + cuda::std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::std::bit_xor{}); c2h::host_vector h_output = d_output; REQUIRE(h_expected == h_output); @@ -381,7 +380,8 @@ CUB_TEST("Deterministic Device reduce works with integral types on gpu with diff c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::std::logical_or<>{}); + h_expected[0] = + cuda::std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::std::logical_or<>{}); c2h::host_vector h_output = d_output; REQUIRE(h_expected == h_output); @@ -400,7 +400,7 @@ CUB_TEST("Deterministic Device reduce works with integral types on gpu with diff c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::minimum<>{}); + h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::minimum<>{}); c2h::host_vector h_output = d_output; REQUIRE(h_expected == h_output); @@ -418,7 +418,7 @@ CUB_TEST("Deterministic Device reduce works with integral types on gpu with diff c2h::host_vector h_input = d_input; c2h::host_vector h_expected(1); - h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::maximum<>{}); + h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::maximum<>{}); c2h::host_vector h_output = d_output; REQUIRE(h_expected == h_output); diff --git a/libcudacxx/include/cuda/std/__memory/addressof.h b/libcudacxx/include/cuda/std/__memory/addressof.h index aa650837d4f9..f103621561c5 100644 --- a/libcudacxx/include/cuda/std/__memory/addressof.h +++ b/libcudacxx/include/cuda/std/__memory/addressof.h @@ -39,6 +39,12 @@ # define _CCCL_HAS_BUILTIN_STD_ADDRESSOF() 0 #endif // _CCCL_FREESTANDING() +// In tile mode the builtin is tile annotated which can have unintended consequences in SIMT code with e.g int128 +#if defined(__CUDACC_TILE__) +# undef _CCCL_HAS_BUILTIN_STD_ADDRESSOF +# define _CCCL_HAS_BUILTIN_STD_ADDRESSOF() 0 +#endif // defined(__CUDACC_TILE__) + // 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/__utility/as_const.h b/libcudacxx/include/cuda/std/__utility/as_const.h index f58912717c08..b888ac9edc03 100644 --- a/libcudacxx/include/cuda/std/__utility/as_const.h +++ b/libcudacxx/include/cuda/std/__utility/as_const.h @@ -40,6 +40,12 @@ # define _CCCL_HAS_BUILTIN_STD_AS_CONST() 0 #endif // _CCCL_FREESTANDING() +// In tile mode the builtin is tile annotated which can have unintended consequences in SIMT code with e.g int128 +#if defined(__CUDACC_TILE__) +# undef _CCCL_HAS_BUILTIN_STD_AS_CONST +# define _CCCL_HAS_BUILTIN_STD_AS_CONST() 0 +#endif // defined(__CUDACC_TILE__) + // include minimal std:: headers #if _CCCL_HAS_BUILTIN_STD_AS_CONST() # if _CCCL_HOST_STD_LIB(LIBCXX) && __has_include(<__utility/as_const.h>) diff --git a/libcudacxx/include/cuda/std/__utility/forward.h b/libcudacxx/include/cuda/std/__utility/forward.h index 1b91ca4ee994..9a25547b4f66 100644 --- a/libcudacxx/include/cuda/std/__utility/forward.h +++ b/libcudacxx/include/cuda/std/__utility/forward.h @@ -45,6 +45,12 @@ # define _CCCL_HAS_BUILTIN_STD_FORWARD() 0 #endif // _CCCL_ENABLE_FREESTANDING +// In tile mode the builtin is tile annotated which can have unintended consequences in SIMT code with e.g int128 +#if defined(__CUDACC_TILE__) +# undef _CCCL_HAS_BUILTIN_STD_FORWARD +# define _CCCL_HAS_BUILTIN_STD_FORWARD() 0 +#endif // defined(__CUDACC_TILE__) + // 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() diff --git a/libcudacxx/include/cuda/std/__utility/forward_like.h b/libcudacxx/include/cuda/std/__utility/forward_like.h index f462349207f0..9ef2e78f6c9d 100644 --- a/libcudacxx/include/cuda/std/__utility/forward_like.h +++ b/libcudacxx/include/cuda/std/__utility/forward_like.h @@ -43,6 +43,12 @@ # define _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE() 0 #endif // _CCCL_FREESTANDING() +// In tile mode the builtin is tile annotated which can have unintended consequences in SIMT code with e.g int128 +#if defined(__CUDACC_TILE__) +# undef _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE +# define _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE() 0 +#endif // defined(__CUDACC_TILE__) + // include minimal std:: headers #if _CCCL_HAS_BUILTIN_STD_FORWARD_LIKE() # if _CCCL_HOST_STD_LIB(LIBSTDCXX) && __has_include() diff --git a/libcudacxx/include/cuda/std/__utility/move.h b/libcudacxx/include/cuda/std/__utility/move.h index defff1bf8c45..16116e4812f9 100644 --- a/libcudacxx/include/cuda/std/__utility/move.h +++ b/libcudacxx/include/cuda/std/__utility/move.h @@ -46,6 +46,12 @@ # define _CCCL_HAS_BUILTIN_STD_MOVE() 0 #endif // _CCCL_ENABLE_FREESTANDING +// In tile mode the builtin is tile annotated which can have unintended consequences in SIMT code with e.g int128 +#if defined(__CUDACC_TILE__) +# undef _CCCL_HAS_BUILTIN_STD_MOVE +# define _CCCL_HAS_BUILTIN_STD_MOVE() 0 +#endif // defined(__CUDACC_TILE__) + #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 @@ -64,6 +70,12 @@ # define _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT() 0 #endif // _CCCL_FREESTANDING() +// In tile mode the builtin is tile annotated which can have unintended consequences in SIMT code with e.g int128 +#if defined(__CUDACC_TILE__) +# undef _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT +# define _CCCL_HAS_BUILTIN_STD_MOVE_IF_NOEXCEPT() 0 +#endif // defined(__CUDACC_TILE__) + // 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() 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 eab6eb3c3e68..687895207c9e 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 @@ -207,27 +207,27 @@ void test_host_types() # if _CCCL_HAS_LONG_DOUBLE() test_constructor_from_host_tuple_like(); # endif // _CCCL_HAS_LONG_DOUBLE() -# if _CCCL_HAS_FLOAT128() +# if _CCCL_HAS_FLOAT128() && !defined(__CUDACC_TILE__) test_constructor_from_host_tuple_like<__float128>(); -# endif // _CCCL_HAS_FLOAT128() +# endif // _CCCL_HAS_FLOAT128() && !defined(__CUDACC_TILE__) test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); -# if _CCCL_HAS_INT128() +# if _CCCL_HAS_INT128() && !defined(__CUDACC_TILE__) test_constructor_from_host_tuple_like<__int128_t>(); -# endif // _CCCL_HAS_INT128() +# endif // _CCCL_HAS_INT128() && !defined(__CUDACC_TILE__) test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); test_constructor_from_host_tuple_like(); -# if _CCCL_HAS_INT128() +# if _CCCL_HAS_INT128() && !defined(__CUDACC_TILE__) test_constructor_from_host_tuple_like<__uint128_t>(); -# endif // _CCCL_HAS_INT128() +# endif // _CCCL_HAS_INT128() && !defined(__CUDACC_TILE__) } #endif // _CCCL_HOSTED() diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/expected/expected.void/tile_only_types.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/expected/expected.void/tile_only_types.pass.cpp index b5d2273760ca..ec98075571fe 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/expected/expected.void/tile_only_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/expected/expected.void/tile_only_types.pass.cpp @@ -10,8 +10,7 @@ // REQUIRES: enable-tile || force-tile // We cannot suppress execution checks in cuda::std::construct_at -// UNSUPPORTED: clang-14 -// UNSUPPORTED: enable-tile && !c++17 +// UNSUPPORTED: clang-14 && !c++17 // UNSUPPORTED: force-tile && !c++17 #include @@ -21,8 +20,9 @@ #include "host_device_types.h" #include "test_macros.h" -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() using expected = cuda::std::expected; { // default construction expected default_constructed{}; @@ -160,6 +160,7 @@ TEST_TILE_FUNC void test() assert(lhs.has_value()); assert(rhs.error() == 1337); } +#endif // _CCCL_TILE_COMPILATION() } __tile_global__ void test_kernel() diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/expected/tile_only_types.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/expected/tile_only_types.pass.cpp index 6a427e78e0e8..13a665baca2c 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/expected/tile_only_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/expected/tile_only_types.pass.cpp @@ -10,8 +10,7 @@ // REQUIRES: enable-tile || force-tile // We cannot suppress execution checks in cuda::std::construct_at -// UNSUPPORTED: clang-14 -// UNSUPPORTED: enable-tile && !c++17 +// UNSUPPORTED: clang-14 && !c++17 // UNSUPPORTED: force-tile && !c++17 #include @@ -21,8 +20,9 @@ #include "host_device_types.h" #include "test_macros.h" -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() using expected = cuda::std::expected; { // default construction expected default_constructed{}; @@ -195,6 +195,7 @@ TEST_TILE_FUNC void test() assert(*lhs == 42); assert(rhs.error() == 1337); } +#endif // _CCCL_TILE_COMPILATION() } __tile_global__ void test_kernel() diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/optional/tile_only_types.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/optional/tile_only_types.pass.cpp index 9cf51d3bb8e9..058bf6eb5c4c 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/optional/tile_only_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/optional/tile_only_types.pass.cpp @@ -20,8 +20,9 @@ #include "test_macros.h" template -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() using optional = cuda::std::optional; { // default construction optional default_constructed{}; @@ -133,12 +134,15 @@ TEST_TILE_FUNC void test() assert(*lhs == 42); assert(*rhs == 1337); } +#endif // _CCCL_TILE_COMPILATION() } -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() test(); test(); +#endif // _CCCL_TILE_COMPILATION() } __tile_global__ void test_kernel() diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/tuple/tile_only_types.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/tuple/tile_only_types.pass.cpp index 815479c3019f..b131b662ab8d 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/tuple/tile_only_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/tuple/tile_only_types.pass.cpp @@ -15,8 +15,9 @@ #include "host_device_types.h" #include "test_macros.h" -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() using tuple = cuda::std::tuple; { // default construction tuple default_constructed{}; @@ -74,6 +75,7 @@ TEST_TILE_FUNC void test() assert(cuda::std::get<0>(lhs) == 42); assert(cuda::std::get<0>(rhs) == 1337); } +#endif // _CCCL_TILE_COMPILATION() } __tile_global__ void test_kernel() diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/unexpected/tile_only_types.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/unexpected/tile_only_types.pass.cpp index c294cf7f4402..8dbf3317980a 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/unexpected/tile_only_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/unexpected/tile_only_types.pass.cpp @@ -15,8 +15,9 @@ #include "host_device_types.h" #include "test_macros.h" -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() using unexpected = cuda::std::unexpected; { // in_place zero initialization unexpected in_place_zero_initialization{cuda::std::in_place}; @@ -75,6 +76,7 @@ TEST_TILE_FUNC void test() assert(lhs.error() == 42); assert(rhs.error() == 1337); } +#endif // _CCCL_TILE_COMPILATION() } __tile_global__ void test_kernel() diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/utility/pair/tile_only_types.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/utility/pair/tile_only_types.pass.cpp index 67f07a5f096d..a98e57678f67 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/utility/pair/tile_only_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/utility/pair/tile_only_types.pass.cpp @@ -15,8 +15,9 @@ #include "host_device_types.h" #include "test_macros.h" -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() using pair = cuda::std::pair; { // default construction pair default_constructed{}; @@ -84,6 +85,7 @@ TEST_TILE_FUNC void test() assert(rhs.first == 1337); assert(rhs.second == 42); } +#endif // _CCCL_TILE_COMPILATION() } __tile_global__ void test_kernel() diff --git a/libcudacxx/test/libcudacxx/cuda/utilities/variant/tile_only_types.pass.cpp b/libcudacxx/test/libcudacxx/cuda/utilities/variant/tile_only_types.pass.cpp index 1d5218ca64e3..fffff4cbfb86 100644 --- a/libcudacxx/test/libcudacxx/cuda/utilities/variant/tile_only_types.pass.cpp +++ b/libcudacxx/test/libcudacxx/cuda/utilities/variant/tile_only_types.pass.cpp @@ -15,8 +15,9 @@ #include "host_device_types.h" #include "test_macros.h" -TEST_TILE_FUNC void test() +__tile__ void test() { +#if _CCCL_TILE_COMPILATION() using variant = cuda::std::variant; { // default construction variant default_constructed{}; @@ -113,6 +114,7 @@ TEST_TILE_FUNC void test() assert(cuda::std::get<0>(lhs) == 42); assert(cuda::std::get<0>(rhs) == 1337); } +#endif // _CCCL_TILE_COMPILATION() } __tile_global__ void test_kernel()