Merge branch 'oneapi-src:SYCLomatic' into migrate_utils_groups

clang/include/clang/DPCT/DPCTOptions.inc

@@ -342,11 +342,12 @@ DPCT_ENUM_OPTION(
     llvm::cl::value_desc("value"), llvm::cl::cat(DPCTReportGenCat),
     llvm::cl::Optional)
 
-DPCT_OPTION(clang::dpct::opt, std::string, SuppressWarnings,
+DPCT_OPTION(clang::dpct::list, std::string, SuppressWarnings,
             clang::dpct::DpctOptionClass::OC_Attribute,
             DPCT_OPTION_ACTIONS(clang::dpct::DpctActionKind::DAK_Migration),
             "suppress-warnings", llvm::cl::desc(SuppressWarningsMessage),
-            llvm::cl::value_desc("value"), llvm::cl::cat(DPCTWarningsCat))
+            llvm::cl::value_desc("value"), llvm::cl::CommaSeparated,
+            llvm::cl::cat(DPCTWarningsCat))
 
 DPCT_FLAG_OPTION(
     SuppressWarningsAll, clang::dpct::DpctOptionClass::OC_Attribute,

clang/include/clang/DPCT/DpctOptions.h

@@ -82,7 +82,7 @@ class DpctOptionBase {
                  std::initializer_list<DpctActionKind>);
 
 public:
-  ~DpctOptionBase() = default;
+  virtual ~DpctOptionBase() = default;
   static void init();
   static void check();
 };

clang/lib/DPCT/AnalysisInfo.cpp

@@ -3485,6 +3485,33 @@ void TextureObjectInfo::addParamDeclReplacement() {
                                          getParamDeclType(), nullptr));
   }
 }
