[CUDA][HIP] allow trivial ctor/dtor in device var init #73140
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Yaxun (Sam) Liu (yxsamliu) ChangesTreat ctor/dtor in device var init as host device function We cannot add implicit host device attrs to non-trivial Fixes: #72261 Fixes: SWDEV-432412 Full diff: https://github.com/llvm/llvm-project/pull/73140.diff 11 Files Affected:
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@-1 2 {{'S' declared here}}
~S() { host_fn(); }
// expected-note@-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@-3 {{candidate constructor (the implicit copy constructor) not viable}}
// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
@@ -39,7 +39,6 @@ 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;
@@ -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@-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;
|
|
Could you first land the two reverts (511cecff7f76958ebfe713189bc106615763b64a and e9a8e906d4c14eb4b317a7420b9bba3dc7321ba2) and then have the third commit properly reviewed? @Artem-B may be unavailable for a few more days, but we'd like a fix/revert to land very soon. |
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: llvm#72261 Fixes: SWDEV-432412
yxsamliu
force-pushed
the
fix-trivial2
branch
from
df2b64e to
2dc8bda
Compare
|
Thank you! |
|
ping |
Artem-B
approved these changes
Nov 30, 2023
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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