[clang] 2b76e20 - [CUDA][HIP] allow trivial ctor/dtor in device var init (#73140)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Dec 1 13:24:10 PST 2023
Author: Yaxun (Sam) Liu
Date: 2023-12-01T16:24:01-05:00
New Revision: 2b76e20ea782790a78ec58d5f94ce88a173bab7f
URL: https://github.com/llvm/llvm-project/commit/2b76e20ea782790a78ec58d5f94ce88a173bab7f
DIFF: https://github.com/llvm/llvm-project/commit/2b76e20ea782790a78ec58d5f94ce88a173bab7f.diff
LOG: [CUDA][HIP] allow trivial ctor/dtor in device var init (#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
Added:
clang/test/SemaCUDA/trivial-ctor-dtor.cu
Modified:
clang/lib/Sema/SemaCUDA.cpp
Removed:
################################################################################
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