Author: Yaxun (Sam) Liu Date: 2023-11-20T16:06:48-05:00 New Revision: a1e2c6566305061c115954b048f2957c8d55cb5b
URL: https://github.com/llvm/llvm-project/commit/a1e2c6566305061c115954b048f2957c8d55cb5b DIFF: https://github.com/llvm/llvm-project/commit/a1e2c6566305061c115954b048f2957c8d55cb5b.diff LOG: [CUDA][HIP] ignore implicit host/device attr for override (#72815) 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. Added: Modified: clang/lib/Sema/SemaOverload.cpp clang/test/SemaCUDA/implicit-member-target-inherited.cu clang/test/SemaCUDA/trivial-ctor-dtor.cu Removed: ################################################################################ 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
