[rocm-libraries] ROCm/rocm-libraries#1720 (commit 0098e5e)

[rocThrust] fix: HIP compiler rejects THRUST_NODISCARD after
 THRUST_DEPRECATED (#1720)

## Motivation

- `THRUST_CPP_DIALECT` is used before being defined causing the code to
use `__attribute__((...))` instead of `[[...]]` under c++ 17 standard,
leading to the compiler throwing an error when putting
`THRUST_DEPRECATED` in front of `THRUST_NODISCARD`. This will be fixed
by this merge/pull request.
- A separate issue is that rocm compiler (both on the host and device
sides) rejects placing `[[nodiscard]]` in front of
`__attribute__((deprecated))` for member functions and operators in
classes. But this syntax is accepted for normal functions. Nvcc accepts
this syntax in either case. In my opinion, this is not strictly a "bug",
but an inconsistency. This can be reproduced on

[godbolt](https://godbolt.org/#g:!((g:!((g:!((h:codeEditor,i:(filename:'1',fontScale:14,fontUsePx:'0',j:1,lang:cuda,selection:(endColumn:1,endLineNumber:6,positionColumn:1,positionLineNumber:6,selectionStartColumn:1,selectionStartLineNumber:6,startColumn:1,startLineNumber:6),source:'%0Astruct+test%7B%0A++++__attribute__((deprecated))+%5B%5Bnodiscard%5D%5D%0A++++int+oeprator()%7B+return+0%3B%7D%0A%7D%3B%0A'),l:'5',n:'0',o:'CUDA+C%2B%2B+source+%231',t:'0')),k:37.61119832548404,l:'4',n:'0',o:'',s:0,t:'0'),(g:!((h:compiler,i:(compiler:hipclang-rocm-60400,filters:(b:'0',binary:'1',binaryObject:'1',commentOnly:'0',debugCalls:'1',demangle:'0',directives:'0',execute:'1',intel:'0',libraryCode:'0',trim:'1',verboseDemangling:'0'),flagsViewOpen:'1',fontScale:14,fontUsePx:'0',j:1,lang:cuda,libs:!(),options:'',overrides:!(),selection:(endColumn:1,endLineNumber:1,positionColumn:1,positionLineNumber:1,selectionStartColumn:1,selectionStartLineNumber:1,startColumn:1,startLineNumber:1),source:1),l:'5',n:'0',o:'+clang+rocm-6.4.0+(Editor+%231)',t:'0')),header:(),k:31.010672395779828,l:'4',n:'0',o:'',s:0,t:'0'),(g:!((h:output,i:(compilerName:'clang+rocm-6.4.0',editorid:1,fontScale:14,fontUsePx:'0',j:1,wrap:'1'),l:'5',n:'0',o:'Output+of+clang+rocm-6.4.0+(Compiler+%231)',t:'0')),k:31.378129278736132,l:'4',n:'0',o:'',s:0,t:'0')),l:'2',n:'0',o:'',t:'0')),version:4)
- `__host__` or `__device__` between 2 attribute is also not allowed by
rocm compiler, but is allow by nvcc. see reproducer on
[godbolt](https://godbolt.org/#g:!((g:!((g:!((h:codeEditor,i:(filename:'1',fontScale:14,fontUsePx:'0',j:1,lang:cuda,selection:(endColumn:1,endLineNumber:6,positionColumn:1,positionLineNumber:6,selectionStartColumn:1,selectionStartLineNumber:6,startColumn:1,startLineNumber:6),source:'%0Astruct+test%7B%0A+++++%5B%5Bnodiscard%5D%5D+__host__+%5B%5Bdeprecated%5D%5D%0A++++int+oeprator()%7B+return+0%3B%7D%0A%7D%3B%0A'),l:'5',n:'0',o:'CUDA+C%2B%2B+source+%231',t:'0')),k:28.20839874411303,l:'4',n:'0',o:'',s:0,t:'0'),(g:!((h:compiler,i:(compiler:hipclang-rocm-60400,filters:(b:'0',binary:'1',binaryObject:'1',commentOnly:'0',debugCalls:'1',demangle:'0',directives:'0',execute:'1',intel:'0',libraryCode:'0',trim:'1',verboseDemangling:'0'),flagsViewOpen:'1',fontScale:14,fontUsePx:'0',j:1,lang:cuda,libs:!(),options:'',overrides:!(),selection:(endColumn:1,endLineNumber:1,positionColumn:1,positionLineNumber:1,selectionStartColumn:1,selectionStartLineNumber:1,startColumn:1,startLineNumber:1),source:1),l:'5',n:'0',o:'+clang+rocm-6.4.0+(Editor+%231)',t:'0')),header:(),k:23.25800429683487,l:'4',n:'0',o:'',s:0,t:'0'),(g:!((h:output,i:(compilerName:'clang+rocm-6.4.0',editorid:1,fontScale:14,fontUsePx:'0',j:1,wrap:'1'),l:'5',n:'0',o:'Output+of+clang+rocm-6.4.0+(Compiler+%231)',t:'0')),k:23.533596959052097,l:'4',n:'0',o:'',s:0,t:'0'),(g:!((h:device,i:(compilerName:'NVCC+12.9.1',device:PTX,editorid:1,fontScale:14,fontUsePx:'0',j:1,selection:(endColumn:1,endLineNumber:1,positionColumn:1,positionLineNumber:1,selectionStartColumn:1,selectionStartLineNumber:1,startColumn:1,startLineNumber:1),treeid:0),l:'5',n:'0',o:'Device+Viewer+clang+rocm-6.4.0+(Editor+%231,+Compiler+%231)',t:'0')),k:25,l:'4',n:'0',o:'',s:0,t:'0')),l:'2',n:'0',o:'',t:'0')),version:4)

## Technical Details

- Fix the include recursion in rocThrust `cpp_dialect.h` and
`deprecated.h`
- Change the order of `THRUST_NODISCARD` and `THRUST_DEPRECATED` for
clarity and consistency with CCCL.

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Author: cenxuantian

thrust/async/copy.h

@@ -115,7 +115,7 @@ struct copy_fn final
           THRUST_FWD(output)))
 
           template <typename... Args>
-          THRUST_NODISCARD THRUST_DEPRECATED THRUST_HOST auto operator()(Args&&... args) const
+          THRUST_DEPRECATED THRUST_NODISCARD THRUST_HOST auto operator()(Args&&... args) const
     THRUST_RETURNS(call(THRUST_FWD(args)...))
 };
 

