[clang] [CUDA][HIP] allow trivial ctor/dtor in device var init (PR #73140)

[Yaxun Liu via cfe-commits]

https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/73140

>From 2dc8bda89483ee655e7a76deac19b8ea9e463c7b 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] [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;