Revert "[CUDA][HIP] Fix CTAD for host/device constructors (llvm#168711)"

breaks rocRAND

This reverts commit e719e93.

Author: ronlieb

clang/docs/HIPSupport.rst

@@ -412,54 +412,6 @@ Example Usage
    __host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));
    __host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));
 
-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 semantically behave
-as ``__host__ __device__``. This ensures they are available for deduction in both
-host and device contexts, and CTAD continues to respect any constraints on the
-corresponding constructors in the usual C++ way.
-
-.. code-block:: c++
-
-   template <typename T>
-   struct MyType {
-     T value;
-     __device__ MyType(T v) : value(v) {}
-   };
-
-   MyType(float) -> MyType<double>;
-
-   __device__ void deviceFunc() {
-     MyType m(1.0f); // Deduces MyType<double>
-   }
-
-.. note::
-
-   Explicit HIP target attributes such as ``__host__`` or ``__device__``
-   are not allowed on deduction guides. Clang treats all deduction guides
-   as if they were ``__host__ __device__`` and diagnoses any explicit
-   target attributes on them as errors.
 
 Host and Device Attributes of Default Destructors
 ===================================================

clang/docs/ReleaseNotes.rst

@@ -897,20 +897,6 @@ RISC-V Support
 CUDA/HIP Language Changes
 ^^^^^^^^^^^^^^^^^^^^^^^^^
 
-- Clang now supports C++17 Class Template Argument Deduction (CTAD) in CUDA/HIP
-  device code by treating deduction guides as if they were ``__host__ __device__``.
-
-- Clang avoids ambiguous CTAD in CUDA/HIP by not synthesizing duplicate implicit
-  deduction guides when ``__host__`` and ``__device__`` constructors differ only
-  in CUDA target attributes (same signature and constraints).
-
-- Clang diagnoses CUDA/HIP target attributes written on deduction guides as errors,
-  since deduction guides do not participate in code generation.
-
-- Clang preserves distinct implicit deduction guides for constructors that differ
-  by constraints, so constraint-based CTAD works in CUDA/HIP device code as in
-  standard C++.
-
 CUDA Support
 ^^^^^^^^^^^^
 

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -2765,9 +2765,6 @@ def err_deduction_guide_name_not_class_template : Error<
   "cannot specify deduction guide for "
   "%select{<error>|function template|variable template|alias template|"
   "template template parameter|concept|dependent template name}0 %1">;
-def err_deduction_guide_target_attr : Error<
-  "in CUDA/HIP, target attributes are not allowed on deduction guides; "
-  "deduction guides are implicitly enabled for both host and device">;
 def err_deduction_guide_wrong_scope : Error<
   "deduction guide must be declared in the same scope as template %q0">;
 def err_deduction_guide_defines_function : Error<

clang/lib/Sema/SemaCUDA.cpp

@@ -215,12 +215,6 @@ CUDAFunctionTarget SemaCUDA::IdentifyTarget(const FunctionDecl *D,
   if (D == nullptr)
     return CurCUDATargetCtx.Target;
 
-  // C++ deduction guides are never codegen'ed and only participate in template
-  // argument deduction.  Treat them as if they were always host+device so that
-  // CUDA/HIP target checking never rejects their use based solely on target.
-  if (isa<CXXDeductionGuideDecl>(D))
-    return CUDAFunctionTarget::HostDevice;
-
   if (D->hasAttr<CUDAInvalidTargetAttr>())
     return CUDAFunctionTarget::InvalidTarget;
 
@@ -992,12 +986,6 @@ bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) {
   if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated())
     return true;
 
-  // C++ deduction guides participate in overload resolution but are not
-  // callable functions and are never codegen'ed.  Treat them as always
-  // allowed for CUDA/HIP compatibility checking.
-  if (isa<CXXDeductionGuideDecl>(Callee))
-    return true;
-
   // FIXME: Is bailing out early correct here?  Should we instead assume that
   // the caller is a global initializer?
   FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);

clang/lib/Sema/SemaDeclAttr.cpp

@@ -7987,19 +7987,6 @@ void Sema::ProcessDeclAttributeList(
     }
   }
 
