[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

@@ -25,6 +25,7 @@
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/TypeLoc.h"
 #include "llvm/Support/raw_ostream.h"
+#include <optional>
 
 extern clang::tooling::UnifiedPath DpctInstallPath;
 namespace clang {
@@ -553,7 +554,22 @@ 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 /*FIXME: && IsDependent*/) {
+    if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) {
+      ExprAnalysis EA(VD->getInit());
+      auto TDSI = EA.getTemplateDependentStringInfo();
+      std::string VDStr = VD->getNameAsString();
+      std::string VDInitStr = EA.getReplacedString();
+      auto Loc = ReplSet.getSourceStr().find(VDStr);
+      // TODO: more than 1 substrings matched
+      if (Loc != std::string::npos) {
+        addReplacement(Loc, VDStr.size(), VDInitStr);
+        addReplacement(Loc, "ThreadsPerBlock" /*FIXME*/, TDSI);
+      }
+    }
+  } 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

@@ -40,6 +40,7 @@ class StringReplacement {
   }
 
   inline const std::string &getReplacedText() { return Text; }
+  inline size_t getLength() { return Length; }
 
 private:
   // SourceStr is the string which need replaced.
@@ -77,6 +78,7 @@ class TemplateDependentReplacement {
   }
   inline size_t getOffset() const { return Offset; }
   inline size_t getLength() const { return Length; }
+  inline size_t getTemplateIndex() const { return TemplateIndex; }
   const TemplateArgumentInfo &
   getTargetArgument(const std::vector<TemplateArgumentInfo> &TemplateList);
   void replace(const std::vector<TemplateArgumentInfo> &TemplateList);
@@ -117,6 +119,10 @@ class TemplateDependentStringInfo {
     HelperFeatureSet = Set;
   }
   bool containsTemplateDependentMacro() const { return ContainsTemplateDependentMacro; }
+  const std::vector<std::shared_ptr<TemplateDependentReplacement>> &
+  getTDRs() const {
+    return TDRs;
+  }
 };
 
 /// Store an expr source string which may need replaced and its replacements
@@ -132,10 +138,36 @@ class StringReplacements {
   // Add a template dependent replacement
   inline void addTemplateDependentReplacement(size_t Offset, size_t Length,
                                               unsigned TemplateIndex) {
-    TDRs.insert(
-        std::make_pair(Offset, std::make_shared<TemplateDependentReplacement>(
-                                   SourceStr, Offset, Length, TemplateIndex)));
+    // Find items in ReplMap whose offset <= Offset.
+    // Then check length, is there is overlap, ignore this insertion.
+    // Finally calculate the shift length.
+    int Shift = 0;
+    auto UpperBound = ReplMap.upper_bound(Offset);
+    for (auto It = ReplMap.begin(); It != UpperBound; ++It) {
+      if ((It->first + It->second->getLength()) > Offset) {
+        // overlap
+        return;
+      }
+      Shift +=
+          (It->second->getReplacedText().length() - It->second->getLength());
+    }
+    auto TDR = std::make_shared<TemplateDependentReplacement>(
+        SourceStr, Offset, Length, TemplateIndex);
+    TDR->shift(Shift);
+    TDRs.insert(std::make_pair(Offset + Shift, TDR));
+  }
+
+  inline void addTemplateDependentReplacement(
+      size_t Offset, std::string String,
+      std::shared_ptr<TemplateDependentStringInfo> TDSI) {
+    for (const auto &Item : TDSI->getTDRs()) {
+      Offset += Item->getOffset();
+      auto TDR = std::make_shared<TemplateDependentReplacement>(
+          String, Offset, String.size(), Item->getTemplateIndex());
+      TDRs.insert(std::make_pair(Offset, TDR));
+    }
   }
+
   // Add a string replacement
   void addStringReplacement(size_t Offset, size_t Length, std::string Text) {
     auto Result = ReplMap.insert(std::make_pair(
@@ -163,6 +195,7 @@ class StringReplacements {
     replaceString();
     return SourceStr;
   }
+  inline const std::string &getSourceStr() { return SourceStr; }
 
 private:
   StringReplacements(const StringReplacements &) = delete;
@@ -586,6 +619,12 @@ class ExprAnalysis {
     ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex);
   }
 
+  inline void
+  addReplacement(size_t Offset, std::string String,
+                 std::shared_ptr<TemplateDependentStringInfo> TDSI) {
+    ReplSet.addTemplateDependentReplacement(Offset, String, TDSI);
+  }
+
   // Analyze the expression, jump to corresponding analysis function according
   // to its class
   // Precondition: Expression != nullptr
@@ -691,6 +730,12 @@ class ExprAnalysis {
   std::string RewritePrefix;
   std::string RewritePostfix;
   std::set<HelperFeatureEnum> HelperFeatureSet;
+  std::optional<std::pair<std::string /*constexpr definitaion*/,
+                          unsigned /*constexpr offset in original str*/>>
+      ConstExprExpansionInfo = std::nullopt;
+
+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 = (2 * ThreadsPerBlock / kWarpSize);
+  __shared__ int smem[kNumWarps * NumWarpQ];
+}
+
+void foo2() {
+  // CHECK: sycl::local_accessor<int, 1> smem_acc_ct1(sycl::range<1>((2 * 128 / kWarpSize) * 8), cgh);
+  kerfunc<128, 8><<<32, 32>>>();
+}