[clang] 6b3470b - Revert "[CUDA][HIP] make trivial ctor/dtor host device (#72394)"
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 22 18:21:19 PST 2023
Author: Yaxun (Sam) Liu
Date: 2023-11-22T21:20:53-05:00
New Revision: 6b3470b4b83195aeeda60b101e8d3bf8800c321c
URL: https://github.com/llvm/llvm-project/commit/6b3470b4b83195aeeda60b101e8d3bf8800c321c
DIFF: https://github.com/llvm/llvm-project/commit/6b3470b4b83195aeeda60b101e8d3bf8800c321c.diff
LOG: Revert "[CUDA][HIP] make trivial ctor/dtor host device (#72394)"
This reverts commit 27e6e4a4d0e3296cebad8db577ec0469a286795e.
This patch is reverted due to regression. A testcase is:
`template <class T>
struct ptr {
~ptr() { static int x = 1;}
};
template <class T>
struct Abc : ptr<T> {
public:
Abc();
~Abc() {}
};
template
class Abc<int>;
`
Added:
Modified:
clang/include/clang/Sema/Sema.h
clang/lib/Sema/SemaCUDA.cpp
clang/lib/Sema/SemaDecl.cpp
clang/test/SemaCUDA/call-host-fn-from-device.cu
clang/test/SemaCUDA/default-ctor.cu
clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
clang/test/SemaCUDA/implicit-member-target-collision.cu
clang/test/SemaCUDA/implicit-member-target-inherited.cu
clang/test/SemaCUDA/implicit-member-target.cu
Removed:
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}}
More information about the cfe-commits
mailing list