[clang] [CUDA][HIP] make trivial ctor/dtor host device (PR #72394)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 15 06:29:26 PST 2023
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Yaxun (Sam) Liu (yxsamliu)
<details>
<summary>Changes</summary>
Make trivial ctor/dtor implicitly host device functions so that they can be used to initialize file-scope
device variables to match nvcc behavior.
Fixes: https://github.com/llvm/llvm-project/issues/72261
Fixes: SWDEV-432412
---
Full diff: https://github.com/llvm/llvm-project/pull/72394.diff
10 Files Affected:
- (modified) clang/include/clang/Sema/Sema.h (+4)
- (modified) clang/lib/Sema/SemaCUDA.cpp (+20)
- (modified) clang/lib/Sema/SemaDecl.cpp (+3)
- (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 (+2-1)
- (modified) clang/test/SemaCUDA/implicit-member-target-collision.cu (+2-1)
- (modified) clang/test/SemaCUDA/implicit-member-target-inherited.cu (+6-4)
- (modified) clang/test/SemaCUDA/implicit-member-target.cu (+4-2)
- (added) clang/test/SemaCUDA/trivial-ctor-dtor.cu (+40)
``````````diff
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index a35a3c2c26c22ad..44dcbbf7605a557 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -13450,6 +13450,10 @@ 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 318174f7be8fa95..c376ab56dbef0e8 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -772,6 +772,26 @@ 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) {
+ auto IsTrivialCtor = [&](auto *D) {
+ if (auto *CD = dyn_cast<CXXConstructorDecl>(D))
+ return isEmptyCudaConstructor(SourceLocation(), CD);
+ return false;
+ };
+ auto IsTrivialDtor = [&](auto *D) {
+ if (auto *DD = dyn_cast<CXXDestructorDecl>(D))
+ return isEmptyCudaDestructor(SourceLocation(), DD);
+ return false;
+ };
+ if ((IsTrivialCtor(FD) || IsTrivialDtor(FD)) &&
+ !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 3876eb501083acb..a6cd0bb9ea2a829 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16232,6 +16232,9 @@ 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 acdd291b664579b..203f4fcbdf1efa0 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() {}
+ S() { x = 1; }
// 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 cbad7a1774c1501..31971fe6b3863c7 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; // expected-error{{no matching constructor for initialization of 'Out'}}
+ Out 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 06015ed0d6d8edc..0ee2e0963e40d59 100644
--- a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
@@ -6,7 +6,8 @@
// Test 1: collision between two bases
struct A1_with_host_ctor {
- A1_with_host_ctor() {}
+ int x;
+ A1_with_host_ctor() { x = 1; }
};
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 a50fddaa4615b22..060443c639924fb 100644
--- a/clang/test/SemaCUDA/implicit-member-target-collision.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-collision.cu
@@ -6,7 +6,8 @@
// Test 1: collision between two bases
struct A1_with_host_ctor {
- A1_with_host_ctor() {}
+ int x;
+ A1_with_host_ctor() { int x = 1; }
};
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 2178172ed01930d..8784135c0d6b66e 100644
--- a/clang/test/SemaCUDA/implicit-member-target-inherited.cu
+++ b/clang/test/SemaCUDA/implicit-member-target-inherited.cu
@@ -6,10 +6,11 @@
// Test 1: infer inherited default ctor to be host.
struct A1_with_host_ctor {
- A1_with_host_ctor() {}
+ A1_with_host_ctor() { x = 1; }
+ int x;
};
-// 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 constructor (the implicit copy constructor) not viable}}
+// expected-note at -5 {{candidate constructor (the implicit move constructor) not viable}}
// The inherited default constructor is inferred to be host, so we'll encounter
// an error when calling it from a __device__ function, but not from a __host__
@@ -83,7 +84,8 @@ void hostfoo3() {
// Test 4: infer inherited default ctor from a field, not a base
struct A4_with_host_ctor {
- A4_with_host_ctor() {}
+ int x;
+ A4_with_host_ctor() { int x = 1; }
};
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 d87e69624043419..2d260c64636ac84 100644
--- a/clang/test/SemaCUDA/implicit-member-target.cu
+++ b/clang/test/SemaCUDA/implicit-member-target.cu
@@ -6,7 +6,8 @@
// Test 1: infer default ctor to be host.
struct A1_with_host_ctor {
- A1_with_host_ctor() {}
+ int x;
+ A1_with_host_ctor() { x = 1; }
};
// The implicit default constructor is inferred to be host because it only needs
@@ -75,7 +76,8 @@ void hostfoo3() {
// Test 4: infer default ctor from a field, not a base
struct A4_with_host_ctor {
- A4_with_host_ctor() {}
+ int x;
+ A4_with_host_ctor() { int x = 1; }
};
struct B4_with_implicit_default_ctor {
diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
new file mode 100644
index 000000000000000..c7c0d33fe4c2d2e
--- /dev/null
+++ b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
@@ -0,0 +1,40 @@
+// 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() { x = 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}}
``````````
</details>
https://github.com/llvm/llvm-project/pull/72394
More information about the cfe-commits
mailing list