+
+// class TempStorageVarInfo //
+void TempStorageVarInfo::addAccessorDecl(StmtList &AccessorList,
+                                         StringRef LocalSize) const {
+  std::string Accessor;
+  llvm::raw_string_ostream OS(Accessor);
+  OS << MapNames::getClNamespace() << "local_accessor<std::byte, 1> " << Name
+     << "_acc(";
+  DpctGlobalInfo::printCtadClass(OS, MapNames::getClNamespace() + "range", 1);
+  OS << '(' << LocalSize << ".size() * sizeof(" << Type->getSourceString()
+     << ")), cgh);";
+  AccessorList.emplace_back(Accessor);
+}
+void TempStorageVarInfo::applyTemplateArguments(
+    const std::vector<TemplateArgumentInfo> &TAList) {
+  Type = Type->applyTemplateArguments(TAList);
+}
+ParameterStream &TempStorageVarInfo::getFuncDecl(ParameterStream &PS) {
+  return PS << MapNames::getClNamespace() << "local_accessor<std::byte, 1> "
+            << Name;
+}
+ParameterStream &TempStorageVarInfo::getFuncArg(ParameterStream &PS) {
+  return PS << Name;
+}
+ParameterStream &TempStorageVarInfo::getKernelArg(ParameterStream &PS) {
+  return PS << Name << "_acc";
+}
 ///// class CudaLaunchTextureObjectInfo /////
 std::string
 CudaLaunchTextureObjectInfo::getAccessorDecl(const std::string &QueueString) {
@@ -3664,6 +3691,9 @@ void TemplateArgumentInfo::setArgFromExprAnalysis(const T &Arg,
 void MemVarMap::addTexture(std::shared_ptr<TextureInfo> Tex) {
   TextureMap.insert(std::make_pair(Tex->getOffset(), Tex));
 }
+void MemVarMap::addCUBTempStorage(std::shared_ptr<TempStorageVarInfo> Tmp) {
+  TempStorageMap.insert(std::make_pair(Tmp->getOffset(), Tmp));
+}
 void MemVarMap::addVar(std::shared_ptr<MemVarInfo> Var) {
   auto Attr = Var->getAttr();
   if (Var->isGlobal() && (Attr == MemVarInfo::VarAttrKind::Device ||
@@ -3688,6 +3718,7 @@ void MemVarMap::merge(const MemVarMap &VarMap,
   merge(LocalVarMap, VarMap.LocalVarMap, TemplateArgs);
   merge(GlobalVarMap, VarMap.GlobalVarMap, TemplateArgs);
   merge(ExternVarMap, VarMap.ExternVarMap, TemplateArgs);
+  merge(TempStorageMap, VarMap.TempStorageMap, TemplateArgs);
   dpct::merge(TextureMap, VarMap.TextureMap);
 }
 int MemVarMap::calculateExtraArgsSize() const {
@@ -3721,6 +3752,7 @@ MemVarMap::getArgumentsOrParameters(int PreParams, int PostParams, LocInfo LI,
   getArgumentsOrParametersFromMap<MemVarInfo, COD>(PS, GlobalVarMap, LI);
   getArgumentsOrParametersFromMap<MemVarInfo, COD>(PS, LocalVarMap, LI);
   getArgumentsOrParametersFromoTextureInfoMap<COD>(PS, TextureMap);
+  getArgumentsOrParametersFromMap<TempStorageVarInfo, COD>(PS, TempStorageMap);
   std::string Result = PS.Str;
   return (Result.empty() || PostParams != 0) && PreParams == 0
              ? Result
@@ -3804,13 +3836,17 @@ const MemVarInfoMap &MemVarMap::getMap(MemVarInfo::VarScope Scope) const {
 const GlobalMap<TextureInfo> &MemVarMap::getTextureMap() const {
   return TextureMap;
 }
+const GlobalMap<TempStorageVarInfo> &MemVarMap::getTempStorageMap() const {
+  return TempStorageMap;
+}
 void MemVarMap::removeDuplicateVar() {
   std::unordered_set<std::string> VarNames{getItemName(),
                                            DpctGlobalInfo::getStreamName()};
   dpct::removeDuplicateVar(GlobalVarMap, VarNames);
   dpct::removeDuplicateVar(LocalVarMap, VarNames);
   dpct::removeDuplicateVar(ExternVarMap, VarNames);
   dpct::removeDuplicateVar(TextureMap, VarNames);
+  dpct::removeDuplicateVar(TempStorageMap, VarNames);
 }
 MemVarInfoMap &MemVarMap::getMap(MemVarInfo::VarScope Scope) {
   switch (Scope) {
@@ -3884,16 +3920,6 @@ unsigned int MemVarMap::getHeadNodeDim() const {
   else
     return 3;
 }
-void MemVarMap::merge(MemVarInfoMap &Master, const MemVarInfoMap &Branch,
-                      const std::vector<TemplateArgumentInfo> &TemplateArgs) {
-  if (TemplateArgs.empty())
-    return dpct::merge(Master, Branch);
-  for (auto &VarInfoPair : Branch)
-    Master
-        .insert(std::make_pair(VarInfoPair.first, std::make_shared<MemVarInfo>(
-                                                      *VarInfoPair.second)))
-        .first->second->applyTemplateArguments(TemplateArgs);
-}
 int MemVarMap::calculateExtraArgsSize(const MemVarInfoMap &Map) const {
   int Size = 0;
   for (auto &VarInfoPair : Map) {
@@ -3907,16 +3933,18 @@ void MemVarMap::getArgumentsOrParametersFromMap(ParameterStream &PS,
                                                 const GlobalMap<T> &VarMap,
                                                 LocInfo LI) {
   for (const auto &VI : VarMap) {
-    if (!VI.second->isUseHelperFunc()) {
-      continue;
-    }
-    if (!VI.second->getType()->SharedVarInfo.TypeName.empty() &&
-        !LI.first.getCanonicalPath().empty() && LI.second) {
-      DiagnosticsUtils::report(
-          LI.first.getCanonicalPath().str(), LI.second,
-          Warnings::MOVE_TYPE_DEFINITION_DEVICE_FUNC, true, false,
-          VI.second->getType()->SharedVarInfo.TypeName,
-          VI.second->getType()->SharedVarInfo.DefinitionFuncName);
+    if constexpr (!std::is_same_v<T, TempStorageVarInfo>) {
+      if (!VI.second->isUseHelperFunc()) {
+        continue;
+      }
+      if (!VI.second->getType()->SharedVarInfo.TypeName.empty() &&
+          !LI.first.getCanonicalPath().empty() && LI.second) {
+        DiagnosticsUtils::report(
+            LI.first.getCanonicalPath().str(), LI.second,
+            Warnings::MOVE_TYPE_DEFINITION_DEVICE_FUNC, true, false,
+            VI.second->getType()->SharedVarInfo.TypeName,
+            VI.second->getType()->SharedVarInfo.DefinitionFuncName);
+      }
     }
     if (PS.FormatInformation.EnableFormat) {
       ParameterStream TPS;
@@ -3970,6 +3998,8 @@ void MemVarMap::getArgumentsOrParametersForDecl(ParameterStream &PS,
       PS, LocalVarMap, LI);
   getArgumentsOrParametersFromoTextureInfoMap<MemVarMap::DeclParameter>(
       PS, TextureMap);
+  getArgumentsOrParametersFromMap<TempStorageVarInfo, MemVarMap::DeclParameter>(
+      PS, TempStorageMap);
 }
 ///// class CallFunctionExpr /////
 void CallFunctionExpr::buildCallExprInfo(const CXXConstructExpr *Ctor) {
@@ -5478,6 +5508,10 @@ void KernelCallExpr::addAccessorDecl() {
     Tex.second->addDecl(SubmitStmts.TextureList, SubmitStmts.SamplerList,
                         getQueueStr());
   }
+  for (auto &Tmp : VM.getTempStorageMap()) {
+    Tmp.second->addAccessorDecl(SubmitStmts.AccessorList,
+                                ExecutionConfig.LocalSize);
+  }
 }
 void KernelCallExpr::buildInfo() {
   CallFunctionExpr::buildInfo();

clang/lib/DPCT/AnalysisInfo.h

@@ -21,6 +21,7 @@
 #include "Utility.h"
 #include "ValidateArguments.h"
 #include <bitset>
+#include <memory>
 #include <optional>
 #include <unordered_set>
 #include <vector>
@@ -2139,6 +2140,24 @@ class TemplateArgumentInfo {
   bool IsWritten = true;
 };
 
+class TempStorageVarInfo {
+  unsigned Offset;
+  std::string Name;
+  std::shared_ptr<TemplateDependentStringInfo> Type;
+
+public:
+  TempStorageVarInfo(unsigned Off, StringRef Name,
+                     std::shared_ptr<TemplateDependentStringInfo> T)
+      : Offset(Off), Name(Name.str()), Type(T) {}
+  const std::string &getName() const { return Name; }
+  unsigned getOffset() const { return Offset; }
+  void addAccessorDecl(StmtList &AccessorList, StringRef LocalSize) const;
+  void applyTemplateArguments(const std::vector<TemplateArgumentInfo> &TA);
+  ParameterStream &getFuncDecl(ParameterStream &PS);
+  ParameterStream &getFuncArg(ParameterStream &PS);
+  ParameterStream &getKernelArg(ParameterStream &PS);
+};
+
 // memory variable map includes memory variable used in __global__/__device__
 // function and call expression.
 class MemVarMap {
@@ -2162,6 +2181,7 @@ class MemVarMap {
   void setBF64(bool Has = true) { HasBF64 = Has; }
   void setBF16(bool Has = true) { HasBF16 = Has; }
   void setGlobalMemAcc(bool Has = true) { HasGlobalMemAcc = Has; }
+  void addCUBTempStorage(std::shared_ptr<TempStorageVarInfo> Tmp);
   void addTexture(std::shared_ptr<TextureInfo> Tex);
   void addVar(std::shared_ptr<MemVarInfo> Var);
   void merge(const MemVarMap &OtherMap);
@@ -2197,6 +2217,7 @@ class MemVarMap {
                                  const clang::tooling::UnifiedPath &Path) const;
   const MemVarInfoMap &getMap(MemVarInfo::VarScope Scope) const;
   const GlobalMap<TextureInfo> &getTextureMap() const;
+  const GlobalMap<TempStorageVarInfo> &getTempStorageMap() const;
   void removeDuplicateVar();
 
   MemVarInfoMap &getMap(MemVarInfo::VarScope Scope);
@@ -2208,8 +2229,17 @@ class MemVarMap {
   unsigned int getHeadNodeDim() const;
 
 private:
-  static void merge(MemVarInfoMap &Master, const MemVarInfoMap &Branch,
-                    const std::vector<TemplateArgumentInfo> &TemplateArgs);
+  template <class VarT>
+  static void merge(GlobalMap<VarT> &Master, const GlobalMap<VarT> &Branch,
+                    const std::vector<TemplateArgumentInfo> &TemplateArgs) {
+    if (TemplateArgs.empty())
+      return dpct::merge(Master, Branch);
+    for (auto &VarInfoPair : Branch)
+      Master
+          .insert(std::make_pair(VarInfoPair.first,
+                                 std::make_shared<VarT>(*VarInfoPair.second)))
+          .first->second->applyTemplateArguments(TemplateArgs);
+  }
   int calculateExtraArgsSize(const MemVarInfoMap &Map) const;
 
   template <CallOrDecl COD>
@@ -2258,6 +2288,7 @@ class MemVarMap {
   MemVarInfoMap GlobalVarMap;
   MemVarInfoMap ExternVarMap;
   GlobalMap<TextureInfo> TextureMap;
+  GlobalMap<TempStorageVarInfo> TempStorageMap;
 };
 
 template <>

clang/lib/DPCT/CUBAPIMigration.cpp

@@ -12,14 +12,22 @@
 #include "CallExprRewriter.h"
 #include "ExprAnalysis.h"
 #include "MigrationRuleManager.h"
+#include "TextModification.h"
+#include "Utility.h"
+#include "clang/AST/Attrs.inc"
 #include "clang/AST/Decl.h"
 #include "clang/AST/DeclCXX.h"
+#include "clang/AST/DeclTemplate.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
+#include "clang/AST/OperationKinds.h"
 #include "clang/AST/Stmt.h"
+#include "clang/AST/Type.h"
 #include "clang/AST/TypeLoc.h"
+#include "clang/ASTMatchers/ASTMatchFinder.h"
 #include "clang/ASTMatchers/ASTMatchers.h"
 #include "clang/Analysis/AnalysisDeclContext.h"
+#include "clang/Basic/AttrKinds.h"
 #include "clang/Basic/CharInfo.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Tooling/Tooling.h"
@@ -32,6 +40,9 @@
 #include "llvm/Support/Path.h"
 #include "llvm/Support/raw_ostream.h"
 #include <iterator>
+#include <memory>
+#include <optional>
+#include <vector>
 
 using namespace clang;
 using namespace dpct;
@@ -1164,6 +1175,99 @@ void CubRule::processBlockLevelMemberCall(const CXXMemberCallExpr *BlockMC) {
       OpRepl = getOpRepl(FuncArgs[1]);
       IsPartialReduce = NumArgs == 3;
       ValidItemParamIdx = 2;
+      const auto *CK = dyn_cast<ImplicitCastExpr>(FuncArgs[1]);
+      if (DpctGlobalInfo::useUserDefineReductions() && OpRepl.empty() && CK &&
+          CK->getCastKind() == CK_FunctionToPointerDecay) {
+        ExprAnalysis EA;
+        EA.analyze(CK);
+        OpRepl =
+            "[](auto&& x, auto&& y) { return " + EA.getReplacedString() +
+            "(std::forward<decltype(x)>(x), std::forward<decltype(y)>(y)); }";
+
+        NewFuncName = MapNames::getClNamespace() +
+                      "ext::oneapi::experimental::reduce_over_group";
+        Expr *Obj = BlockMC->getImplicitObjectArgument();
+        const VarDecl *TempStorage = nullptr;
+
+        auto FindTempStorageVarInCtor = [&](const Expr *E) -> const VarDecl * {
+          if (auto *Ctor = dyn_cast<CXXConstructExpr>(E)) {
+            if (auto *DRE = dyn_cast<DeclRefExpr>(Ctor->getArg(0))) {
+              if (auto *VD = dyn_cast<VarDecl>(DRE->getDecl())) {
+                if (VD->hasAttr<CUDASharedAttr>() && isCubVar(VD)) {
+                  return VD;
+                }
+              }
+            }
+          }
+          return nullptr;
+        };
+
+        auto HandleTypeLoc = [&](TypeLoc Loc) -> TypeLoc {
+          if (Loc.isNull())
+            return Loc;
+          while (true) {
+            switch (Loc.getTypeLocClass()) {
+            case TypeLoc::Elaborated:
+              Loc = Loc.getNextTypeLoc();
+              break;
+            case TypeLoc::Typedef: {
+              auto NewLoc = Loc.castAs<TypedefTypeLoc>();
+              Loc = NewLoc.getTypedefNameDecl()
+                        ->getTypeSourceInfo()
+                        ->getTypeLoc();
+              break;
+            }
+            case TypeLoc::TemplateSpecialization: {
+              auto NewLoc = Loc.getAs<TemplateSpecializationTypeLoc>();
+              return NewLoc.getArgLocInfo(0).getAsTypeSourceInfo()->getTypeLoc();
+              break;
+            }
+            default:
+              return Loc;
+            }
+          }
+        };
+
+        TypeLoc DataTypeLoc;
+        if (const auto *MTE = dyn_cast<MaterializeTemporaryExpr>(Obj)) {
+          if (auto *TOE = dyn_cast<CXXTemporaryObjectExpr>(MTE->getSubExpr())) {
+            DataTypeLoc = HandleTypeLoc(TOE->getTypeSourceInfo()->getTypeLoc());
+          } else if (auto *FC =
+                         dyn_cast<CXXFunctionalCastExpr>(MTE->getSubExpr())) {
+            DataTypeLoc =
+                HandleTypeLoc(FC->getTypeInfoAsWritten()->getTypeLoc());
+          }
+          TempStorage =
+              FindTempStorageVarInCtor(MTE->getSubExpr()->IgnoreCasts());
+        } else if (const auto *DRE = dyn_cast<DeclRefExpr>(Obj)) {
+          if (const auto *VD = dyn_cast<VarDecl>(DRE->getDecl())) {
+            DataTypeLoc = HandleTypeLoc(VD->getTypeSourceInfo()->getTypeLoc());
+            if (isCubCollectiveRecordType(VD->getType()) && VD->hasInit()) {
+              emplaceTransformation(new ReplaceVarDecl(VD, ""));
+              TempStorage = FindTempStorageVarInCtor(VD->getInit());
+            }
+          }
+        }
+
+        auto *FD = DpctGlobalInfo::findAncestor<FunctionDecl>(TempStorage);
+        if (!FD || !TempStorage || DataTypeLoc.isNull())
+          return;
+        if (auto FuncInfo = DeviceFunctionDecl::LinkRedecls(FD)) {
+          auto LocInfo = DpctGlobalInfo::getLocInfo(TempStorage);
+          ExprAnalysis EA;
+          EA.analyze(DataTypeLoc);
+          FuncInfo->getVarMap().addCUBTempStorage(
+              std::make_shared<TempStorageVarInfo>(
+                  LocInfo.second, TempStorage->getName(),
+                  EA.getTemplateDependentStringInfo()));
+        }
+        std::string Span = MapNames::getClNamespace() + "span<std::byte, 1>" +
+                           "(&" + TempStorage->getNameAsString() + "[0], " +
+                           TempStorage->getNameAsString() + ".size())";
+        GroupOrWorkitem = MapNames::getClNamespace() +
+                          "ext::oneapi::experimental::group_with_scratchpad(" +
+                          GroupOrWorkitem + ", " + Span + ")";
+      }
     } else if (FuncName == "Sum") {
       OpRepl = getOpRepl(nullptr);
       IsPartialReduce = NumArgs == 2;

clang/lib/DPCT/DPCT.cpp

@@ -871,6 +871,12 @@ int runDPCT(int argc, const char **argv) {
     }
 
     Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("-w"));
+#ifdef _WIN32 // Avoid some error on windows platform.
+    if (DpctGlobalInfo::getSDKVersion() <= CudaVersion::CUDA_100) {
+      Tool.appendArgumentsAdjuster(
+          getInsertArgumentAdjuster("-D_MSC_VER=1900"));
+    }
+#endif
     NoIncrementalMigration.setValue(true);
     StopOnParseErr.setValue(true);
     Tool.setPrintErrorMessage(false);