From 70360268a6ac4f7a484ece2be8ec931d5656a3ab Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 9 Jan 2026 11:53:47 -0800 Subject: [PATCH 1/6] Add support for `[[lifetimebound]]` --- .../include/cuda/std/__cccl/attributes.h | 12 +++++++ .../macros/lifetime_bound.compile.fail.cpp | 34 +++++++++++++++++++ .../macros/lifetime_bound.compile.pass.cpp | 25 ++++++++++++++ 3 files changed, 71 insertions(+) create mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp create mode 100644 libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.pass.cpp diff --git a/libcudacxx/include/cuda/std/__cccl/attributes.h b/libcudacxx/include/cuda/std/__cccl/attributes.h index 4d029f6d810a..57b773e40e5c 100644 --- a/libcudacxx/include/cuda/std/__cccl/attributes.h +++ b/libcudacxx/include/cuda/std/__cccl/attributes.h @@ -154,6 +154,18 @@ #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]] +#elif _CCCL_HAS_CPP_ATTRIBUTE(lifetimebound) +# define _CCCL_LIFETIMEBOUND [[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/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 000000000000..24dda24165b9 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp @@ -0,0 +1,34 @@ +//===----------------------------------------------------------------------===// +// +// 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 +{ + 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 000000000000..23f6cc85565c --- /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; +} From af2a8c613ec2942ac11ae22ad816da8af1812ec1 Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 9 Jan 2026 11:57:27 -0800 Subject: [PATCH 2/6] add documentation --- docs/cccl/development/macro.rst | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/docs/cccl/development/macro.rst b/docs/cccl/development/macro.rst index 1f61f2670a1f..c2fbe8c7fa5a 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**: From 12c9415d23d6809da83fbe3a337ca12d0a7a22fa Mon Sep 17 00:00:00 2001 From: fbusato Date: Fri, 9 Jan 2026 12:03:02 -0800 Subject: [PATCH 3/6] disable test when _CCCL_LIFETIMEBOUND is not supported --- .../libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp index 24dda24165b9..ebba56b8e27c 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp @@ -11,6 +11,11 @@ #include "test_macros.h" +#if (_CCCL_LIFETIMEBOUND + 0) +#else +# error "lifetimebound attribute not supported" +#endif + struct S { char data[32]; From 556ec3124aa1338f4e0e2bebae384bea4cc3400e Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 14 Jan 2026 10:18:45 -0800 Subject: [PATCH 4/6] address comments --- libcudacxx/include/cuda/std/__cccl/attributes.h | 4 +--- libcudacxx/include/cuda/std/__cccl/epilogue.h | 8 ++++++++ libcudacxx/include/cuda/std/__cccl/prologue.h | 6 ++++++ .../libcxx/macros/lifetime_bound.compile.fail.cpp | 3 ++- 4 files changed, 17 insertions(+), 4 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cccl/attributes.h b/libcudacxx/include/cuda/std/__cccl/attributes.h index 57b773e40e5c..6721e6c9ee19 100644 --- a/libcudacxx/include/cuda/std/__cccl/attributes.h +++ b/libcudacxx/include/cuda/std/__cccl/attributes.h @@ -158,10 +158,8 @@ #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) +#elif _CCCL_HAS_CPP_ATTRIBUTE(msvc::lifetimebound) || _CCCL_COMPILER(MSVC, >=, 19, 37) # define _CCCL_LIFETIMEBOUND [[msvc::lifetimebound]] -#elif _CCCL_HAS_CPP_ATTRIBUTE(lifetimebound) -# define _CCCL_LIFETIMEBOUND [[lifetimebound]] #else # define _CCCL_LIFETIMEBOUND #endif diff --git a/libcudacxx/include/cuda/std/__cccl/epilogue.h b/libcudacxx/include/cuda/std/__cccl/epilogue.h index b5b790bad99e..9d1b0ad7a6d3 100644 --- a/libcudacxx/include/cuda/std/__cccl/epilogue.h +++ b/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 and " +#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 and " diff --git a/libcudacxx/include/cuda/std/__cccl/prologue.h b/libcudacxx/include/cuda/std/__cccl/prologue.h index 8850045da399..982f7c139675 100644 --- a/libcudacxx/include/cuda/std/__cccl/prologue.h +++ b/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 diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp index ebba56b8e27c..45c87a2f3411 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp @@ -11,7 +11,8 @@ #include "test_macros.h" -#if (_CCCL_LIFETIMEBOUND + 0) +#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 From 52e91843ce685e9a142c1ac088051ecfe68a0155 Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 14 Jan 2026 10:52:43 -0800 Subject: [PATCH 5/6] use generate_prologue_epilogue.py --- libcudacxx/codegen/generate_prologue_epilogue.py | 9 +++++---- libcudacxx/include/cuda/std/__cccl/epilogue.h | 16 ++++++++-------- libcudacxx/include/cuda/std/__cccl/prologue.h | 12 ++++++------ 3 files changed, 19 insertions(+), 18 deletions(-) diff --git a/libcudacxx/codegen/generate_prologue_epilogue.py b/libcudacxx/codegen/generate_prologue_epilogue.py index 6478ed677009..b5a7babf3c55 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/epilogue.h b/libcudacxx/include/cuda/std/__cccl/epilogue.h index 9d1b0ad7a6d3..70a1c9994d87 100644 --- a/libcudacxx/include/cuda/std/__cccl/epilogue.h +++ b/libcudacxx/include/cuda/std/__cccl/epilogue.h @@ -105,14 +105,6 @@ _CCCL_DIAG_POP # undef _CCCL_POP_MACRO_hybrid_patchable #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(jitintrinsic) # error \ "cccl internal error: macro `jitintrinsic` was redefined between and " @@ -121,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 982f7c139675..a15331e391d0 100644 --- a/libcudacxx/include/cuda/std/__cccl/prologue.h +++ b/libcudacxx/include/cuda/std/__cccl/prologue.h @@ -84,18 +84,18 @@ # 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 # 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 From 58608674a044a0980dbfd6fb67b0d9099080e824 Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 14 Jan 2026 13:23:52 -0800 Subject: [PATCH 6/6] add FORCE_ALL_WARNINGS --- .../libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp index 45c87a2f3411..036623e073d1 100644 --- a/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp +++ b/libcudacxx/test/libcudacxx/libcxx/macros/lifetime_bound.compile.fail.cpp @@ -6,6 +6,7 @@ // SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// +// FORCE_ALL_WARNINGS. #include