[clang] [CUDA][HIP] Fix CTAD for host/device constructors (PR #168711)

[Artem Belevich via cfe-commits]

================
@@ -287,6 +287,46 @@ Example Usage
       basePtr->virtualFunction(); // Allowed since obj is constructed in device code
    }
 
+C++17 Class Template Argument Deduction (CTAD) Support
+======================================================
+
+Clang supports C++17 Class Template Argument Deduction (CTAD) in both host and device code for HIP.
+This allows you to omit template arguments when creating class template instances, letting the compiler
+deduce them from constructor arguments.
+
+.. code-block:: c++
+
+   #include <tuple>
+
+   __host__ __device__ void func() {
+     std::tuple<int, int> t = std::tuple(1, 1);
+   }
+
+In the above example, ``std::tuple(1, 1)`` automatically deduces the type to be ``std::tuple<int, int>``.
+
+Deduction Guides
+----------------
+
+User-defined deduction guides are also supported. Since deduction guides are not executable code and only
+participate in type deduction, they are treated as ``__host__ __device__`` by the compiler, regardless of
+explicit target attributes. This ensures they are available for deduction in both host and device contexts.
+
+.. code-block:: c++
+
+   template <typename T>
+   struct MyType {
+     T value;
+     MyType(T v) : value(v) {}
+   };
+
+   // User-defined deduction guide
+   template <typename T>
+   MyType(T) -> MyType<T>;
----------------
Artem-B wrote:

With `__device__` constructor, this example compiler regardless of the target attributes of the guide. https://godbolt.org/z/nxhGjh7dr

To reproduce the same issue you demonstrated with std::tuple deviceFunc must also be HD. Or, perhaps the example should be a bit more complex.

In any case, the immediate issue we need to deal with is deduction failures on attempts to use host-side deduction guides supplied by the standard library from D or HD functions. A big-picture issue is to figure out a consistent way to mix target attributes and deduction guides.

Proposed solution is to treat the guides as HD, and *ignore* explicit target attributes. It allows the code to compile, but ignoring target attributes is the best idea. If we do not want to use them for the guides, then we should diagnose it. 

I'm not sure if there's any existing CUDA code that may rely on the deduction guides with attributes. @miscco -- you may be the best person to answer this. Would target-specific deduction guides be useful in practice? 

My preference would be, in order:
a) do not allow target attributes, and treat guides as HD. It's simple and unambiguous.
b) treat guides as regular functions, so target attributes apply. Treat guides in system headers as HD. I think we already have pragmas for doing that for other standard library parts. This allows standard library to work, but keeps target attribute use in the user code consistent between guides and functions.
c) always treat guides as HD, but diagnose as a warning when an explicit host or device attribute is used.










https://github.com/llvm/llvm-project/pull/168711