Skip to content

Commit 1f48a0e

Browse files
committed
[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 llvm#72394 sine the template specialization is treated as overload due to implicit host/device attrs are considered for overload/override differentiation.
1 parent c4c52d4 commit 1f48a0e

File tree

3 files changed

+21
-2
lines changed

3 files changed

+21
-2
lines changed

clang/lib/Sema/SemaOverload.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1491,8 +1491,10 @@ static bool IsOverloadOrOverrideImpl(Sema &SemaRef, FunctionDecl *New,
14911491
// Don't allow overloading of destructors. (In theory we could, but it
14921492
// would be a giant change to clang.)
14931493
if (!isa<CXXDestructorDecl>(New)) {
1494-
Sema::CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget(New),
1495-
OldTarget = SemaRef.IdentifyCUDATarget(Old);
1494+
Sema::CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget(
1495+
New, isa<CXXConstructorDecl>(New)),
1496+
OldTarget = SemaRef.IdentifyCUDATarget(
1497+
Old, isa<CXXConstructorDecl>(New));
14961498
if (NewTarget != Sema::CFT_InvalidTarget) {
14971499
assert((OldTarget != Sema::CFT_InvalidTarget) &&
14981500
"Unexpected invalid target.");

clang/test/SemaCUDA/implicit-member-target-inherited.cu

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -39,6 +39,7 @@ struct A2_with_device_ctor {
3939
};
4040
// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
4141
// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
42+
// expected-note@-4 {{candidate inherited constructor not viable: call to __device__ function from __host__ function}}
4243

4344
struct B2_with_implicit_default_ctor : A2_with_device_ctor {
4445
using A2_with_device_ctor::A2_with_device_ctor;

clang/test/SemaCUDA/trivial-ctor-dtor.cu

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,3 +38,19 @@ struct TC : TB<T> {
3838
};
3939

4040
__device__ TC<int> tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
41+
42+
// Check trivial ctor specialization
43+
template <typename T>
44+
struct C { //expected-note {{candidate constructor (the implicit copy constructor) not viable}}
45+
//expected-note@-1 {{candidate constructor (the implicit move constructor) not viable}}
46+
explicit C() {};
47+
};
48+
49+
template <> C<int>::C() {};
50+
__device__ C<int> ci_d;
51+
C<int> ci_h;
52+
53+
// Check non-trivial ctor specialization
54+
template <> C<float>::C() { static int nontrivial_ctor = 1; } //expected-note {{candidate constructor not viable: call to __host__ function from __device__ function}}
55+
__device__ C<float> cf_d; //expected-error {{no matching constructor for initialization of 'C<float>'}}
56+
C<float> cf_h;

0 commit comments

Comments
 (0)