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

libcudacxx/include/cuda/__execution/guarantee.h

@@ -0,0 +1,72 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 constexpr auto guarantee(_Guarantees... __guarantees_args) noexcept
+{
+  static_assert((::cuda::std::is_base_of_v<__guarantee, _Guarantees> && ...),
+                "Only guarantees can be passed to guarantee");
+  // static_assert((::cuda::std::is_empty_v<_Guarantees> && ...), "Stateful guarantees are not implemented");
+
+  ::cuda::std::execution::env<_Guarantees...> __env{__guarantees_args...};
+
+  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,110 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+
+struct __get_max_segment_size_t;
+
+_CCCL_GLOBAL_CONSTANT auto dynamic_max_segment_size = static_cast<size_t>(-1);
+
+//! 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 = dynamic_max_segment_size>
+struct max_segment_size : __guarantee
+{
+  using value_type             = size_t;
+  static constexpr size_t size = _Size;
+
+  constexpr max_segment_size() = default;
+
+  _CCCL_API constexpr max_segment_size(size_t) noexcept {}
+
+  _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;
+  }
+};
+
+template <>
+struct max_segment_size<dynamic_max_segment_size> : __guarantee
+{
+  using value_type = size_t;
+
+  static constexpr size_t size = dynamic_max_segment_size;
+
+  _CCCL_API constexpr max_segment_size(size_t __s)
+      : __val(__s)
+  {}
+
+  _CCCL_API constexpr operator value_type() const noexcept
+  {
+    return __val;
+  }
+
+  [[nodiscard]] _CCCL_NODEBUG_API constexpr auto query(const __get_max_segment_size_t&) const noexcept
+  {
+    return *this;
+  }
+
+private:
+  size_t __val;
+};
+
+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{};
+
+_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,38 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+#include <cuda/__execution/max_segment_size.h>
+#include <cuda/std/__execution/env.h>
+
+__host__ __device__ void test(size_t dynamic_val)
+{
+  // Test for state-less guarantee
+  auto static_genv           = cuda::execution::guarantee(cuda::execution::max_segment_size<42>{});
+  auto static_env            = ::cuda::std::execution::env{static_genv};
+  auto static_genv_extracted = cuda::execution::__get_guarantees(static_env);
+  (void) static_genv_extracted;
+
+  // Test for stateful guarantee
+  auto dynamic_genv           = cuda::execution::guarantee(cuda::execution::max_segment_size<>{dynamic_val});
+  auto dynamic_env            = ::cuda::std::execution::env{dynamic_genv};
+  auto dynamic_genv_extracted = cuda::execution::__get_guarantees(dynamic_env);
+  (void) dynamic_genv_extracted;
+
+  // Test that max_segment_size is a guarantee
+  static_assert(cuda::std::is_base_of_v<cuda::execution::__guarantee, cuda::execution::max_segment_size<42>>);
+}
+
+int main(int argc, char**)
+{
+  test(argc);
+
+  return 0;
+}

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

@@ -0,0 +1,58 @@
+//===----------------------------------------------------------------------===//
+//
+// 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>
+#include <cuda/__execution/max_segment_size.h>
+#include <cuda/std/__execution/env.h>
+
+__host__ __device__ void test(size_t dynamic_val)
+{
+  namespace exec = cuda::execution;
+  static_assert(cuda::std::is_base_of_v<exec::__guarantee, exec::max_segment_size<42>>);
+
+  static_assert(cuda::std::is_same_v<decltype(exec::__get_max_segment_size(exec::max_segment_size<42>{})),
+                                     exec::max_segment_size<42>>);
+
+  static_assert(cuda::std::is_same_v<decltype(exec::__get_max_segment_size(exec::max_segment_size<>{dynamic_val})),
+                                     exec::max_segment_size<exec::dynamic_max_segment_size>>);
+
+  static_assert(
+    cuda::std::is_same_v<
+      decltype(exec::__get_max_segment_size(exec::__get_guarantees(exec::guarantee(exec::max_segment_size<42>{})))),
+      exec::max_segment_size<42>>);
+
+  static_assert(cuda::std::is_same_v<decltype(exec::__get_max_segment_size(
+                                       exec::__get_guarantees(exec::guarantee(exec::max_segment_size<>{dynamic_val})))),
+                                     exec::max_segment_size<exec::dynamic_max_segment_size>>);
+
+  constexpr exec::max_segment_size<42> static_size{};
+  static_assert(static_cast<int>(static_size) == 42);
+
+  constexpr exec::max_segment_size<42> static_size_with_runtime_value{100};
+  // ignore runtime value in case of static size
+  static_assert(static_cast<int>(static_size_with_runtime_value) == 42);
+
+  exec::max_segment_size dynamic_size{dynamic_val};
+  (void) (static_cast<size_t>(dynamic_size) == dynamic_val);
+
+  auto g_env = exec::guarantee(dynamic_size);
+  (void) g_env;
+
+  auto dynamic_size_extracted =
+    ::cuda::std::execution::__query_or(g_env, exec::__get_max_segment_size_t{}, exec::max_segment_size<0>{});
+  (void) (static_cast<size_t>(dynamic_size_extracted) == dynamic_val);
+}
+
+int main(int argc, char**)
+{
+  test(argc);
+
+  return 0;
+}