thrust/async/for_each.h

@@ -92,7 +92,7 @@ struct for_each_fn final
         THRUST_FWD(f)))
 
         template <typename... Args>
-        THRUST_NODISCARD THRUST_DEPRECATED THRUST_HOST auto operator()(Args&&... args) const
+        THRUST_DEPRECATED THRUST_NODISCARD THRUST_HOST auto operator()(Args&&... args) const
     THRUST_RETURNS(call(THRUST_FWD(args)...))
 };
 

thrust/async/reduce.h

@@ -162,7 +162,7 @@ struct reduce_fn final
             ::internal::remove_cvref_t<typename iterator_traits<::internal::remove_cvref_t<ForwardIt>>::value_type>>{}))
 
           template <typename... Args>
-          THRUST_NODISCARD THRUST_DEPRECATED THRUST_HOST auto operator()(Args&&... args) const
+          THRUST_DEPRECATED THRUST_NODISCARD THRUST_HOST auto operator()(Args&&... args) const
     THRUST_RETURNS(call(THRUST_FWD(args)...))
 };
 
@@ -310,7 +310,7 @@ struct reduce_into_fn final
         thrust::is_execution_policy<::internal::remove_cvref_t<T1>>{}))
 
         template <typename... Args>
-        THRUST_NODISCARD THRUST_DEPRECATED THRUST_HOST auto operator()(Args&&... args) const
+        THRUST_DEPRECATED THRUST_NODISCARD THRUST_HOST auto operator()(Args&&... args) const
     THRUST_RETURNS(call(THRUST_FWD(args)...))
 };
 

thrust/async/sort.h

@@ -144,11 +144,12 @@ THRUST_SUPPRESS_DEPRECATED_PUSH
   )
 
   template <typename... Args>