-  // CUDA/HIP: disallow explicit CUDA target attributes on deduction guides.
-  // Deduction guides are not callable functions and never participate in
-  // codegen; they are always treated as host+device for CUDA/HIP semantic
-  // checks, so explicit target attributes on them would be misleading.
-  if (getLangOpts().CUDA)
-    if (auto *Guide = dyn_cast<CXXDeductionGuideDecl>(D);
-        Guide &&
-        (Guide->hasAttr<CUDAHostAttr>() || Guide->hasAttr<CUDADeviceAttr>() ||
-         Guide->hasAttr<CUDAGlobalAttr>())) {
-      Diag(Guide->getLocation(), diag::err_deduction_guide_target_attr);
-      Guide->setInvalidDecl();
-    }
-
   // Do not permit 'constructor' or 'destructor' attributes on __device__ code.
   if (getLangOpts().CUDAIsDevice && D->hasAttr<CUDADeviceAttr>() &&
       (D->hasAttr<ConstructorAttr>() || D->hasAttr<DestructorAttr>()) &&

clang/lib/Sema/SemaTemplateDeductionGuide.cpp

@@ -54,26 +54,6 @@ using namespace clang;
 using namespace sema;
 
 namespace {
-
-/// Return true if two associated-constraint sets are semantically equal.
-static bool HaveSameAssociatedConstraints(
-    Sema &SemaRef, const NamedDecl *Old, ArrayRef<AssociatedConstraint> OldACs,
-    const NamedDecl *New, ArrayRef<AssociatedConstraint> NewACs) {
-  if (OldACs.size() != NewACs.size())
-    return false;
-  if (OldACs.empty())
-    return true;
-
-  // General case: pairwise compare each associated constraint expression.
-  Sema::TemplateCompareNewDeclInfo NewInfo(New);
-  for (size_t I = 0, E = OldACs.size(); I != E; ++I)
-    if (!SemaRef.AreConstraintExpressionsEqual(
-            Old, OldACs[I].ConstraintExpr, NewInfo, NewACs[I].ConstraintExpr))
-      return false;
-
-  return true;
-}
-
 /// Tree transform to "extract" a transformed type from a class template's
 /// constructor to a deduction guide.
 class ExtractTypeForDeductionGuide
@@ -238,51 +218,9 @@ buildDeductionGuide(Sema &SemaRef, TemplateDecl *OriginalTemplate,
       TInfo->getTypeLoc().castAs<FunctionProtoTypeLoc>().getParams();
 
   // Build the implicit deduction guide template.
-  QualType GuideType = TInfo->getType();
-
-  // In CUDA/HIP mode, avoid duplicate implicit guides that differ only in CUDA
-  // target attributes (same constructor signature and constraints).
-  if (IsImplicit && Ctor && SemaRef.getLangOpts().CUDA) {
-    SmallVector<AssociatedConstraint, 4> NewACs;
-    Ctor->getAssociatedConstraints(NewACs);
-
-    for (NamedDecl *Existing : DC->lookup(DeductionGuideName)) {
-      auto *ExistingFT = dyn_cast<FunctionTemplateDecl>(Existing);
-      auto *ExistingGuide =
-          ExistingFT
-              ? dyn_cast<CXXDeductionGuideDecl>(ExistingFT->getTemplatedDecl())
-              : dyn_cast<CXXDeductionGuideDecl>(Existing);
-      if (!ExistingGuide)
-        continue;
-
-      // Only consider guides that were also synthesized from a constructor.
-      auto *ExistingCtor = ExistingGuide->getCorrespondingConstructor();
-      if (!ExistingCtor)
-        continue;
-
-      // If the underlying constructors are overloads (different signatures once
-      // CUDA attributes are ignored), they should each get their own guides.
-      if (SemaRef.IsOverload(Ctor, ExistingCtor,
-                             /*UseMemberUsingDeclRules=*/false,
-                             /*ConsiderCudaAttrs=*/false))
-        continue;
-
-      // At this point, the constructors have the same signature ignoring CUDA
-      // attributes. Decide whether their associated constraints are also the
-      // same; only in that case do we treat one guide as a duplicate of the
-      // other.
-      SmallVector<AssociatedConstraint, 4> ExistingACs;
-      ExistingCtor->getAssociatedConstraints(ExistingACs);
-
-      if (HaveSameAssociatedConstraints(SemaRef, ExistingCtor, ExistingACs,
-                                        Ctor, NewACs))
-        return Existing;
-    }
-  }
-
   auto *Guide = CXXDeductionGuideDecl::Create(
-      SemaRef.Context, DC, LocStart, ES, Name, GuideType, TInfo, LocEnd, Ctor,
-      DeductionCandidate::Normal, FunctionTrailingRC);
+      SemaRef.Context, DC, LocStart, ES, Name, TInfo->getType(), TInfo, LocEnd,
+      Ctor, DeductionCandidate::Normal, FunctionTrailingRC);
   Guide->setImplicit(IsImplicit);
   Guide->setParams(Params);