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/codegen/generate_prologue_epilogue.py

@@ -54,6 +54,7 @@
         "empty_bases",
         "hybrid_patchable",
         "jitintrinsic",
+        "lifetimebound",
         "naked",
         "noalias",
         "noinline",
@@ -122,12 +123,12 @@ def make_prologue(file):
         for macro in macros:
             write_section(
                 file,
-                f'''\
+                f"""\
 #if defined({macro})
 #  pragma push_macro("{macro}")
 #  undef {macro}
 #  define _CCCL_POP_MACRO_{macro}
-#endif // defined({macro})''',
+#endif // defined({macro})""",
             )
 
     # Add warnings suppressions.
@@ -214,14 +215,14 @@ def make_epilogue(file):
         for macro in macros:
             write_section(
                 file,
-                f'''\
+                f"""\
 #if defined({macro})
 #  error \\
     "cccl internal error: macro `{macro}` was redefined between <cuda/std/__cccl/prologue.h> and <cuda/std/__cccl/epilogue.h>"
 #elif defined(_CCCL_POP_MACRO_{macro})
 #  pragma pop_macro("{macro}")
 #  undef _CCCL_POP_MACRO_{macro}
-#endif''',
+#endif""",
             )
 
     # Write the common footer.

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

@@ -113,6 +113,14 @@ _CCCL_DIAG_POP
 #  undef _CCCL_POP_MACRO_jitintrinsic
 #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(naked)
 #  error \
     "cccl internal error: macro `naked` was redefined between <cuda/std/__cccl/prologue.h> and <cuda/std/__cccl/epilogue.h>"

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

@@ -90,6 +90,12 @@
 #  define _CCCL_POP_MACRO_jitintrinsic
 #endif // defined(jitintrinsic)
 
+#if defined(lifetimebound)
+#  pragma push_macro("lifetimebound")
+#  undef lifetimebound
+#  define _CCCL_POP_MACRO_lifetimebound
+#endif // defined(lifetimebound)
+
 #if defined(naked)
 #  pragma push_macro("naked")
 #  undef naked

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

@@ -0,0 +1,41 @@
+//===----------------------------------------------------------------------===//
+//
+// 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.
+//
+//===----------------------------------------------------------------------===//
+// FORCE_ALL_WARNINGS.
+
+#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;
+}