Skip to content

Commit

Permalink
[CUDA][HIP] ignore implicit host/device attr for override (#72815)
Browse files Browse the repository at this point in the history
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.
  • Loading branch information
yxsamliu authored Nov 20, 2023
1 parent f42482d commit a1e2c65
Show file tree
Hide file tree
Showing 3 changed files with 21 additions and 2 deletions.
6 changes: 4 additions & 2 deletions clang/lib/Sema/SemaOverload.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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.");
Expand Down
1 change: 1 addition & 0 deletions clang/test/SemaCUDA/implicit-member-target-inherited.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
16 changes: 16 additions & 0 deletions clang/test/SemaCUDA/trivial-ctor-dtor.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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;

0 comments on commit a1e2c65

Please sign in to comment.