Add support for [[lifetimebound]]

docs/cccl/development/macro.rst

@@ -294,7 +294,8 @@ Usage example:
 +----------------------------------+------------------------------------------------------------------------------+
 | ``_CCCL_CONST``                  | Portable "constant" function attribute                                       |
 +----------------------------------+------------------------------------------------------------------------------+
-
+| ``_CCCL_LIFETIMEBOUND``          | Portable "lifetime bound" function attribute                                 |
++----------------------------------+------------------------------------------------------------------------------+
 
 **Portable Builtin Macros**:
 

libcudacxx/include/cuda/std/__cccl/attributes.h

@@ -154,6 +154,16 @@
 #define _CCCL_NO_SPECIALIZATIONS \
   _CCCL_NO_SPECIALIZATIONS_BECAUSE("Users are not allowed to specialize this cccl entity")
 
+// _CCCL_LIFETIMEBOUND
+
+#if _CCCL_HAS_CPP_ATTRIBUTE(clang::lifetimebound) || _CCCL_COMPILER(CLANG)
+#  define _CCCL_LIFETIMEBOUND [[clang::lifetimebound]]
+#elif _CCCL_HAS_CPP_ATTRIBUTE(msvc::lifetimebound) || _CCCL_COMPILER(MSVC, >=, 19, 37)
+#  define _CCCL_LIFETIMEBOUND [[msvc::lifetimebound]]
+#else
+#  define _CCCL_LIFETIMEBOUND
+#endif
+
 // _CCCL_NO_UNIQUE_ADDRESS
 
 #if _CCCL_COMPILER(MSVC) || _CCCL_HAS_CPP_ATTRIBUTE(no_unique_address) < 201803L

libcudacxx/include/cuda/std/__cccl/epilogue.h

@@ -105,6 +105,14 @@ _CCCL_DIAG_POP
 #  undef _CCCL_POP_MACRO_hybrid_patchable
 #endif
 
+#if defined(lifetimebound)
+#  error \
+    "cccl internal error: macro `lifetimebound` was redefined between <cuda/std/__cccl/prologue.h> and <cuda/std/__cccl/epilogue.h>"
+#elif defined(_CCCL_POP_MACRO_lifetimebound)
+#  pragma pop_macro("lifetimebound")
+#  undef _CCCL_POP_MACRO_lifetimebound
+#endif
+
 #if defined(jitintrinsic)
 #  error \
     "cccl internal error: macro `jitintrinsic` was redefined between <cuda/std/__cccl/prologue.h> and <cuda/std/__cccl/epilogue.h>"

libcudacxx/include/cuda/std/__cccl/prologue.h

@@ -84,6 +84,12 @@
 #  define _CCCL_POP_MACRO_hybrid_patchable
 #endif // defined(hybrid_patchable)
 
+#if defined(lifetimebound)
+#  pragma push_macro("lifetimebound")
+#  undef lifetimebound
+#  define _CCCL_POP_MACRO_lifetimebound
+#endif // defined(lifetimebound)
+
 #if defined(jitintrinsic)
 #  pragma push_macro("jitintrinsic")
 #  undef jitintrinsic

libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp

@@ -0,0 +1,40 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__cccl/attributes.h>
+
+#include "test_macros.h"
+
+#if _CCCL_HAS_CPP_ATTRIBUTE(clang::lifetimebound) || _CCCL_COMPILER(CLANG)
+#elif _CCCL_HAS_CPP_ATTRIBUTE(msvc::lifetimebound) || _CCCL_COMPILER(MSVC, >=, 19, 37)
+#else
+#  error "lifetimebound attribute not supported"
+#endif
+
+struct S
+{
+  char data[32];
+
+  __host__ __device__ const char* get() const _CCCL_LIFETIMEBOUND
+  {
+    return data;
+  }
+};
+
+__host__ __device__ bool test()
+{
+  auto sv = S{"abc"}.get();
+  return true;
+}
+
+int main(int, char**)
+{
+  assert(test());
+  return 0;
+}

libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.pass.cpp

@@ -0,0 +1,25 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__cccl/attributes.h>
+
+#include "test_macros.h"
+
+struct S
+{
+  __host__ __device__ int get() const _CCCL_LIFETIMEBOUND
+  {
+    return 3;
+  }
+};
+
+int main(int, char**)
+{
+  return 0;
+}