https://github.com/yxsamliu updated
https://github.com/llvm/llvm-project/pull/72815
>From 1f48a0eb04df6ccc2d8117ebd6b68cc45484a0f3 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun....@amd.com>
Date: Sun, 19 Nov 2023 19:34:35 -0500
Subject: [PATCH] [CUDA][HIP] ignore implicit host/device attr for override
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 https://github.com/llvm/llvm-project/pull/72394
sine the template specialization is treated as overload
due to implicit host/device attrs are considered for
overload/override differentiation.
---
clang/lib/Sema/SemaOverload.cpp | 6 ++++--
.../SemaCUDA/implicit-member-target-inherited.cu | 1 +
clang/test/SemaCUDA/trivial-ctor-dtor.cu | 16 ++++++++++++++++
3 files changed, 21 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 9800d7f1c9cfee9..64607e28b8b35e6 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, isa<CXXConstructorDecl>(New)),
+ OldTarget = SemaRef.IdentifyCUDATarget(
+ Old, isa<CXXConstructorDecl>(New));
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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
