Add cuda::execution::guarantee and cuda::execution::segment_size::max_segment_size

libcudacxx/include/cuda/__execution/guarantee.h

@@ -0,0 +1,75 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CUDA___EXECUTION_GUARANTEE_H
+#define __CUDA___EXECUTION_GUARANTEE_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__execution/env.h>
+#include <cuda/std/__type_traits/is_base_of.h>
+#include <cuda/std/__type_traits/is_empty.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA_EXECUTION
+
+class __guarantee
+{};
+
+struct __get_guarantees_t
+{
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_TEMPLATE(class _Env)
+  _CCCL_REQUIRES(::cuda::std::execution::__queryable_with<_Env, __get_guarantees_t>)
+  [[nodiscard]] _CCCL_NODEBUG_API constexpr auto operator()(const _Env& __env) const noexcept
+  {
+    static_assert(noexcept(__env.query(*this)));
+    return __env.query(*this);
+  }
+
+  [[nodiscard]]
+  _CCCL_NODEBUG_API static constexpr auto query(::cuda::std::execution::forwarding_query_t) noexcept -> bool
+  {
+    return true;
+  }
+};
+
+_CCCL_GLOBAL_CONSTANT auto __get_guarantees = __get_guarantees_t{};
+
+template <class... _Guarantees>
+[[nodiscard]] _CCCL_NODEBUG_API auto guarantee(_Guarantees...)
+{
+  static_assert((::cuda::std::is_base_of_v<__guarantee, _Guarantees> && ...),
+                "Only guarantees can be passed to require");
+  static_assert((::cuda::std::is_empty_v<_Guarantees> && ...), "Stateful guarantees are not implemented");
+
+  // clang < 19 doesn't like this code
+  // since the only guarantees we currently allow are in max_segment_size.h and
+  // all of them are stateless, let's ignore incoming parameters
+  ::cuda::std::execution::env<_Guarantees...> __env{};
+
+  return ::cuda::std::execution::prop{__get_guarantees_t{}, __env};
+}
+
+_CCCL_END_NAMESPACE_CUDA_EXECUTION
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // __CUDA___EXECUTION_REQUIRE_H

libcudacxx/include/cuda/__execution/max_segment_size.h

@@ -0,0 +1,85 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CUDA___EXECUTION_MAX_SEG_SIZE_H
+#define __CUDA___EXECUTION_MAX_SEG_SIZE_H
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/__execution/guarantee.h>
+#include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__execution/env.h>
+#include <cuda/std/__type_traits/is_one_of.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA_EXECUTION
+
+namespace segment_size
+{
+struct __get_max_segment_size_t;
+
+// TODO : add dynamic extent specialization later, when stateful guarantees (env) are supported
+
+//! A class template that can be used to specify the maximum segment size
+//! for segmented algorithms.
+//! \tparam _Size The maximum segment size.
+template <size_t _Size>
+struct max_segment_size : __guarantee
+{
+  using value_type = size_t;
+
+  constexpr max_segment_size() = default;
+
+  _CCCL_API constexpr operator value_type() const noexcept
+  {
+    return _Size;
+  }
+
+  [[nodiscard]] _CCCL_NODEBUG_API constexpr auto query(const __get_max_segment_size_t&) const noexcept
+  {
+    return *this;
+  }
+};
+
+struct __get_max_segment_size_t
+{
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_TEMPLATE(class _Env)
+  _CCCL_REQUIRES(::cuda::std::execution::__queryable_with<_Env, __get_max_segment_size_t>)
+  [[nodiscard]] _CCCL_NODEBUG_API constexpr auto operator()(const _Env& __env) const noexcept
+  {
+    static_assert(noexcept(__env.query(*this)));
+    return __env.query(*this);
+  }
+
+  [[nodiscard]]
+  _CCCL_NODEBUG_API static constexpr auto query(::cuda::std::execution::forwarding_query_t) noexcept -> bool
+  {
+    return true;
+  }
+};
+
+_CCCL_GLOBAL_CONSTANT auto __get_max_segment_size = __get_max_segment_size_t{};
+} // namespace segment_size
+
+_CCCL_END_NAMESPACE_CUDA_EXECUTION
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // __CUDA___EXECUTION_MAX_SEG_SIZE_H

libcudacxx/test/libcudacxx/cuda/execution/guarantee.fail.cpp

@@ -0,0 +1,29 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/__execution/guarantee.h>
+
+[[maybe_unused]] _CCCL_GLOBAL_CONSTANT struct query_t
+{
+} query{};
+
+__host__ __device__ void test()
+{
+  // not every environment is a requirement
+  cuda::std::execution::prop p{query, 42};
+  cuda::execution::guarantee(p);
+}
+
+int main(int, char**)
+{
+  test();
+
+  return 0;
+}

libcudacxx/test/libcudacxx/cuda/execution/guarantee.pass.cpp

@@ -0,0 +1,26 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/__execution/max_segment_size.h>
+
+__host__ __device__ void test()
+{
+  static_assert(
+    cuda::std::is_same_v<decltype(cuda::execution::segment_size::__get_max_segment_size(cuda::execution::__get_guarantees(
+                           cuda::execution::guarantee(cuda::execution::segment_size::max_segment_size<42>{})))),
+                         cuda::execution::segment_size::max_segment_size<42>>);
+}
+
+int main(int, char**)
+{
+  test();
+
+  return 0;
+}

libcudacxx/test/libcudacxx/cuda/execution/max_segment_size.pass.cpp

@@ -0,0 +1,28 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/__execution/max_segment_size.h>
+
+__host__ __device__ void test()
+{
+  namespace exec = cuda::execution;
+  static_assert(cuda::std::is_base_of_v<exec::__guarantee, exec::segment_size::max_segment_size<42>>);
+
+  static_assert(
+    cuda::std::is_same_v<decltype(exec::segment_size::__get_max_segment_size(exec::segment_size::max_segment_size<42>{})),
+                         exec::segment_size::max_segment_size<42>>);
+}
+
+int main(int, char**)
+{
+  test();
+
+  return 0;
+}