[clang] [CUDA][HIP] allow trivial ctor/dtor in device var init (PR #73140)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 22 07:55:01 PST 2023
https://github.com/yxsamliu created https://github.com/llvm/llvm-project/pull/73140
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
>From 511cecff7f76958ebfe713189bc106615763b64a Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 22 Nov 2023 09:53:17 -0500
Subject: [PATCH 1/3] Revert "[CUDA][HIP] ignore implicit host/device attr for
override (#72815)"
This reverts commit a1e2c6566305061c115954b048f2957c8d55cb5b.
---
clang/lib/Sema/SemaOverload.cpp | 6 ++----
.../SemaCUDA/implicit-member-target-inherited.cu | 1 -
clang/test/SemaCUDA/trivial-ctor-dtor.cu | 16 ----------------
3 files changed, 2 insertions(+), 21 deletions(-)
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/implicit-member-target-inherited.cu b/clang/test/SemaCUDA/implicit-member-target-inherited.cu
index ceca0891fc9b03c..781199bba6b5a11 100644
--- a/clang/test/SemaCUDA/implicit-member-target-inherited.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-inherited.cu
@@ -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;
diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
index 21d698d28492ac3..1df8adc62bab590 100644
--- a/clang/test/SemaCUDA/trivial-ctor-dtor.cu
+++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
@@ -38,19 +38,3 @@ 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 at -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;
>From e9a8e906d4c14eb4b317a7420b9bba3dc7321ba2 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 22 Nov 2023 09:53:20 -0500
Subject: [PATCH 2/3] Revert "[CUDA][HIP] make trivial ctor/dtor host device
(#72394)"
This reverts commit 27e6e4a4d0e3296cebad8db577ec0469a286795e.
---
clang/include/clang/Sema/Sema.h | 4 --
clang/lib/Sema/SemaCUDA.cpp | 16 --------
clang/lib/Sema/SemaDecl.cpp | 3 --
.../test/SemaCUDA/call-host-fn-from-device.cu | 2 +-
clang/test/SemaCUDA/default-ctor.cu | 2 +-
.../implicit-member-target-collision-cxx11.cu | 2 +-
.../implicit-member-target-collision.cu | 2 +-
.../implicit-member-target-inherited.cu | 4 +-
clang/test/SemaCUDA/implicit-member-target.cu | 4 +-
clang/test/SemaCUDA/trivial-ctor-dtor.cu | 40 -------------------
10 files changed, 8 insertions(+), 71 deletions(-)
delete mode 100644 clang/test/SemaCUDA/trivial-ctor-dtor.cu
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..318174f7be8fa95 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -772,22 +772,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/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 781199bba6b5a11..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}}
@@ -83,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
deleted file mode 100644
index 1df8adc62bab590..000000000000000
--- a/clang/test/SemaCUDA/trivial-ctor-dtor.cu
+++ /dev/null
@@ -1,40 +0,0 @@
-// RUN: %clang_cc1 -isystem %S/Inputs -fsyntax-only -verify %s
-// RUN: %clang_cc1 -isystem %S/Inputs -fcuda-is-device -fsyntax-only -verify %s
-
-#include <cuda.h>
-
-// Check trivial ctor/dtor
-struct A {
- int x;
- A() {}
- ~A() {}
-};
-
-__device__ A a;
-
-// Check trivial ctor/dtor of template class
-template<typename T>
-struct TA {
- T x;
- TA() {}
- ~TA() {}
-};
-
-__device__ TA<int> ta;
-
-// Check non-trivial ctor/dtor in parent template class
-template<typename T>
-struct TB {
- T x;
- TB() { static int nontrivial_ctor = 1; }
- ~TB() {}
-};
-
-template<typename T>
-struct TC : TB<T> {
- T x;
- TC() {}
- ~TC() {}
-};
-
-__device__ TC<int> tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
>From df2b64e19df24b2c2a3256af501b7190b32ebf36 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Wed, 22 Nov 2023 10:02:59 -0500
Subject: [PATCH 3/3] [CUDA][HIP] allow trivial ctor/dtor in device var init
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
---
clang/lib/Sema/SemaCUDA.cpp | 9 ++++
clang/test/SemaCUDA/trivial-ctor-dtor.cu | 57 ++++++++++++++++++++++++
2 files changed, 66 insertions(+)
create mode 100644 clang/test/SemaCUDA/trivial-ctor-dtor.cu
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 318174f7be8fa95..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);
diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
new file mode 100644
index 000000000000000..34142bcc621200f
--- /dev/null
+++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -isystem %S/Inputs -fsyntax-only -verify %s
+// RUN: %clang_cc1 -isystem %S/Inputs -fcuda-is-device -fsyntax-only -verify %s
+
+#include <cuda.h>
+
+// Check trivial ctor/dtor
+struct A {
+ int x;
+ A() {}
+ ~A() {}
+};
+
+__device__ A a;
+
+// Check trivial ctor/dtor of template class
+template<typename T>
+struct TA {
+ T x;
+ TA() {}
+ ~TA() {}
+};
+
+__device__ TA<int> ta;
+
+// Check non-trivial ctor/dtor in parent template class
+template<typename T>
+struct TB {
+ T x;
+ TB() { static int nontrivial_ctor = 1; }
+ ~TB() {}
+};
+
+template<typename T>
+struct TC : TB<T> {
+ T x;
+ TC() {}
+ ~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 {
+ 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; }
+__device__ C<float> cf_d; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+C<float> cf_h;
More information about the cfe-commits
mailing list