[Tile] Disable STD builtins in tile mode

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

cub/test/catch2_test_device_find.cu

@@ -10,6 +10,7 @@
 #include <thrust/tabulate.h>
 
 #include <cuda/iterator>
+#include <cuda/std/algorithm>
 
 #include "catch2_test_device_reduce.cuh"
 #include "catch2_test_launch_helper.h"
@@ -43,8 +44,8 @@ enum class gen_data_t
 template <typename OffsetT, typename InputIt, typename Predicate>
 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<OffsetT>(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<OffsetT>(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<input_t>(::cuda::std::numeric_limits<input_t>::max() - 1));
+      c2h::gen(C2H_SEED(1), in_items, input_t{0}, static_cast<input_t>(cuda::std::numeric_limits<input_t>::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<input_t>::max();
+    val_to_find = cuda::std::numeric_limits<input_t>::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<offset_t> out_result(1, thrust::no_init);
     auto predicate = cuda::equal_to_value<input_t>{-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<offset_t> out_result(1, thrust::no_init);
     auto predicate = cuda::equal_to_value<input_t>{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 <typename RangeIteratorT, typename T, typename CompareOpT>
   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 <typename RangeIteratorT, typename T, typename CompareOpT>
   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;
 

cub/test/catch2_test_device_reduce_deterministic.cu

@@ -11,8 +11,7 @@
 #include <cuda/__execution/determinism.h>
 #include <cuda/__execution/require.h>
 #include <cuda/iterator>
-
-#include <numeric>
+#include <cuda/std/numeric>
 
 #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<type> h_input = d_input;
 
   c2h::host_vector<type> 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<type>());
+  h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus<type>());
 
   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<size_t>(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<size_t>(cuda::std::numeric_limits<cuda::std::int32_t>::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<type>{});
+    cuda::std::negate<type>{});
 
   cyclic_chunk_accessor<type, decltype(d_chunk.data())> wrapper{d_chunk.data(), chunk_size};
   auto d_input = cuda::transform_iterator(cuda::counting_iterator<size_t>{}, wrapper);
@@ -167,7 +165,7 @@ CUB_TEST("Deterministic Device reduce works with float and double and is determi
 
   c2h::host_vector<type> h_input = d_input;
   c2h::host_vector<type> h_expected(1);
-  h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus<type>());
+  h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus<type>());
 
   // 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<type> h_input = d_input;
 
     c2h::host_vector<type> 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<type>());
+    // NOTE: `cuda::std::reduce` is not equivalent
+    h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), type{}, cuda::std::plus<type>());
     c2h::host_vector<type> 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<type> 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<type>());
+    // NOTE: `cuda::std::reduce` is not equivalent
+    h_expected[0] = cuda::std::accumulate(input, input + num_items, type{}, cuda::std::plus<type>());
 
     c2h::host_vector<type> 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<type> 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<type>());
+  // 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<type>());
 
   // 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<type> h_input = d_input;
   c2h::host_vector<type> 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<type>());
+  // NOTE: `cuda::std::reduce` is not equivalent
+  h_expected[0] = cuda::std::accumulate(h_input.begin(), h_input.end(), init_value, cuda::std::plus<type>());
 
   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<type> d_input(num_items);
   c2h::gen(C2H_SEED(2), d_input, min_value, max_value);
 
-  if constexpr (::cuda::std::is_integral_v<type>)
+  if constexpr (cuda::std::is_integral_v<type>)
   {
     SECTION("plus")
     {
@@ -342,10 +340,10 @@ CUB_TEST("Deterministic Device reduce works with integral types on gpu with diff
       c2h::host_vector<type> h_input = d_input;
 
       c2h::host_vector<type> 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<type>{});
+      // 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<type>{});
 
       c2h::host_vector<type> 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<type> h_input = d_input;
       c2h::host_vector<type> h_expected(1);
-      h_expected[0] = std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::std::bit_xor<type>{});
+      h_expected[0] =
+        cuda::std::accumulate(h_input.begin(), h_input.end(), type{init_value}, cuda::std::bit_xor<type>{});
 
       c2h::host_vector<type> 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<type> h_input = d_input;
       c2h::host_vector<type> 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<type> 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<type> h_input = d_input;
     c2h::host_vector<type> 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<type> 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<type> h_input = d_input;
     c2h::host_vector<type> 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<type> h_output = d_output;
     REQUIRE(h_expected == h_output);

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(<bits/move.h>)

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>)

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(<bits/move.h>)

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(<bits/move.h>)

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(<bits/move.h>)

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<long double>();
 #  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<signed char>();
   test_constructor_from_host_tuple_like<signed short>();
   test_constructor_from_host_tuple_like<signed int>();
   test_constructor_from_host_tuple_like<signed long>();
   test_constructor_from_host_tuple_like<signed long long>();
-#  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<unsigned char>();
   test_constructor_from_host_tuple_like<unsigned short>();
   test_constructor_from_host_tuple_like<unsigned int>();
   test_constructor_from_host_tuple_like<unsigned long>();
   test_constructor_from_host_tuple_like<unsigned long long>();
-#  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()

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 <cuda/std/cassert>
@@ -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<void, tile_only_type>;
   { // 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()