[SYCLomatic] Enhance the migration of __shared__ memory

clang/lib/DPCT/AnalysisInfo.cpp

@@ -2811,6 +2811,7 @@ void CtTypeInfo::setArrayInfo(const DependentSizedArrayTypeLoc &TL,
                               bool NeedSizeFold) {
   ContainSizeofType = containSizeOfType(TL.getSizeExpr());
   ExprAnalysis EA;
+  EA.IsAnalyzingCtTypeInfo = true;
   EA.analyze(TL.getSizeExpr());
   auto TDSI = EA.getTemplateDependentStringInfo();
   if (TDSI->containsTemplateDependentMacro())

clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp

@@ -553,7 +553,16 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) {
   }
   if (auto TemplateDecl = dyn_cast<NonTypeTemplateParmDecl>(DRE->getDecl()))
     addReplacement(DRE, TemplateDecl->getIndex());
-  else if (auto ECD = dyn_cast<EnumConstantDecl>(DRE->getDecl())) {
+  else if (const auto *VD = dyn_cast<VarDecl>(DRE->getDecl());
+           VD && VD->isConstexpr() &&
+           IsAnalyzingCtTypeInfo /*&& IsDependent*/) {
+    if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) {
+      ExprAnalysis EA(VD->getInit());
+      std::string VDInitStr = EA.getReplacedString();
+      std::string VDStr = VD->getNameAsString();
+      addReplacement(0, VDStr.size(), VDInitStr);
+    }
+  } else if (auto ECD = dyn_cast<EnumConstantDecl>(DRE->getDecl())) {
     std::unordered_set<std::string> targetStr = {
         "thread_scope_system",  "thread_scope_device",  "thread_scope_block",
         "memory_order_relaxed", "memory_order_acq_rel", "memory_order_release",

clang/lib/DPCT/RuleInfra/ExprAnalysis.h

@@ -77,6 +77,7 @@ class TemplateDependentReplacement {
   }
   inline size_t getOffset() const { return Offset; }
   inline size_t getLength() const { return Length; }
+  inline unsigned getTemplateIndex() const { return TemplateIndex; }
   const TemplateArgumentInfo &
   getTargetArgument(const std::vector<TemplateArgumentInfo> &TemplateList);
   void replace(const std::vector<TemplateArgumentInfo> &TemplateList);
@@ -136,6 +137,15 @@ class StringReplacements {
         std::make_pair(Offset, std::make_shared<TemplateDependentReplacement>(
                                    SourceStr, Offset, Length, TemplateIndex)));
   }
+
+  inline void addTemplateDependentReplacement(
+      size_t Offset, size_t Length,
+      std::shared_ptr<TemplateDependentStringInfo> TDSI) {
+   // TDRs.insert(
+   //     std::make_pair(Offset, std::make_shared<TemplateDependentReplacement>(
+   //                                SourceStr, Offset, Length, TemplateIndex)));
+  }
+
   // Add a string replacement
   void addStringReplacement(size_t Offset, size_t Length, std::string Text) {
     auto Result = ReplMap.insert(std::make_pair(
@@ -163,6 +173,10 @@ class StringReplacements {
     replaceString();
     return SourceStr;
   }
+  inline std::map<size_t, std::shared_ptr<TemplateDependentReplacement>>
+  getTDRs() {
+    return TDRs;
+  }
 
 private:
   StringReplacements(const StringReplacements &) = delete;
@@ -230,6 +244,10 @@ class ExprAnalysis {
   inline const std::string &getReplacedString() {
     return ReplSet.getReplacedString();
   }
+  inline std::map<size_t, std::shared_ptr<TemplateDependentReplacement>>
+  getReplSetTDRs() {
+    return ReplSet.getTDRs();
+  }
   inline std::shared_ptr<TemplateDependentStringInfo>
   getTemplateDependentStringInfo() {
     auto Res = ReplSet.getTemplateDependentStringInfo();
@@ -586,6 +604,12 @@ class ExprAnalysis {
     ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex);
   }
 
+  inline void
+  addReplacement(size_t Offset, size_t Length,
+                 std::shared_ptr<TemplateDependentStringInfo> TDSI) {
+    ReplSet.addTemplateDependentReplacement(Offset, Length, TDSI);
+  }
+
   // Analyze the expression, jump to corresponding analysis function according
   // to its class
   // Precondition: Expression != nullptr
@@ -691,6 +715,9 @@ class ExprAnalysis {
   std::string RewritePrefix;
   std::string RewritePostfix;
   std::set<HelperFeatureEnum> HelperFeatureSet;
+
+public:
+  bool IsAnalyzingCtTypeInfo = false;
 };
 
 // Analyze pointer allocated by cudaMallocManaged.

clang/test/dpct/sharedmem_var_static.cu

@@ -4,11 +4,12 @@
 // RUN: dpct --format-range=none --usm-level=none -out-root %T/sharedmem_var_static %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only
 // RUN: FileCheck %s --match-full-lines --input-file %T/sharedmem_var_static/sharedmem_var_static.dp.cpp
 // RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST  %T/sharedmem_var_static/sharedmem_var_static.dp.cpp -o %T/sharedmem_var_static/sharedmem_var_static.dp.o %}
-#ifndef NO_BUILD_TEST
+
 #include <stdio.h>
 #include <complex>
 #define SIZE 64
 
+#ifndef NO_BUILD_TEST
 class TestObject{
 public:
   // CHECK: static void run(int *in, int *out, int &a0) {
@@ -224,3 +225,15 @@ void fooh() {
   fook<SZ><<<1, 1>>>();
 }
 #endif
+
+constexpr int kWarpSize = 32;
+
+template <int ThreadsPerBlock, int NumWarpQ> __global__ void kerfunc() {
+  constexpr int kNumWarps = ThreadsPerBlock / kWarpSize;
+  __shared__ int smem[kNumWarps * NumWarpQ];
+}
+
+void foo2() {
+  // CHECK: sycl::local_accessor<int, 1> smem_acc_ct1(sycl::range<1>(128 / kWarpSize * 8), cgh);
+  kerfunc<128, 8><<<32, 32>>>();
+}