+  THRUST_DEPRECATED
   #if !(defined(__CUDA__) && THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_CLANG)
   // clang in CUDA mode can only handle one attribute
   THRUST_NODISCARD THRUST_HOST
   #endif
- THRUST_DEPRECATED auto operator()(Args&&... args) const
+  auto operator()(Args&&... args) const
   THRUST_RETURNS(
     call(THRUST_FWD(args)...)
   )
@@ -267,11 +268,12 @@ struct sort_fn final
   )
 
   template <typename... Args>
+  THRUST_DEPRECATED
   #if !(defined(__CUDA__) && THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_CLANG)
   // clang in CUDA mode can only handle one attribute
   THRUST_NODISCARD THRUST_HOST
   #endif
- THRUST_DEPRECATED auto operator()(Args&&... args) const
+  auto operator()(Args&&... args) const
   THRUST_RETURNS(
     call(THRUST_FWD(args)...)
   )

thrust/async/transform.h

@@ -123,7 +123,7 @@ struct transform_fn final
   )
 
   template <typename... Args>
-  THRUST_NODISCARD THRUST_DEPRECATED THRUST_HOST
+  THRUST_DEPRECATED THRUST_NODISCARD THRUST_HOST
  auto operator()(Args&&... args) const
   THRUST_RETURNS(
     call(THRUST_FWD(args)...)

thrust/detail/config/cpp_dialect.h

@@ -30,7 +30,6 @@
 #endif // no system header
 
 #include <thrust/detail/config/compiler.h> // IWYU pragma: export
-#include <thrust/detail/config/deprecated.h> // IWYU pragma: export
 
 // Deprecation warnings may be silenced by defining the following macros. These
 // may be combined.
@@ -87,6 +86,38 @@
 
 #endif // !THRUST_CPP_DIALECT
 
+// Macros to suppress deprecation compiler warnings, from "deprecated.h"
+// TODO: These macros start with `LIBCUDACXX`. So, when libhipcxx is
+// available in this scope, we should remove these macros and use the ones
+// from libhipcxx.
+// Check for deprecation opt-outs
+#if defined(LIBCUDACXX_IGNORE_DEPRECATED_CPP_DIALECT) || defined(CCCL_IGNORE_DEPRECATED_CPP_DIALECT) \
+  || defined(CUB_IGNORE_DEPRECATED_CPP_DIALECT)
+#  if !defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT)
+#    define THRUST_IGNORE_DEPRECATED_CPP_DIALECT
+#  endif
+#endif // suppress all dialect deprecation warnings
+#if defined(LIBCUDACXX_IGNORE_DEPRECATED_CPP_14) || defined(CCCL_IGNORE_DEPRECATED_CPP_14) \
+  || defined(CUB_IGNORE_DEPRECATED_CPP_14) || defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT)
+#  if !defined(THRUST_IGNORE_DEPRECATED_CPP_14)
+#    define THRUST_IGNORE_DEPRECATED_CPP_14
+#  endif
+#endif // suppress all c++14 dialect deprecation warnings
+#if defined(LIBCUDACXX_IGNORE_DEPRECATED_CPP_11) || defined(CCCL_IGNORE_DEPRECATED_CPP_11)  \
+  || defined(CUB_IGNORE_DEPRECATED_CPP_11) || defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT) \
+  || defined(THRUST_IGNORE_DEPRECATED_CPP_14)
+#  if !defined(THRUST_IGNORE_DEPRECATED_CPP_11)
+#    define THRUST_IGNORE_DEPRECATED_CPP_11
+#  endif
+#endif // suppress all c++11 dialect deprecation warnings
+#if defined(LIBCUDACXX_IGNORE_DEPRECATED_COMPILER) || defined(CCCL_IGNORE_DEPRECATED_COMPILER) \
+  || defined(CUB_IGNORE_DEPRECATED_COMPILER) || defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT)  \
+  || defined(THRUST_IGNORE_DEPRECATED_CPP_14) || defined(THRUST_IGNORE_DEPRECATED_CPP_11)
+#  if !defined(THRUST_IGNORE_DEPRECATED_COMPILER)
+#    define THRUST_IGNORE_DEPRECATED_COMPILER
+#  endif
+#endif // suppress all compiler deprecation warnings
+
 // Constexpr feature macros:
 #if THRUST_CPP_DIALECT >= 2023
 #    define THRUST_CONSTEXPR_SINCE_CXX23 constexpr
