[SYCLomatic] Use the SYCL extension device_global to support more migration scenarios (#2757)

Signed-off-by: intwanghao <hao3.wang@intel.com>

Author: intwanghao

clang/lib/DPCT/AnalysisInfo.cpp

@@ -2947,7 +2947,7 @@ void MemVarInfo::migrateToDeviceGlobal(const VarDecl *MemVar) {
   auto &Ctx = DpctGlobalInfo::getContext();
   auto &MacroArgMap = DpctGlobalInfo::getMacroArgRecordMap();
   auto TSI = MemVar->getTypeSourceInfo();
-  auto OriginTL = TSI->getTypeLoc();
+  auto OriginTL = TSI->getTypeLoc().getUnqualifiedLoc().getAs<TypeLoc>();
   auto TL = OriginTL;
   auto BegLoc = MemVar->getBeginLoc();
   if (BegLoc.isMacroID()) {

clang/lib/DPCT/RulesLang/RulesLangNoneAPIAndType.cpp

@@ -809,7 +809,6 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) {
   auto MemVarRef = getNodeAsType<DeclRefExpr>(Result, "used");
   auto Func = getAssistNodeAsType<FunctionDecl>(Result, "func");
   auto Decl = getAssistNodeAsType<VarDecl>(Result, "decl");
-  DpctGlobalInfo &Global = DpctGlobalInfo::getInstance();
   if (MemVarRef && Func && Decl) {
     if (isCubVar(Decl)) {
       return;
@@ -820,8 +819,11 @@ void MemVarAnalysisRule::runRule(const MatchFinder::MatchResult &Result) {
       return;
     if (VD == nullptr)
       return;
-
-    auto Var = Global.findMemVarInfo(VD);
+    std::string CanonicalType = VD->getType().getCanonicalType().getAsString();
+    if (CanonicalType.find("block_tile_memory") != std::string::npos) {
+      return;
+    }
+    auto Var = MemVarInfo::buildMemVarInfo(VD);
     if (Func->hasAttr<CUDAGlobalAttr>() || Func->hasAttr<CUDADeviceAttr>()) {
       if (!(DpctGlobalInfo::useGroupLocalMemory() &&
             VD->hasAttr<CUDASharedAttr>() &&

clang/test/dpct/device_global/device_global2.cu

@@ -0,0 +1,20 @@
+// RUN: dpct --format-range=none --use-experimental-features=device_global -in-root %S -out-root %T/device_global2 %S/device_global2.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/device_global2/device_global2.dp.cpp --match-full-lines %s
+
+#include <cuda_runtime.h>
+#include <iostream>
+#include <vector>
+
+// CHECK: static sycl::ext::oneapi::experimental::device_global<int> var_a;
+__device__ int var_a;
+
+// CHECK: static constexpr sycl::ext::oneapi::experimental::device_global<int8_t[2]> var_b  {-1, -1};
+static constexpr __device__ int8_t var_b[2] = {-1, -1};
+
+template<typename T>
+__global__ void kernel(T b) {
+  var_a;
+  var_b[0];
+}
+
+template __global__ void kernel<int>(int b);