[clang] [CUDA/HIP] Fix errors for device function used in host global initializers (PR #222338)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Mariya Podchishchaeva (Fznamznon)

<details>
<summary>Changes</summary>

The existing checker was simply looking for a call or `CXXConstructExpr` which would miss a call to device function in case an implicit cast was in place or if a call was a part of an expression. Use a visitor to improve the situation.

---
Full diff: https://github.com/llvm/llvm-project/pull/222338.diff


2 Files Affected:

- (modified) clang/lib/Sema/SemaCUDA.cpp (+33-21) 
- (modified) clang/test/SemaCUDA/global-initializers.cu (+45-2) 


``````````diff
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 13aa06aa97399..13d98b344cd97 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -776,39 +776,51 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
   } else {
     // This is a host-side global variable.  Check that the initializer is
     // callable from the host side.
-    const FunctionDecl *InitFn = nullptr;
-    if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
-      InitFn = CE->getConstructor();
-    } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
-      InitFn = CE->getDirectCallee();
-    }
-    if (InitFn) {
-      CUDAFunctionTarget InitFnTarget = IdentifyTarget(InitFn);
-      if (InitFnTarget != CUDAFunctionTarget::Host &&
-          InitFnTarget != CUDAFunctionTarget::HostDevice) {
-        Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
-            << InitFnTarget << InitFn;
-        Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
-        VD->setInvalidDecl();
-      }
-    }
+
     struct GlobVarInitChecker : ConstEvaluatedExprVisitor<GlobVarInitChecker> {
+    private:
       using Base = ConstEvaluatedExprVisitor<GlobVarInitChecker>;
       SemaCUDA &SCRef;
-      SourceLocation InitLoc;
+      VarDecl *VD;
+      void CheckForWrongSidedCall(const FunctionDecl *FD) {
+        CUDAFunctionTarget InitFnTarget = SCRef.IdentifyTarget(FD);
+        if (InitFnTarget != CUDAFunctionTarget::Host &&
+            InitFnTarget != CUDAFunctionTarget::HostDevice) {
+          SCRef.Diag(VD->getLocation(),
+                     diag::err_ref_bad_target_global_initializer)
+              << InitFnTarget << FD;
+          SCRef.Diag(FD->getLocation(), diag::note_previous_decl) << FD;
+          VD->setInvalidDecl();
+        }
+      }
 
-      GlobVarInitChecker(SemaCUDA &S, SourceLocation L)
-          : Base(S.getASTContext()), SCRef(S), InitLoc(L) {}
+    public:
+      GlobVarInitChecker(SemaCUDA &S, VarDecl *VD)
+          : Base(S.getASTContext()), SCRef(S), VD(VD) {}
       void VisitDeclRefExpr(const DeclRefExpr *DRE) {
         if (auto *VarD = dyn_cast<VarDecl>(DRE->getDecl());
             VarD && VarD->hasAttr<HIPManagedAttr>()) {
           SCRef.Diag(DRE->getLocation(),
                      diag::err_cuda_invalid_use_of_managedvar);
-          SCRef.Diag(InitLoc, diag::note_cuda_managed_var_in_glob_init);
+          SCRef.Diag(VD->getLocation(),
+                     diag::note_cuda_managed_var_in_glob_init);
         }
+        Base::VisitDeclRefExpr(DRE);
+      }
+      void VisitCallExpr(const CallExpr *CE) {
+        const FunctionDecl *InitFn = CE->getDirectCallee();
+        if (InitFn)
+          CheckForWrongSidedCall(InitFn);
+        Base::VisitCallExpr(CE);
+      }
+      void VisitCXXConstructExpr(const CXXConstructExpr *CE) {
+        const FunctionDecl *InitFn = CE->getConstructor();
+        if (InitFn)
+          CheckForWrongSidedCall(InitFn);
+        Base::VisitCXXConstructExpr(CE);
       }
     };
-    GlobVarInitChecker Checker(*this, VD->getLocation());
+    GlobVarInitChecker Checker(*this, VD);
     Checker.Visit(Init);
   }
 }
diff --git a/clang/test/SemaCUDA/global-initializers.cu b/clang/test/SemaCUDA/global-initializers.cu
index 29e386134a3dd..3cd46b90638c7 100644
--- a/clang/test/SemaCUDA/global-initializers.cu
+++ b/clang/test/SemaCUDA/global-initializers.cu
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify
-// RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify
+// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify=expected,host
+// RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify=expected,device
 
 #include "Inputs/cuda.h"
 
@@ -70,3 +70,46 @@ __device__ double AY = a.pow(2.0, 2); // expected-error{{dynamic initialization
 const A ca;
 const double CAX = ca.cpow(1.0, 1);
 const __device__ double CAY = ca.cpow(2.0, 2);
+
+namespace ns1 {
+  // host-note at +3 {{'value_func' declared here}}
+  // expected-note at +2 5{{'value_func' declared here}}
+  // expected-note at +1 {{candidate function not viable: call to __device__ function from __host__ function}}
+__device__ constexpr inline int value_func() {
+  return 32;
+}
+}
+
+namespace ns2 {
+  using namespace ns1;
+  // diagnosed via overloading.
+  constexpr static unsigned var0 = value_func();
+  // expected-error at -1 {{no matching function for call to 'value_func'}}
+
+  // diagnosed via SemaCUDA::checkAllowedInitializer
+  constexpr static unsigned var1 = ns1::value_func();
+  // host-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+  // FIXME: Inconsistency with var1 - non constexpr cases are diagnosed for both host and device.
+  static int var2 = 1 + ns1::value_func();
+  // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+  static int var3 {ns1::value_func()};
+  // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+  struct A {
+    unsigned b;
+  };
+  A b{ns1::value_func()};
+  // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+  int foo(int);
+  int nested = foo(ns1::value_func());
+  // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+  void foobar() {
+    // diagnosed via SemaCUDA::checkCall
+    static int var2 = 1 + ns1::value_func();
+  // host-error at -1 {{reference to __device__ function 'value_func' in __host__ function}}
+  // device-error at -2 {{reference to __device__ function 'value_func' in global initializer}}
+  }
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/222338