@@ -149,7 +180,7 @@ THRUST_COMPILER_DEPRECATION_SOFT(C++ 17, C++ 11);
 #elif THRUST_CPP_DIALECT == 2014 && !defined(THRUST_IGNORE_DEPRECATED_CPP_14)
 // =C++14. Soft upgrade message:
 THRUST_COMPILER_DEPRECATION_SOFT(C++ 17, C++ 14);
-#endif // THRUST_CPP_DIALECT >= 2017
+#endif // THRUST_CPP_DIALECT < 2011
 
 #undef THRUST_COMPILER_DEPRECATION_SOFT
 #undef THRUST_COMPILER_DEPRECATION

thrust/detail/config/deprecated.h

@@ -12,7 +12,7 @@
 #ifndef THRUST_DETAIL_CONFIG_DEPRECATED_H
 #define THRUST_DETAIL_CONFIG_DEPRECATED_H
 
-#include <thrust/detail/config/compiler.h>
+#include <thrust/detail/config/cpp_dialect.h>
 
 #if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
 #  pragma GCC system_header
@@ -22,33 +22,6 @@
 #  pragma system_header
 #endif // no system header
 
-// Check for deprecation opt-outs
-#if defined(LIBCUDACXX_IGNORE_DEPRECATED_CPP_DIALECT) || defined(CCCL_IGNORE_DEPRECATED_CPP_DIALECT) \
-  || defined(CUB_IGNORE_DEPRECATED_CPP_DIALECT)
-#  if !defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT)
-#    define THRUST_IGNORE_DEPRECATED_CPP_DIALECT
-#  endif
-#endif // suppress all dialect deprecation warnings
-#if defined(LIBCUDACXX_IGNORE_DEPRECATED_CPP_14) || defined(CCCL_IGNORE_DEPRECATED_CPP_14) \
-  || defined(CUB_IGNORE_DEPRECATED_CPP_14) || defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT)
-#  if !defined(THRUST_IGNORE_DEPRECATED_CPP_14)
-#    define THRUST_IGNORE_DEPRECATED_CPP_14
-#  endif
-#endif // suppress all c++14 dialect deprecation warnings
-#if defined(LIBCUDACXX_IGNORE_DEPRECATED_CPP_11) || defined(CCCL_IGNORE_DEPRECATED_CPP_11)  \
-  || defined(CUB_IGNORE_DEPRECATED_CPP_11) || defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT) \
-  || defined(THRUST_IGNORE_DEPRECATED_CPP_14)
-#  if !defined(THRUST_IGNORE_DEPRECATED_CPP_11)
-#    define THRUST_IGNORE_DEPRECATED_CPP_11
-#  endif
-#endif // suppress all c++11 dialect deprecation warnings
-#if defined(LIBCUDACXX_IGNORE_DEPRECATED_COMPILER) || defined(CCCL_IGNORE_DEPRECATED_COMPILER) \
-  || defined(CUB_IGNORE_DEPRECATED_COMPILER) || defined(THRUST_IGNORE_DEPRECATED_CPP_DIALECT)  \
-  || defined(THRUST_IGNORE_DEPRECATED_CPP_14) || defined(THRUST_IGNORE_DEPRECATED_CPP_11)
-#  if !defined(THRUST_IGNORE_DEPRECATED_COMPILER)
-#    define THRUST_IGNORE_DEPRECATED_COMPILER
-#  endif
-#endif // suppress all compiler deprecation warnings
 #if defined(LIBCUDACXX_IGNORE_DEPRECATED_API) || defined(CCCL_IGNORE_DEPRECATED_API) \
   || defined(CUB_IGNORE_DEPRECATED_API)
 #  if !defined(THRUST_IGNORE_DEPRECATED_API)