diff --git a/docs/cccl/development/macro.rst b/docs/cccl/development/macro.rst index 1f61f2670a1..c2fbe8c7fa5 100644 --- a/docs/cccl/development/macro.rst +++ b/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**: diff --git a/libcudacxx/codegen/generate_prologue_epilogue.py b/libcudacxx/codegen/generate_prologue_epilogue.py index 6478ed67700..b5a7babf3c5 100755 --- a/libcudacxx/codegen/generate_prologue_epilogue.py +++ b/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 and " #elif defined(_CCCL_POP_MACRO_{macro}) # pragma pop_macro("{macro}") # undef _CCCL_POP_MACRO_{macro} -#endif''', +#endif""", ) # Write the common footer. diff --git a/libcudacxx/include/cuda/std/__cccl/attributes.h b/libcudacxx/include/cuda/std/__cccl/attributes.h index 4d029f6d810..6721e6c9ee1 100644 --- a/libcudacxx/include/cuda/std/__cccl/attributes.h +++ b/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 diff --git a/libcudacxx/include/cuda/std/__cccl/epilogue.h b/libcudacxx/include/cuda/std/__cccl/epilogue.h index b5b790bad99..70a1c9994d8 100644 --- a/libcudacxx/include/cuda/std/__cccl/epilogue.h +++ b/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 and " +#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 and " diff --git a/libcudacxx/include/cuda/std/__cccl/prologue.h b/libcudacxx/include/cuda/std/__cccl/prologue.h index 8850045da39..a15331e391d 100644 --- a/libcudacxx/include/cuda/std/__cccl/prologue.h +++ b/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 diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp new file mode 100644 index 00000000000..036623e073d --- /dev/null +++ b/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 + +#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; +} diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.pass.cpp new file mode 100644 index 00000000000..23f6cc85565 --- /dev/null +++ b/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 + +#include "test_macros.h" + +struct S +{ + __host__ __device__ int get() const _CCCL_LIFETIMEBOUND + { + return 3; + } +}; + +int main(int, char**) +{ + return 0; +}