[clang] [CUDA][HIP] allow trivial ctor/dtor in device var init (PR #73140)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 22 07:55:32 PST 2023
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Yaxun (Sam) Liu (yxsamliu)
<details>
<summary>Changes</summary>
Treat ctor/dtor in device var init as host device function
so that they can be used to initialize file-scope
device variables to match nvcc behavior. If they are non-trivial
they will be diagnosed.
We cannot add implicit host device attrs to non-trivial
ctor/dtor since determining whether they are non-trivial
needs to know whether they have a trivial body and all their
member and base classes' ctor/dtor have trivial body, which
is affected by where their bodies are defined or instantiated.
Fixes: #<!-- -->72261
Fixes: SWDEV-432412
---
Full diff: https://github.com/llvm/llvm-project/pull/73140.diff
11 Files Affected:
- (modified) clang/include/clang/Sema/Sema.h (-4)
- (modified) clang/lib/Sema/SemaCUDA.cpp (+9-16)
- (modified) clang/lib/Sema/SemaDecl.cpp (-3)
- (modified) clang/lib/Sema/SemaOverload.cpp (+2-4)
- (modified) clang/test/SemaCUDA/call-host-fn-from-device.cu (+1-1)
- (modified) clang/test/SemaCUDA/default-ctor.cu (+1-1)
- (modified) clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu (+1-1)
- (modified) clang/test/SemaCUDA/implicit-member-target-collision.cu (+1-1)
- (modified) clang/test/SemaCUDA/implicit-member-target-inherited.cu (+2-3)
- (modified) clang/test/SemaCUDA/implicit-member-target.cu (+2-2)
- (modified) clang/test/SemaCUDA/trivial-ctor-dtor.cu (+5-4)
``````````diff
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 59806bcbcbb2dbc..e8914f5fcddf19e 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -13466,10 +13466,6 @@ class Sema final {
void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
const LookupResult &Previous);
- /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to a
- /// trivial cotr/dtor that does not have host and device attributes.
- void maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FunctionDecl *FD);
-
/// May add implicit CUDAConstantAttr attribute to VD, depending on VD
/// and current compilation settings.
void MaybeAddCUDAConstantAttr(VarDecl *VD);
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index b94f448dabe7517..6a66ecf6f94c178 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -225,6 +225,15 @@ Sema::CUDAFunctionPreference
Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
assert(Callee && "Callee must be valid.");
+
+ // Treat ctor/dtor as host device function in device var initializer to allow
+ // trivial ctor/dtor without device attr to be used. Non-trivial ctor/dtor
+ // will be diagnosed by checkAllowedCUDAInitializer.
+ if (Caller == nullptr && CurCUDATargetCtx.Kind == CTCK_InitGlobalVar &&
+ CurCUDATargetCtx.Target == CFT_Device &&
+ (isa<CXXConstructorDecl>(Callee) || isa<CXXDestructorDecl>(Callee)))
+ return CFP_HostDevice;
+
CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
@@ -772,22 +781,6 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
}
-// If a trivial ctor/dtor has no host/device
-// attributes, make it implicitly host device function.
-void Sema::maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FunctionDecl *FD) {
- bool IsTrivialCtor = false;
- if (auto *CD = dyn_cast<CXXConstructorDecl>(FD))
- IsTrivialCtor = isEmptyCudaConstructor(SourceLocation(), CD);
- bool IsTrivialDtor = false;
- if (auto *DD = dyn_cast<CXXDestructorDecl>(FD))
- IsTrivialDtor = isEmptyCudaDestructor(SourceLocation(), DD);
- if ((IsTrivialCtor || IsTrivialDtor) && !FD->hasAttr<CUDAHostAttr>() &&
- !FD->hasAttr<CUDADeviceAttr>()) {
- FD->addAttr(CUDAHostAttr::CreateImplicit(Context));
- FD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
- }
-}
-
// TODO: `__constant__` memory may be a limited resource for certain targets.
// A safeguard may be needed at the end of compilation pipeline if
// `__constant__` memory usage goes beyond limit.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 4e1857b931cc868..23dd8ae15c16583 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16255,9 +16255,6 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
if (FD && !FD->isDeleted())
checkTypeSupport(FD->getType(), FD->getLocation(), FD);
- if (LangOpts.CUDA)
- maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FD);
-
return dcl;
}
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 64607e28b8b35e6..9800d7f1c9cfee9 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -1491,10 +1491,8 @@ 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, isa<CXXConstructorDecl>(New)),
- OldTarget = SemaRef.IdentifyCUDATarget(
- Old, isa<CXXConstructorDecl>(New));
+ Sema::CUDAFunctionTarget NewTarget = SemaRef.IdentifyCUDATarget(New),
+ OldTarget = SemaRef.IdentifyCUDATarget(Old);
if (NewTarget != Sema::CFT_InvalidTarget) {
assert((OldTarget != Sema::CFT_InvalidTarget) &&
"Unexpected invalid target.");
diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu
index b62de92db02d6de..acdd291b664579b 100644
--- a/clang/test/SemaCUDA/call-host-fn-from-device.cu
+++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu
@@ -12,7 +12,7 @@ extern "C" void host_fn() {}
struct Dummy {};
struct S {
- S() { static int nontrivial_ctor = 1; }
+ S() {}
// expected-note at -1 2 {{'S' declared here}}
~S() { host_fn(); }
// expected-note at -1 {{'~S' declared here}}
diff --git a/clang/test/SemaCUDA/default-ctor.cu b/clang/test/SemaCUDA/default-ctor.cu
index 31971fe6b3863c7..cbad7a1774c1501 100644
--- a/clang/test/SemaCUDA/default-ctor.cu
+++ b/clang/test/SemaCUDA/default-ctor.cu
@@ -25,7 +25,7 @@ __device__ void fd() {
InD ind;
InH inh; // expected-error{{no matching constructor for initialization of 'InH'}}
InHD inhd;
- Out out;
+ Out out; // expected-error{{no matching constructor for initialization of 'Out'}}
OutD outd;
OutH outh; // expected-error{{no matching constructor for initialization of 'OutH'}}
OutHD outhd;
diff --git a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu b/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
index edb543f637ccc18..06015ed0d6d8edc 100644
--- a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
@@ -6,7 +6,7 @@
// Test 1: collision between two bases
struct A1_with_host_ctor {
- A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+ A1_with_host_ctor() {}
};
struct B1_with_device_ctor {
diff --git a/clang/test/SemaCUDA/implicit-member-target-collision.cu b/clang/test/SemaCUDA/implicit-member-target-collision.cu
index 16b5978af40872b..a50fddaa4615b22 100644
--- a/clang/test/SemaCUDA/implicit-member-target-collision.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-collision.cu
@@ -6,7 +6,7 @@
// Test 1: collision between two bases
struct A1_with_host_ctor {
- A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+ A1_with_host_ctor() {}
};
struct B1_with_device_ctor {
diff --git a/clang/test/SemaCUDA/implicit-member-target-inherited.cu b/clang/test/SemaCUDA/implicit-member-target-inherited.cu
index ceca0891fc9b03c..2178172ed01930d 100644
--- a/clang/test/SemaCUDA/implicit-member-target-inherited.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-inherited.cu
@@ -6,7 +6,7 @@
// Test 1: infer inherited default ctor to be host.
struct A1_with_host_ctor {
- A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+ A1_with_host_ctor() {}
};
// expected-note at -3 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note at -4 {{candidate constructor (the implicit move constructor) not viable}}
@@ -39,7 +39,6 @@ struct A2_with_device_ctor {
};
// expected-note at -3 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note at -4 {{candidate constructor (the implicit move constructor) not viable}}
-// expected-note at -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;
@@ -84,7 +83,7 @@ void hostfoo3() {
// Test 4: infer inherited default ctor from a field, not a base
struct A4_with_host_ctor {
- A4_with_host_ctor() { static int nontrivial_ctor = 1; }
+ A4_with_host_ctor() {}
};
struct B4_with_inherited_host_ctor : A4_with_host_ctor{
diff --git a/clang/test/SemaCUDA/implicit-member-target.cu b/clang/test/SemaCUDA/implicit-member-target.cu
index 552f8f2ebd94fd5..d87e69624043419 100644
--- a/clang/test/SemaCUDA/implicit-member-target.cu
+++ b/clang/test/SemaCUDA/implicit-member-target.cu
@@ -6,7 +6,7 @@
// Test 1: infer default ctor to be host.
struct A1_with_host_ctor {
- A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+ A1_with_host_ctor() {}
};
// The implicit default constructor is inferred to be host because it only needs
@@ -75,7 +75,7 @@ void hostfoo3() {
// Test 4: infer default ctor from a field, not a base
struct A4_with_host_ctor {
- A4_with_host_ctor() { static int nontrivial_ctor = 1; }
+ A4_with_host_ctor() {}
};
struct B4_with_implicit_default_ctor {
diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
index 21d698d28492ac3..34142bcc621200f 100644
--- a/clang/test/SemaCUDA/trivial-ctor-dtor.cu
+++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
@@ -37,12 +37,13 @@ struct TC : TB<T> {
~TC() {}
};
+template class TC<int>;
+
__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 at -1 {{candidate constructor (the implicit move constructor) not viable}}
+struct C {
explicit C() {};
};
@@ -51,6 +52,6 @@ __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>'}}
+template <> C<float>::C() { static int nontrivial_ctor = 1; }
+__device__ C<float> cf_d; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
C<float> cf_h;
``````````
</details>
https://github.com/llvm/llvm-project/pull/73140
More information about the cfe-commits
mailing list