[CUDA][HIP] ignore implicit host/device attr for override

[yxsamliu]

When deciding whether a previous function declaration is an overload or override, implicit host/device attrs should not be considered.

This fixes the failure for the following code:

`template
class C {
explicit C() {};
};

template <> C::C() {};
`

The issue was introduced by #72394

sine the template specialization is treated as overload due to implicit host/device attrs are considered for overload/override differentiation.

[llvmbot]

@llvm/pr-subscribers-clang

Author: Yaxun (Sam) Liu (yxsamliu)

Changes

When deciding whether a previous function declaration is an overload or override, implicit host/device attrs should not be considered.

This fixes the failure for the following code:

`template <typename T>
class C {
explicit C() {};
};

template <> C<int>::C() {};
`

The issue was introduced by #72394

sine the template specialization is treated as overload due to implicit host/device attrs are considered for overload/override differentiation.

---

Full diff: https://github.com/llvm/llvm-project/pull/72815.diff

4 Files Affected:

• (modified) clang/lib/Sema/SemaCUDA.cpp (+2-6)
• (modified) clang/lib/Sema/SemaOverload.cpp (+4-2)
• (modified) clang/test/SemaCUDA/implicit-member-target-inherited.cu (+1)
• (modified) clang/test/SemaCUDA/trivial-ctor-dtor.cu (+16)

diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index b94f448dabe7517..9dc63ff315a3926 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -1000,13 +1000,9 @@ void Sema::checkCUDATargetOverload(FunctionDecl *NewFD,
     // should have the same implementation on both sides.
     if (NewTarget != OldTarget &&
         ((NewTarget == CFT_HostDevice &&
-          !(LangOpts.OffloadImplicitHostDeviceTemplates &&
-            isCUDAImplicitHostDeviceFunction(NewFD) &&
-            OldTarget == CFT_Device)) ||
+          !isCUDAImplicitHostDeviceFunction(NewFD)) ||
          (OldTarget == CFT_HostDevice &&
-          !(LangOpts.OffloadImplicitHostDeviceTemplates &&
-            isCUDAImplicitHostDeviceFunction(OldFD) &&
-            NewTarget == CFT_Device)) ||
+          !isCUDAImplicitHostDeviceFunction(OldFD)) ||
          (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) &&
         !IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
                     /* ConsiderCudaAttrs = */ false)) {
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 9800d7f1c9cfee9..a33acc184b75274 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -1491,8 +1491,10 @@ static bool IsOverloadOrOverrideImpl(Sema &SemaRef, FunctionDecl *New,
     // Don't allow overloading of destructors.  (In theory we could, but it
     // would be a giant change to clang.)
     if (!isa<CXXDestructorDecl>(New)) {
-      Sema::CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget(New),
-                               OldTarget = SemaRef.IdentifyCUDATarget(Old);
+      Sema::CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget(
+                                   New, /*IgnoreImplicitHDAttr=*/true),
+                               OldTarget = SemaRef.IdentifyCUDATarget(
+                                   Old, /*IgnoreImplicitHDAttr=*/true);
       if (NewTarget != Sema::CFT_InvalidTarget) {
         assert((OldTarget != Sema::CFT_InvalidTarget) &&
                "Unexpected invalid target.");
diff --git a/clang/test/SemaCUDA/implicit-member-target-inherited.cu b/clang/test/SemaCUDA/implicit-member-target-inherited.cu
index 781199bba6b5a11..ceca0891fc9b03c 100644
--- a/clang/test/SemaCUDA/implicit-member-target-inherited.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-inherited.cu
@@ -39,6 +39,7 @@ struct A2_with_device_ctor {
 };
 // expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
 // expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
+// expected-note@-4 {{candidate inherited constructor not viable: call to __device__ function from __host__ function}}
 
 struct B2_with_implicit_default_ctor : A2_with_device_ctor {
   using A2_with_device_ctor::A2_with_device_ctor;
diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
index 1df8adc62bab590..21d698d28492ac3 100644
--- a/clang/test/SemaCUDA/trivial-ctor-dtor.cu
+++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
@@ -38,3 +38,19 @@ struct TC : TB<T> {
 };
 
 __device__ TC<int> tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+
+// Check trivial ctor specialization
+template <typename T>
+struct C { //expected-note {{candidate constructor (the implicit copy constructor) not viable}}
+           //expected-note@-1 {{candidate constructor (the implicit move constructor) not viable}}
+    explicit C() {};
+};
+
+template <> C<int>::C() {};
+__device__ C<int> ci_d;
+C<int> ci_h;
+
+// Check non-trivial ctor specialization
+template <> C<float>::C() { static int nontrivial_ctor = 1; } //expected-note {{candidate constructor not viable: call to __host__ function from __device__ function}}
+__device__ C<float> cf_d; //expected-error {{no matching constructor for initialization of 'C<float>'}}
+C<float> cf_h;

[Artem-B]

LGTM, with one question.