[SYCLomatic] Fix the issue that auto type is used in template scenario after migration for auto type global device memory (#2763)

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

Author: intwanghao

clang/lib/DPCT/AnalysisInfo.cpp

@@ -2625,6 +2625,9 @@ CtTypeInfo::CtTypeInfo(const VarDecl *D, bool NeedSizeFold) : CtTypeInfo() {
     auto TL = D->getTypeSourceInfo()->getTypeLoc();
     IsConstantQualified = D->hasAttr<CUDAConstantAttr>();
     setTypeInfo(TL, NeedSizeFold);
+    if (IsAutoSpecified) {
+      DeducedTypeStr = DpctGlobalInfo::getReplacedTypeName(D->getType());
+    }
     if (TL.getTypeLocClass() == TypeLoc::IncompleteArray) {
       if (auto CAT = dyn_cast<ConstantArrayType>(D->getType())) {
         Range[0] = std::to_string(CAT->getSize().getZExtValue());
@@ -2822,6 +2825,12 @@ void CtTypeInfo::setArrayInfo(const IncompleteArrayTypeLoc &TL,
 }
 void CtTypeInfo::setName(const TypeLoc &TL) {
   ExprAnalysis EA;
+  QualType QT = TL.getType();
+  if (const Type *TP = QT.getTypePtr()) {
+    if (isa<AutoType>(TP)) {
+      IsAutoSpecified = true;
+    }
+  }
   EA.analyze(TL);
   TDSI = EA.getTemplateDependentStringInfo();
   auto SetFromTL = EA.getHelperFeatureSet();
@@ -3019,7 +3028,11 @@ void MemVarInfo::migrateToDeviceGlobal(const VarDecl *MemVar) {
     }
   }
   if (BaseTypeStr.empty()) {
-    BaseTypeStr = getType()->getBaseNameWithoutQualifiers();
+    if (getType()->isAutoSpecified()) {
+      BaseTypeStr = getType()->getDeducedTypeStr();
+    } else {
+      BaseTypeStr = getType()->getBaseNameWithoutQualifiers();
+    }
     TypeReplLoc = TL.getBeginLoc();
     TypeReplLen =
         SM.getFileOffset(TL.getEndLoc()) - SM.getFileOffset(TL.getBeginLoc()) +

clang/lib/DPCT/AnalysisInfo.h

@@ -1799,6 +1799,8 @@ class CtTypeInfo {
   }
   bool containsTemplateDependentMacro() const { return TemplateDependentMacro; }
   bool isConstantQualified() const { return IsConstantQualified; }
+  bool isAutoSpecified() { return IsAutoSpecified; }
+  std::string getDeducedTypeStr() { return DeducedTypeStr; }
 
 private:
   // For ConstantArrayType, size in generated code is folded as an integer.
@@ -1858,6 +1860,8 @@ class CtTypeInfo {
   std::vector<std::string> ArraySizeOriginExprs{};
   std::set<HelperFeatureEnum> HelperFeatureSet;
   std::shared_ptr<TemplateDependentStringInfo> TDSI;
+  bool IsAutoSpecified = false;
+  std::string DeducedTypeStr;
 };
 
 // variable info includes name, type and location.

clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp

@@ -611,6 +611,11 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) {
       }
     }
   } else if (auto VD = dyn_cast<VarDecl>(DRE->getDecl())) {
+    DpctGlobalInfo &Global = DpctGlobalInfo::getInstance();
+    auto Info = Global.findMemVarInfo(VD);
+    if (Info && Info->isUseDeviceGlobal()) {
+      addReplacement(DRE, getStmtSpelling(DRE) + ".get()");
+    }
     if (RefString == "warpSize" &&
         !DpctGlobalInfo::isInAnalysisScope(VD->getLocation())) {
       addReplacement(DRE, DpctGlobalInfo::getSubGroup(DRE) +

clang/test/dpct/device_global/device_global.cu

@@ -14,6 +14,7 @@ struct B{
   int data = 1;
 };
 
+const float init_val = 0;
 
 // CHECK: static sycl::ext::oneapi::experimental::device_global<int> var_a;
 // CHECK: /*
@@ -31,13 +32,16 @@ struct B{
 // CHECK: static sycl::ext::oneapi::experimental::device_global<A> var_e;
 // CHECK: static sycl::ext::oneapi::experimental::device_global<B> var_f;
 // CHECK: inline dpct::global_memory<int, 0> var_g;
+// CHECK: static sycl::ext::oneapi::experimental::device_global<float> var_h{init_val};
 __device__ int var_a;
 __device__ int var_b = 0;
 __constant__ float var_c = 2.f;
 __constant__ float var_d = 1.f;
 __device__ A var_e;
 __constant__ B var_f;
 __device__ int var_g;
+__device__ auto var_h = init_val;
+
 // CHECK: static sycl::ext::oneapi::experimental::device_global<float[10]> arr_a;
 // CHECK: /*
 // CHECK: DPCT1127:{{[0-9]+}}: The constant compile-time initialization for device_global is supported when compiling with C++20. You may need to adjust the compile commands.
@@ -82,6 +86,7 @@ TABLE_END()
 // CHECK:   var_a.get() = 0;
 // CHECK:   var_e.get().data = 1;
 // CHECK:   *ptr = var_a.get() + var_b.get() + var_c.get() + var_d.get() + var_e.get().data + var_f.get().data + device_func();
+// CHECK:   sycl::fmax(var_h.get(), (double)0);
 // CHECK: }
 __device__ int device_func() {
   float *p = arr_a;
@@ -94,6 +99,7 @@ __global__ void kernel(int *ptr) {
   var_a = 0;
   var_e.data = 1;
   *ptr = var_a + var_b + var_c + var_d + var_e.data + var_f.data + device_func();
+  fmax(var_h, 0);
 }
 
 // CHECK: void kernel2(int *ptr,