diff --git a/clang/CMakeLists.txt b/clang/CMakeLists.txt index 77f25217fabb..64e6ec0b647a 100644 --- a/clang/CMakeLists.txt +++ b/clang/CMakeLists.txt @@ -911,6 +911,7 @@ if (NOT CLANG_BUILT_STANDALONE) list(APPEND DPCT_RUN install-dpct-intercept-build install-dpct-pattern-rewriter + install-dpct-gitdiff2yaml install-dpct-headers install-clang-resource-headers install-dpct-binary diff --git a/clang/include/clang/DPCT/DPCTOptions.inc b/clang/include/clang/DPCT/DPCTOptions.inc index 9173ba488dc2..4aaad0f86c4c 100644 --- a/clang/include/clang/DPCT/DPCTOptions.inc +++ b/clang/include/clang/DPCT/DPCTOptions.inc @@ -958,6 +958,12 @@ DPCT_FLAG_OPTION( "--out-root directory. Default: off."), llvm::cl::cat(CtHelpCatAll), llvm::cl::cat(CtHelpCatCodeGen)) +DPCT_FLAG_OPTION( + ReMigration, clang::dpct::DpctOptionClass::OC_Attribute, + DPCT_OPTION_ACTIONS(clang::dpct::DpctActionKind::DAK_Migration), + "remigration", llvm::cl::desc("Enable remigration. Default: off."), + llvm::cl::cat(CtHelpCatAll), llvm::cl::cat(CtHelpCatCodeGen)) + DPCT_FLAG_OPTION( PathToHelperFunction, clang::dpct::DpctOptionClass::OC_Action, DPCT_OPTION_ACTIONS(clang::dpct::DpctActionKind::DAK_Help), diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 383a4dea10db..4da711482b51 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2513,6 +2513,7 @@ unsigned DpctGlobalInfo::ExperimentalFlag = 0; unsigned DpctGlobalInfo::HelperFuncPreferenceFlag = 0; bool DpctGlobalInfo::AnalysisModeFlag = false; bool DpctGlobalInfo::UseSYCLCompatFlag = false; +bool DpctGlobalInfo::ReMigrationFlag = false; bool DpctGlobalInfo::CVersionCUDALaunchUsedFlag = false; unsigned int DpctGlobalInfo::ColorOption = 1; std::unordered_map> diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index c4c45a27a769..4446f3a62bc4 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1349,6 +1349,8 @@ class DpctGlobalInfo { static bool isCVersionCUDALaunchUsed() { return CVersionCUDALaunchUsedFlag; } static void setUseSYCLCompat(bool Flag = true) { UseSYCLCompatFlag = Flag; } static bool useSYCLCompat() { return UseSYCLCompatFlag; } + static void setReMigration(bool Flag = true) { ReMigrationFlag = Flag; } + static bool useReMigration() { return ReMigrationFlag; } static bool useEnqueueBarrier() { return getUsingExtensionDE( DPCPPExtensionsDefaultEnabled::ExtDE_EnqueueBarrier); @@ -1682,6 +1684,7 @@ class DpctGlobalInfo { static unsigned HelperFuncPreferenceFlag; static bool AnalysisModeFlag; static bool UseSYCLCompatFlag; + static bool ReMigrationFlag; static bool CVersionCUDALaunchUsedFlag; static unsigned int ColorOption; static std::unordered_map> diff --git a/clang/lib/DPCT/CMakeLists.txt b/clang/lib/DPCT/CMakeLists.txt index b74c77697383..2802c61f9863 100644 --- a/clang/lib/DPCT/CMakeLists.txt +++ b/clang/lib/DPCT/CMakeLists.txt @@ -233,6 +233,7 @@ add_clang_library(DPCT MigrateScript/GenMakefile.cpp RulesInclude/InclusionHeaders.cpp IncMigration/IncrementalMigrationUtility.cpp + IncMigration/ReMigration.cpp UserDefinedRules/UserDefinedRules.cpp UserDefinedRules/PatternRewriter.cpp MigrateScript/MigrateBuildScript.cpp diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp index ca70d9886928..e58bf5cdfda8 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -79,7 +79,7 @@ using namespace clang::tooling; using namespace llvm::cl; extern bool isDPCT; - +extern bool ReMigrationReady; namespace clang { namespace tooling { UnifiedPath getFormatSearchPath(); @@ -516,7 +516,7 @@ static void loadMainSrcFileInfo(clang::tooling::UnifiedPath YamlFilePath) { if (!llvm::sys::fs::exists(YamlFilePath.getCanonicalPath())) return; auto PreTU = std::make_shared(); - if (loadFromYaml(YamlFilePath, *PreTU) != 0) { + if (loadTUFromYaml(YamlFilePath, *PreTU) != 0) { llvm::errs() << getLoadYamlFailWarning(); } DpctGlobalInfo::setMainSourceYamlTUR(PreTU); @@ -804,6 +804,7 @@ int runDPCT(int argc, const char **argv) { CudaIncludePath = CudaInclude; SDKPath = SDKPathOpt; + DpctGlobalInfo::setReMigration(ReMigration); loadMainSrcFileInfo(OutRootPath.getCanonicalPath() + "/MainSourceFiles.yaml"); std::transform( @@ -1391,6 +1392,8 @@ int runDPCT(int argc, const char **argv) { dpctExit(MigrationSucceeded, false); } + ReMigrationReady = tryLoadingUpstreamChangesAndUserChanges(); + ReplTy ReplCUDA, ReplSYCL; volatile int RunCount = 0; do { diff --git a/clang/lib/DPCT/FileGenerator/GenFiles.cpp b/clang/lib/DPCT/FileGenerator/GenFiles.cpp index b0f933aa6d4c..1573ba9fccfe 100644 --- a/clang/lib/DPCT/FileGenerator/GenFiles.cpp +++ b/clang/lib/DPCT/FileGenerator/GenFiles.cpp @@ -48,6 +48,8 @@ extern DpctOption ProcessAll; extern DpctOption BuildScriptFile; extern DpctOption GenBuildScript; extern std::map ErrorCnt; +bool ReMigrationReady = false; +extern std::set MainSrcFilesHasCudaSyntex; namespace clang { namespace tooling { @@ -188,18 +190,10 @@ bool rewriteCanonicalDir(clang::tooling::UnifiedPath &FilePath, return Result; } -void rewriteFileName(clang::tooling::UnifiedPath &FileName) { - rewriteFileName(FileName, FileName); -} - -void rewriteFileName(clang::tooling::UnifiedPath &FileName, - const clang::tooling::UnifiedPath &FullPathName) { - std::string FilePath = FileName.getPath().str(); - rewriteFileName(FilePath, FullPathName.getPath().str()); - FileName = FilePath; -} - -void rewriteFileName(std::string &FileName, const std::string &FullPathName) { +static void +rewriteFileName(std::string &FileName, const std::string &FullPathName, + std::function HasCUDASyntax = + DpctGlobalInfo::hasCUDASyntax) { SmallString<512> CanonicalPathStr(FullPathName); const auto Extension = path::extension(CanonicalPathStr); SourceProcessType FileType = GetSourceFileType(FullPathName); @@ -210,8 +204,7 @@ void rewriteFileName(std::string &FileName, const std::string &FullPathName) { if (FileType & SPT_CudaSource) { path::replace_extension(CanonicalPathStr, DpctGlobalInfo::getSYCLSourceExtension()); - } else if ((FileType & SPT_CppSource) && - DpctGlobalInfo::hasCUDASyntax(FileName)) { + } else if ((FileType & SPT_CppSource) && HasCUDASyntax(FileName)) { path::replace_extension(CanonicalPathStr, Extension + DpctGlobalInfo::getSYCLSourceExtension()); @@ -224,6 +217,17 @@ void rewriteFileName(std::string &FileName, const std::string &FullPathName) { FileName = CanonicalPathStr.c_str(); } +void rewriteFileName(clang::tooling::UnifiedPath &FileName, + const clang::tooling::UnifiedPath &FullPathName) { + std::string FilePath = FileName.getPath().str(); + rewriteFileName(FilePath, FullPathName.getPath().str()); + FileName = FilePath; +} + +void rewriteFileName(clang::tooling::UnifiedPath &FileName) { + rewriteFileName(FileName, FileName); +} + static std::vector FilesNotInCompilationDB; std::map OutFilePath2InFilePath; @@ -506,7 +510,7 @@ static void getReplsFromTUR(const std::string &FilePath, } int writeReplacementsToFiles( - ReplTy &Replset, Rewriter &Rewrite, const std::string &Folder, + ReplTy &Replset2, Rewriter &Rewrite, const std::string &Folder, clang::tooling::UnifiedPath &InRoot, std::unordered_map &MainSrcFileMap, std::vector &AllFilesRepls, @@ -520,6 +524,48 @@ int writeReplacementsToFiles( volatile ProcessStatus status = MigrationSucceeded; clang::tooling::UnifiedPath OutPath; + ReplTy Replset; + if (ReMigrationReady) { + std::vector Repl_B; + std::vector Repl_C1; + for (const auto &Entry : Replset2) { + for (const auto &Repl : Entry.second) { + Repl_B.push_back(Repl); + } + } + for (const auto &Repl : clang::dpct::getLastMigration()->Replacements) { + Repl_C1.push_back(Repl); + } + + std::map FileNameMap; + auto hasCUDASyntax = [](tooling::UnifiedPath Path) -> bool { + if (MainSrcFilesHasCudaSyntex.find(Path.getCanonicalPath().str()) != + MainSrcFilesHasCudaSyntex.end()) + return true; + return false; + }; + for (const auto &Entry : Repl_C1) { + std::string CUDAFilePath = Entry.getFilePath().str(); + std::string SYCLFilePath; + rewriteFileName(SYCLFilePath, CUDAFilePath, hasCUDASyntax); + FileNameMap[SYCLFilePath] = CUDAFilePath; + } + std::map> Result = + clang::dpct::reMigrationMerge(clang::dpct::getUpstreamChanges(), Repl_B, + Repl_C1, clang::dpct::getUserChanges(), + FileNameMap); + for (const auto &Entry : Result) { + clang::tooling::Replacements Repls; + for (const auto &Repl : Entry.second) { + llvm::cantFail(Repls.add(Repl)); + } + Replset.insert(std::make_pair(Entry.first, Repls)); + } + } else { + // If not re-migration, just use the original Replset. + Replset = Replset2; + } + for (auto &Entry : Replset) { OutPath = StringRef(DpctGlobalInfo::removeSymlinks( Rewrite.getSourceMgr().getFileManager(), Entry.first)); @@ -633,6 +679,7 @@ int writeReplacementsToFiles( // We have written a migrated file; Update the output file path info OutFilePath2InFilePath[OutPath.getCanonicalPath().str()] = Entry.first; } + Replset2 = Replset; return status; } diff --git a/clang/lib/DPCT/FileGenerator/GenFiles.h b/clang/lib/DPCT/FileGenerator/GenFiles.h index 828d5c439aa8..1ba8fdda4830 100644 --- a/clang/lib/DPCT/FileGenerator/GenFiles.h +++ b/clang/lib/DPCT/FileGenerator/GenFiles.h @@ -76,9 +76,6 @@ void rewriteFileName(clang::tooling::UnifiedPath &FileName); void rewriteFileName(clang::tooling::UnifiedPath &FileName, const clang::tooling::UnifiedPath &FullPathName); -// Replace file name \p FileName with new migrated name. -void rewriteFileName(std::string &FileName, const std::string &FullPathName); - // A mapping from output file path to it's corresponding input file. extern std::map OutFilePath2InFilePath; diff --git a/clang/lib/DPCT/IncMigration/ExternalReplacement.cpp b/clang/lib/DPCT/IncMigration/ExternalReplacement.cpp index 4d84c58a6169..e5206c9b938f 100644 --- a/clang/lib/DPCT/IncMigration/ExternalReplacement.cpp +++ b/clang/lib/DPCT/IncMigration/ExternalReplacement.cpp @@ -10,19 +10,19 @@ // -Load replacement from external (disk file) // -Merge replacement in current migration with previous migration. +#include "ExternalReplacement.h" #include "AnalysisInfo.h" +#include "IncMigration/IncrementalMigrationUtility.h" #include "Utility.h" + +#include "clang/Tooling/Core/Diagnostic.h" #include "clang/Tooling/Core/Replacement.h" +#include "clang/Tooling/Refactoring.h" +#include "clang/Tooling/ReplacementsYaml.h" #include "llvm/ADT/SmallString.h" #include "llvm/ADT/StringRef.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/Path.h" - -#include "ExternalReplacement.h" -#include "IncMigration/IncrementalMigrationUtility.h" -#include "clang/Tooling/Core/Diagnostic.h" -#include "clang/Tooling/Refactoring.h" -#include "clang/Tooling/ReplacementsYaml.h" #include "llvm/Support/YAMLTraits.h" #include "llvm/Support/raw_os_ostream.h" @@ -77,8 +77,8 @@ int save2Yaml( return 0; } -int loadFromYaml(const clang::tooling::UnifiedPath &Input, - clang::tooling::TranslationUnitReplacements &TU) { +template +static int loadFromYaml(const clang::tooling::UnifiedPath &Input, T &Content) { llvm::ErrorOr> Buffer = llvm::MemoryBuffer::getFile(Input.getCanonicalPath()); if (!Buffer) { @@ -86,9 +86,20 @@ int loadFromYaml(const clang::tooling::UnifiedPath &Input, << Buffer.getError().message() << "\n"; return -1; } - llvm::yaml::Input YAMLIn(Buffer.get()->getBuffer()); - YAMLIn >> TU; + YAMLIn >> Content; + if (YAMLIn.error()) { + Content = T(); + return -1; + } + return 0; +} + +int loadTUFromYaml(const clang::tooling::UnifiedPath &Input, + clang::tooling::TranslationUnitReplacements &TU) { + int Status = loadFromYaml(Input, TU); + if (Status) + return Status; bool IsSrcFileChanged = false; for (const auto &digest : TU.MainSourceFilesDigest) { @@ -100,7 +111,7 @@ int loadFromYaml(const clang::tooling::UnifiedPath &Input, } } - if (IsSrcFileChanged || YAMLIn.error()) { + if (IsSrcFileChanged && !DpctGlobalInfo::useReMigration()) { // File doesn't appear to be a header change description. Ignore it. TU = clang::tooling::TranslationUnitReplacements(); return -1; @@ -109,6 +120,16 @@ int loadFromYaml(const clang::tooling::UnifiedPath &Input, return 0; } +void loadGDCFromYaml(const clang::tooling::UnifiedPath &Input, + clang::dpct::GitDiffChanges &GDC) { + int status = loadFromYaml(Input, GDC); + if (status) { + llvm::errs() << "Failed to load git diff Changes from " + << Input.getCanonicalPath() << "\n"; + GDC = clang::dpct::GitDiffChanges(); + } +} + void mergeAndUniqueReps( Replacements &Replaces, const std::vector &PreRepls) { diff --git a/clang/lib/DPCT/IncMigration/ExternalReplacement.h b/clang/lib/DPCT/IncMigration/ExternalReplacement.h index b672ccfbe8ba..93a4d20191ff 100644 --- a/clang/lib/DPCT/IncMigration/ExternalReplacement.h +++ b/clang/lib/DPCT/IncMigration/ExternalReplacement.h @@ -9,8 +9,8 @@ #ifndef __EXTERNAL_REPLACEMENT_H__ #define __EXTERNAL_REPLACEMENT_H__ -#include "clang/Tooling/Core/Replacement.h" -#include "llvm/ADT/StringRef.h" +#include "IncMigration/ReMigration.h" + #include #include @@ -27,8 +27,10 @@ class Replacements; namespace clang { namespace dpct { -int loadFromYaml(const clang::tooling::UnifiedPath &Input, - clang::tooling::TranslationUnitReplacements &TU); +int loadTUFromYaml(const clang::tooling::UnifiedPath &Input, + clang::tooling::TranslationUnitReplacements &TU); +void loadGDCFromYaml(const clang::tooling::UnifiedPath &Input, + clang::dpct::GitDiffChanges &GDC); int save2Yaml( const std::vector &Replaces, const std::map &PreRepls); +bool tryLoadingUpstreamChangesAndUserChanges(); } // namespace dpct } // namespace clang diff --git a/clang/lib/DPCT/IncMigration/ReMigration.cpp b/clang/lib/DPCT/IncMigration/ReMigration.cpp new file mode 100644 index 000000000000..9ea94180b420 --- /dev/null +++ b/clang/lib/DPCT/IncMigration/ReMigration.cpp @@ -0,0 +1,802 @@ +//===----------------------- ReMigration.cpp ------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Workflow: +// CUDA code v1 +// | +// | dpct (--format-range=off) => MainSourceFiles.yaml +// v +// SYCL code v1 +// | 1. git init and commit +// | 2. manual format with clang-format +// | 3. manual fix +// v 4. gitdiff2yaml => UserChange.yaml +// SYCL code v1.1 +// +// CUDA code v2 +// | +// | dpct (--format-range=off) +// | and MainSourceFiles.yaml +// | and UserChange.yaml +// | and UpStreamChange.yaml +// v +// SYCL code v2 +//===----------------------------------------------------------------------===// + +#include "ReMigration.h" +#include "AnalysisInfo.h" +#include "ExternalReplacement.h" +#include "TextModification.h" + +#include "clang/AST/Expr.h" +#include "clang/Tooling/Core/Replacement.h" +#include "clang/Tooling/Core/UnifiedPath.h" +#include "llvm/ADT/StringRef.h" +#include "llvm/Support/raw_ostream.h" + +#include +#include +#include + +std::optional< + std::function> + getLineStringHook = std::nullopt; +std::optional> + getLineNumberHook = std::nullopt; +std::optional> + getLineBeginOffsetHook = std::nullopt; + +namespace clang::dpct { +using namespace clang::tooling; +static GitDiffChanges UpstreamChanges; +static GitDiffChanges UserChanges; +static std::shared_ptr LastMigration = nullptr; +AddFileHunk::AddFileHunk(std::string NewFilePath) + : Hunk(AddFile), NewFilePath(UnifiedPath(NewFilePath).getCanonicalPath()) {} +DeleteFileHunk::DeleteFileHunk(std::string OldFilePath) + : Hunk(DeleteFile), + OldFilePath(UnifiedPath(OldFilePath).getCanonicalPath()) {} +GitDiffChanges &getUpstreamChanges() { return UpstreamChanges; } +GitDiffChanges &getUserChanges() { return UserChanges; } +std::shared_ptr &getLastMigration() { + return LastMigration; +} +static void dumpGitDiffChanges(const GitDiffChanges &GHC) { + llvm::errs() << "GitDiffChanges:\n"; + llvm::errs() << " ModifyFileHunks:\n"; + for (const auto &Hunk : GHC.ModifyFileHunks) { + llvm::errs() << " - FilePath: " << Hunk.getFilePath() << "\n"; + llvm::errs() << " Offset: " << Hunk.getOffset() << "\n"; + llvm::errs() << " Length: " << Hunk.getLength() << "\n"; + llvm::errs() << " ReplacementText: " << Hunk.getReplacementText() + << "\n"; + } + llvm::errs() << " AddFileHunks:\n"; + for (const auto &Hunk : GHC.AddFileHunks) { + llvm::errs() << " - NewFilePath: " << Hunk.getNewFilePath() << "\n"; + } + llvm::errs() << " DeleteFileHunks:\n"; + for (const auto &Hunk : GHC.DeleteFileHunks) { + llvm::errs() << " - OldFilePath: " << Hunk.getOldFilePath() << "\n"; + } + llvm::errs() << " MoveFileHunks:\n"; + for (const auto &Hunk : GHC.MoveFileHunks) { + llvm::errs() << " - FilePath: " << Hunk.getFilePath() << "\n"; + llvm::errs() << " Offset: " << Hunk.getOffset() << "\n"; + llvm::errs() << " Length: " << Hunk.getLength() << "\n"; + llvm::errs() << " ReplacementText: " << Hunk.getReplacementText() + << "\n"; + llvm::errs() << " NewFilePath: " << Hunk.getNewFilePath() << "\n"; + } +} + +bool tryLoadingUpstreamChangesAndUserChanges() { + unsigned int Found = 0; + llvm::SmallString<128> UpstreamChangesFilePath( + DpctGlobalInfo::getInRoot().getCanonicalPath()); + llvm::SmallString<128> UserChangesFilePath( + DpctGlobalInfo::getInRoot().getCanonicalPath()); + llvm::SmallString<128> LastMigrationFilePath( + DpctGlobalInfo::getInRoot().getCanonicalPath()); + llvm::sys::path::append(UpstreamChangesFilePath, "UpstreamChanges.yaml"); + llvm::sys::path::append(UserChangesFilePath, "UserChanges.yaml"); + llvm::sys::path::append(LastMigrationFilePath, "LastMigration.yaml"); + + if (llvm::sys::fs::exists(UpstreamChangesFilePath)) { + ::loadGDCFromYaml(UpstreamChangesFilePath, getUpstreamChanges()); + Found++; + } + if (llvm::sys::fs::exists(UserChangesFilePath)) { + ::loadGDCFromYaml(UserChangesFilePath, getUserChanges()); + Found++; + } + getLastMigration() = DpctGlobalInfo::getMainSourceYamlTUR(); + return (Found == 2) ? true : false; +} + +/// Below 3 functions only used for merging Repl_B and repl_D. So they only +/// need to read CUDA code v2. +static StringRef getLineString(UnifiedPath FilePath, unsigned LineNumber) { +#ifndef NDEBUG + if (getLineStringHook.has_value()) { + return getLineStringHook.value()(FilePath, LineNumber); + } +#endif + auto FileInfo = DpctGlobalInfo::getInstance().insertFile(FilePath); + StringRef Line = FileInfo->getLineString(LineNumber); + return Line; +} + +static unsigned getLineNumber(UnifiedPath FilePath, unsigned Offset) { +#ifndef NDEBUG + if (getLineNumberHook.has_value()) { + return getLineNumberHook.value()(FilePath, Offset); + } +#endif + auto FileInfo = DpctGlobalInfo::getInstance().insertFile(FilePath); + return FileInfo->getLineNumber(Offset); +} + +static unsigned getLineBeginOffset(UnifiedPath FilePath, unsigned LineNumber) { +#ifndef NDEBUG + if (getLineBeginOffsetHook.has_value()) { + return getLineBeginOffsetHook.value()(FilePath, LineNumber); + } +#endif + auto FileInfo = DpctGlobalInfo::getInstance().insertFile(FilePath); + return FileInfo->getLineInfo(LineNumber).Offset; +} + +/// Calculate the new Repls of the input \p NewRepl after \p Repls is applied to +/// the files. +/// This shfit may have conflicts. +/// Since \p NewRepl (from Repl_C_y) is line based and \p Repls (from Repl_A) is +/// character based, we assume that if there is a conflict, the range from +/// Repl_A is always covering the range from Repl_C_y. Then we just ignore the +/// \p NewRepl (from Repl_C_y) since the old CUDA code is changed, so the +/// migration repl is out-of-date. +/// \param Repls Replacements to apply. +/// \param NewRepl Replacements before applying \p Repls. +/// \return The result Repls. +std::vector +calculateUpdatedRanges(const clang::tooling::Replacements &Repls, + const std::vector &NewRepl) { + // Assumption: no overlap in the each groups. + std::vector Result; + for (const auto &R : NewRepl) { + // Check if the range (BOffset, EOffset - BOffset) is overlapped with any + // repl in Repls + std::optional MaxNotGreater = std::nullopt; + for (const auto &ExistingR : Repls) { + if (ExistingR.getOffset() <= R.getOffset()) + MaxNotGreater = ExistingR; + else + break; + } + if (MaxNotGreater.has_value()) { + if (MaxNotGreater->getOffset() + MaxNotGreater->getLength() > + R.getOffset()) + continue; // has overlap + } + + unsigned int BOffset = Repls.getShiftedCodePosition(R.getOffset()); + unsigned int EOffset = + Repls.getShiftedCodePosition(R.getOffset() + R.getLength()); + if (BOffset > EOffset) + continue; + Result.push_back({R.getFilePath(), BOffset, EOffset - BOffset, + R.getReplacementText(), R.containsManualFix()}); + } + return Result; +} + +std::map> +groupReplcementsByFile(const std::vector &Repls) { + std::map> Result; + for (const auto &R : Repls) { + Result[R.getFilePath().str()].push_back(R); + } + return Result; +} + +// If repl range is cross lines, we treat the \n itself belongs to current line. +// case: +/* +aaaaa +bbbbb +ccccc +ddddd +*/ +// replacement is: offset:6, length:11, text:xxx\nyyy +// +// If the splited repl of line 2 does not end with \n, if there is conflict at +// this line, the mark will not correct. +// In other words, if there is no conflict, the repl is correct since there is +// \n after the repl. While the conflict mark assumes that the repl is endding +// with \n so it can at the beginning of each line, this is the problem. +// So the solution is: +// 1. If the newtext is not ending with \n, we merge the left part of the last +// line into the repl (the first line). +// 2. If the newtext is ending with \n, we still keep the last line at the last +// line. +std::vector +splitReplInOrderToNotCrossLines(const std::vector &InRepls) { + std::string FilePath = InRepls[0].getFilePath().str(); + std::vector Result; + + for (const auto &Repl : InRepls) { + unsigned StartOffset = Repl.getOffset(); + unsigned EndOffset = StartOffset + Repl.getLength(); + unsigned StartLine = getLineNumber(FilePath, StartOffset); + unsigned EndLine = getLineNumber(FilePath, EndOffset); + bool Tag = Repl.containsManualFix(); + + if (StartLine == EndLine) { + // Single line replacement + Result.push_back(Repl); + continue; + } + + // Cross-line replacement + unsigned CurrentOffset = StartOffset; + + // The first line + // We need merge the left part of the last line into the replacement if the + // newText is not ending with \n. + bool IsFisrtLineEndingWithNL = Repl.getReplacementText().ends_with('\n'); + unsigned LineEndOffset = getLineBeginOffset(FilePath, StartLine + 1); + unsigned FirstLineLength = LineEndOffset - StartOffset; + TaggedReplacement ReplFisrtLine(FilePath, CurrentOffset, FirstLineLength, + Repl.getReplacementText(), Tag); + if (IsFisrtLineEndingWithNL) + Result.push_back(ReplFisrtLine); + CurrentOffset += FirstLineLength; + + // middle lines + for (unsigned Line = StartLine + 1; Line < EndLine; ++Line) { + LineEndOffset = getLineBeginOffset(FilePath, Line + 1); + unsigned LineLength = LineEndOffset - CurrentOffset; + Result.emplace_back(Repl.getFilePath(), CurrentOffset, LineLength, "", + Tag); + CurrentOffset += LineLength; + } + + // The last line + std::string LastLineStr = getLineString(FilePath, EndLine).str(); + unsigned LastLineLength = EndOffset - CurrentOffset; + if (IsFisrtLineEndingWithNL) { + if (LastLineLength > 0) { + Result.emplace_back(Repl.getFilePath(), CurrentOffset, LastLineLength, + "", Tag); + } + } else { + Result.emplace_back(Repl.getFilePath(), ReplFisrtLine.getOffset(), + ReplFisrtLine.getLength(), + ReplFisrtLine.getReplacementText().str() + + LastLineStr.substr(LastLineLength), + Tag); + Result.emplace_back(Repl.getFilePath(), CurrentOffset, LastLineStr.size(), + "", Tag); + } + } + + return Result; +} + +std::map> +convertReplcementsLineString(const std::vector &InRepls) { + std::vector Replacements = + splitReplInOrderToNotCrossLines(InRepls); + UnifiedPath FilePath(InRepls[0].getFilePath()); + + // group replacement by line + std::map> ReplacementsByLine; + for (const auto &Repl : Replacements) { + unsigned LineNum = getLineNumber(FilePath, Repl.getOffset()); + ReplacementsByLine[LineNum].push_back(Repl); + } + + // process each line + std::map> Result; + for (auto &[LineNum, Repls] : ReplacementsByLine) { + std::sort(Repls.begin(), Repls.end()); + + std::string OriginalLineStr = getLineString(FilePath, LineNum).str(); + unsigned LineStartOffset = getLineBeginOffset(FilePath, LineNum); + std::string NewLineStr; + bool ContainsManualFix = false; + unsigned Pos = 0; + for (const auto &Repl : Repls) { + unsigned StrOffset = Repl.getOffset() - LineStartOffset; + NewLineStr += OriginalLineStr.substr(Pos, StrOffset - Pos); + NewLineStr += Repl.getReplacementText().str(); + Pos = StrOffset + Repl.getLength(); + ContainsManualFix = ContainsManualFix || Repl.containsManualFix(); + } + NewLineStr += OriginalLineStr.substr(Pos); + Result[LineNum] = std::make_pair(NewLineStr, ContainsManualFix); + } + return Result; +} + +static std::map +convertReplcementsLineString(const std::vector &Repls) { + std::vector ReplsVec; + for (const auto &R : Repls) { + ReplsVec.push_back({R, true}); + } + auto Temp = convertReplcementsLineString(ReplsVec); + std::map Result; + // Drop the tag info + for (const auto &[LineNum, LineStr] : Temp) { + Result[LineNum] = LineStr.first; + } + return Result; +} + +std::vector +mergeMapsByLine(const std::map + &MapA /*CUDA v2 code migration repls*/, + const std::map> + &MapB /*CUDA v1 code migration + manual fix repls*/, + const UnifiedPath &FilePath) { + auto genReplacement = [&](unsigned LineNumber, + const std::string &LineContent) { + unsigned Offset = getLineBeginOffset(FilePath, LineNumber); + unsigned StrLen = getLineString(FilePath, LineNumber).size(); + return Replacement(FilePath.getCanonicalPath(), Offset, StrLen, + LineContent); + }; + + std::vector Result; + auto ItA = MapA.begin(); + auto ItB = MapB.begin(); + + while (ItA != MapA.end() || ItB != MapB.end()) { + if (ItA == MapA.end()) { + Result.push_back(genReplacement(ItB->first, ItB->second.first)); + ++ItB; + } else if (ItB == MapB.end()) { + Result.push_back(genReplacement(ItA->first, ItA->second)); + ++ItA; + } else if (ItA->first < ItB->first) { + Result.push_back(genReplacement(ItA->first, ItA->second)); + ++ItA; + } else if (ItB->first < ItA->first) { + Result.push_back(genReplacement(ItB->first, ItB->second.first)); + ++ItB; + } else { + // ItB->first == ItA->first + // Conflict resolution + + if (ItB->second.first == ItA->second) { + // No conflict, just add one of them + Result.push_back(genReplacement(ItA->first, ItA->second)); + ++ItA; + ++ItB; + continue; + } + if (!ItB->second.second) { + // If ItB->second does not contain manual fix, we can use ItA->second + // directly + Result.push_back(genReplacement(ItA->first, ItA->second)); + ++ItA; + ++ItB; + continue; + } + + // Conflict line(s) + std::vector ConflictA; + std::vector ConflictB; + unsigned ConflictOffset = getLineBeginOffset(FilePath, ItA->first); + unsigned ConflictLength = 0; + + // Collect continuous conflicting lines + unsigned CurrentLineNum = ItA->first; + while (ItA != MapA.end() && ItB != MapB.end() && + CurrentLineNum == ItA->first && CurrentLineNum == ItB->first) { + ConflictA.push_back(ItA->second); + ConflictB.push_back(ItB->second.first); + ConflictLength += getLineString(FilePath, ItA->first).size(); + ++ItA; + ++ItB; + ++CurrentLineNum; + } + + // generate merged string + std::string Merged = "<<<<<<<\n"; + for (const auto &L : ConflictA) + Merged += L; + Merged += "=======\n"; + for (const auto &L : ConflictB) + Merged += L; + Merged += ">>>>>>>\n"; + + Result.emplace_back(FilePath.getCanonicalPath(), ConflictOffset, + ConflictLength, Merged); + } + } + return Result; +} + +static bool hasConflict(const Replacement &R1, const Replacement &R2) { + if (R1.getFilePath() != R2.getFilePath()) + return false; + if (R1.getOffset() == R2.getOffset()) { + if (R1.getLength() && R2.getLength()) { + return true; + } + } + if ((R1.getOffset() < R2.getOffset() && + R1.getOffset() + R1.getLength() > R2.getOffset()) || + (R2.getOffset() < R1.getOffset() && + R2.getOffset() + R2.getLength() > R1.getOffset())) { + return true; + } + return false; +} + +namespace { +// Repl group A is the replacements generated by the first run of dpct +// Repl group B is the replacements written by the user manaully based on the +// SYCL code generated by the first run of dpct. + +// We name the original CUDA code TEXT0, the migrated code TEXT1, the manually +// fixed SYCL code TEXT2. + +// Generate a mapping: (offset in TEXT1) -> (offset in TEXT0). + +// Build the mappings based on the replacements in group A. +std::pair>, + std::vector>> +buildMapping(const std::vector &A) { + std::vector> OffsetMap; + std::vector> ModifiedRangeList; + unsigned OrigPos = 0; + unsigned MigratedCodePos = 0; + + auto SortedA = A; + std::sort(SortedA.begin(), SortedA.end()); + + for (const auto &Repl : SortedA) { + // Add mapping before repl + if (Repl.getOffset() > OrigPos) { + unsigned len = Repl.getOffset() - OrigPos; + OffsetMap.emplace_back(MigratedCodePos, OrigPos); + MigratedCodePos += len; + OrigPos += len; + } + + // record the repl range + ModifiedRangeList.emplace_back(MigratedCodePos, Repl.getLength(), + Repl.getReplacementText().str()); + + // Add mapping after repl + OffsetMap.emplace_back(MigratedCodePos, Repl.getOffset()); + MigratedCodePos += Repl.getReplacementText().size(); + OrigPos += Repl.getLength(); + } + + // the last fragment + OffsetMap.emplace_back(MigratedCodePos, OrigPos); + return {OffsetMap, ModifiedRangeList}; +} + +// Map the offset in TEXT1 to the offset in TEXT0. +unsigned +mapToOriginalOffset(unsigned MigratedCodeOffset, + const std::vector> &Mapping) { + auto It = std::upper_bound( + Mapping.begin(), Mapping.end(), std::make_pair(MigratedCodeOffset, 0), + [](const auto &a, const auto &b) { return a.first < b.first; }); + + if (It != Mapping.begin()) + --It; + return It->second + (MigratedCodeOffset - It->first); +} + +std::vector +mergeC1AndC2Impl(const std::vector &A, + const std::vector &B) { + // Build 2 mappings + auto [OffsetMap, ModifiedRangeList] = buildMapping(A); + std::string FilePath = A[0].getFilePath().str(); + + std::vector Result; + std::vector AGroupRepls; + for (const auto &repl : ModifiedRangeList) { + AGroupRepls.push_back( + {FilePath, std::get<0>(repl), std::get<1>(repl), std::get<2>(repl)}); + } + std::vector BGroupRepls = B; + std::sort( + AGroupRepls.begin(), + AGroupRepls.end()); // The length in replA is based on original text, so + // for below calculation which is based on modified + // text, we need to use text.size() instead. + std::sort(BGroupRepls.begin(), BGroupRepls.end()); + + auto IterA = AGroupRepls.begin(); + auto IterB = BGroupRepls.begin(); + + std::optional< + std::tuple> + UnfinishedRepl = std::nullopt; + while (IterA != AGroupRepls.end() || IterB != BGroupRepls.end()) { + const bool NoA = IterA == AGroupRepls.end(); + const bool NoB = IterB == BGroupRepls.end(); + if (UnfinishedRepl) { + const unsigned Off = std::get<0>(*UnfinishedRepl); + const unsigned Len = std::get<1>(*UnfinishedRepl); + const std::string Text = std::get<2>(*UnfinishedRepl); + if ((NoA || (Off + Len) < IterA->getOffset()) && + (NoB || (Off + Len) < IterB->getOffset())) { + // Finished + Result.push_back({FilePath, mapToOriginalOffset(Off, OffsetMap), + mapToOriginalOffset(Off + Len, OffsetMap) - + mapToOriginalOffset(Off, OffsetMap), + Text, std::get<3>(*UnfinishedRepl)}); + UnfinishedRepl = std::nullopt; + continue; + } + if (!NoB && (Off + Len) >= IterB->getOffset()) { + // Unfinished repl overlapping with next B repl. + // We process B first because it has higher priority. + std::get<1>(*UnfinishedRepl) = + IterB->getLength() + IterB->getOffset() - Off; + std::get<2>(*UnfinishedRepl) = + Text.substr(0, IterB->getOffset() - Off) + + IterB->getReplacementText().str(); + std::get<3>(*UnfinishedRepl) = true; + ++IterB; + continue; + } + if (!NoA && (Off + Len) >= IterA->getOffset()) { + if ((Off + Len) >= + (IterA->getOffset() + IterA->getReplacementText().size())) { + // case 1: totally cover IterA + ++IterA; + continue; + } + // case 2: partially cover IterA + std::get<1>(*UnfinishedRepl) = + IterA->getReplacementText().size() + IterA->getOffset() - Off; + std::get<2>(*UnfinishedRepl) = + Text + IterA->getReplacementText().str().substr(Off + Len - + IterA->getOffset()); + ++IterA; + continue; + } + } + if (NoA) { + // only left B group replacements + Result.push_back( + {FilePath, mapToOriginalOffset(IterB->getOffset(), OffsetMap), + mapToOriginalOffset(IterB->getOffset() + IterB->getLength(), + OffsetMap) - + mapToOriginalOffset(IterB->getOffset(), OffsetMap), + IterB->getReplacementText(), true}); + ++IterB; + continue; + } + if (NoB) { + // only left A group replacements + Result.push_back({*IterA, false}); + ++IterA; + continue; + } + if (IterA->getOffset() < IterB->getOffset()) { + UnfinishedRepl = std::make_tuple( + IterA->getOffset(), IterA->getReplacementText().size(), + IterA->getReplacementText().str(), false); + ++IterA; + } else { + UnfinishedRepl = std::make_tuple(IterB->getOffset(), IterB->getLength(), + IterB->getReplacementText().str(), true); + ++IterB; + } + } + + // TODO: figure out why we need this? + if (UnfinishedRepl) { + Result.push_back( + {FilePath, mapToOriginalOffset(std::get<0>(*UnfinishedRepl), OffsetMap), + mapToOriginalOffset(std::get<0>(*UnfinishedRepl) + + std::get<1>(*UnfinishedRepl), + OffsetMap) - + mapToOriginalOffset(std::get<0>(*UnfinishedRepl), OffsetMap), + std::get<2>(*UnfinishedRepl), std::get<3>(*UnfinishedRepl)}); + } + + std::sort(Result.begin(), Result.end()); + return Result; +} +} // namespace + +// Merge Repl_C1 and Repl_C2. If has conflict, keep repl from Repl_C2. +std::vector mergeC1AndC2( + const std::vector &Repl_C1, const GitDiffChanges &Repl_C2, + const std::map + &FileNameMap) { + std::vector Result; + std::vector Repl_C2_vec; + std::for_each( + Repl_C2.ModifyFileHunks.begin(), Repl_C2.ModifyFileHunks.end(), + [&Repl_C2_vec, FileNameMap](const Replacement &Hunk) { + std::string OldFilePath = FileNameMap.at(Hunk.getFilePath().str()); + Replacement Replacement(OldFilePath, Hunk.getOffset(), Hunk.getLength(), + Hunk.getReplacementText()); + Repl_C2_vec.push_back(Replacement); + }); + + auto C1 = groupReplcementsByFile(Repl_C1); + auto C2 = groupReplcementsByFile(Repl_C2_vec); + for (const auto &[FilePath, Repls_C1] : C1) { + auto ItC2 = C2.find(FilePath); + if (ItC2 == C2.end()) { + // No replacements in C2 for this file, just add C1 + for (const auto &R : Repls_C1) { + Result.push_back({R, false}); // false means not containing manual fix + } + continue; + } + // Merge replacements in C1 and C2 for this file + auto Repls_C2 = ItC2->second; + auto Merged = mergeC1AndC2Impl(Repls_C1, Repls_C2); + Result.insert(Result.end(), Merged.begin(), Merged.end()); + } + + return Result; +} + +// clang-format off +// Repl A +// [CUDA code 1] -----------------------------------------> [CUDA code 2] +// | / | +// | Repl C1 / | Repl B +// | ________/ | +// V / V +// [SYCL code 1] / [SYCL code 2] +// | / | +// | Repl C2 Repl D / | +// | / | +// V shift (may have conlifct) V merge | +// [SYCL code 1.1] ----------------> [SYCL code 1.1] ----------> | +// (based on CUDA code 1) (based on CUDA code 2) | +// V +// [SYCL code 2.1] +// +// Repl_A: Read from gitdiff2yaml generated files. +// Repl_B: Current in-memory migration replacements. +// Repl C1: Read from MainSourceFiles.yaml file. +// Repl_C2: Read from gitdiff2yaml generated files. +// +// Repl_A has 4 parts: +// Repl_A_1: New added files. +// Repl_A_2: Replacements in modified files. +// Repl_A_3: Deleted files. +// Repl_A_4: Replacements in moved files. +// +// clang-format on +// +// Merge process: +// 1. Merge Repl_C1 and Repl_C2 directly, named Repl_C. If there is conlict, +// we keep Repl_C2. +// Repl_C can be divided in to 2 parts: +// Repl_C_x: Replacements which in ranges of Repl_A_3 or delete hunks in +// Repl_A_2/Repl_A_4. +// Repl_C_y: Other replacements. +// Repl_C_x will be ignored during this merge. +// 2. Shfit Repl_C_y with Repl_A, called Repl_D. +// 3. Merge Repl_D and Repl_B. May have conflicts. +std::map> reMigrationMerge( + const GitDiffChanges &Repl_A, const std::vector &Repl_B, + const std::vector &Repl_C1, const GitDiffChanges &Repl_C2, + const std::map + &FileNameMap) { + // TODO: is this assumption true? Will the manual fix delete/add/move file? + assert(Repl_C2.AddFileHunks.empty() && Repl_C2.DeleteFileHunks.empty() && + Repl_C2.MoveFileHunks.empty() && + "Repl_C2 should only have ModifiyFileHunks."); + // Merge Repl_C1 and Repl_C2. If has conflict, keep repl from Repl_C2. + // 1. Repl_C1 has name like file1.cpp, file2.cpp, file3.cu, file4.cuh + // but Repl_C2 has name like file1.cpp, file2.cpp.dp.cpp, file3.dp.cpp, + // file4.dp.hpp. We need convert the filename in Repl_C2 to CUDA style + // 2. Repl_C1 is based on CUDA code, Repl_C2 is based on the migrated SYCL + // code, the result Repl_C should based on the original CUDA code. + std::vector Repl_C = + mergeC1AndC2(Repl_C1, Repl_C2, FileNameMap); + + // Convert vector in Repl_A to map for quick lookup. + std::map> + DeletedParts; + std::map ModifiedParts; + for (const auto &Hunk : Repl_A.ModifyFileHunks) { + if (Hunk.getLength() != 0 && Hunk.getReplacementText().size() == 0) { + DeletedParts[Hunk.getFilePath().str()][Hunk.getOffset()] = + Hunk.getLength(); + } + llvm::cantFail(ModifiedParts[Hunk.getFilePath().str()].add( + Replacement(Hunk.getFilePath().str(), Hunk.getOffset(), + Hunk.getLength(), Hunk.getReplacementText()))); + } + for (const auto &Hunk : Repl_A.MoveFileHunks) { + if (Hunk.getLength() != 0 && Hunk.getReplacementText().size() == 0) { + DeletedParts[Hunk.getFilePath().str()][Hunk.getOffset()] = + Hunk.getLength(); + } + llvm::cantFail(ModifiedParts[Hunk.getFilePath().str()].add( + Replacement(Hunk.getFilePath().str(), Hunk.getOffset(), + Hunk.getLength(), Hunk.getReplacementText()))); + } + for (const auto &Hunk : Repl_A.DeleteFileHunks) { + DeletedParts[Hunk.getOldFilePath()] = std::map(); + } + + // Get Repl_C_y + std::map> Repl_C_y; + + for (const auto &Repl : Repl_C) { + // The gitdiff changes are line-based while clang replacements are + // character-based. So here assume there is no overlap (only repl totally + // covered by delete hunk) between delete hunks and replacements. + const auto &It = DeletedParts.find(Repl.getFilePath().str()); + if (It == DeletedParts.end()) { + Repl_C_y[Repl.getFilePath().str()].push_back(Repl); + continue; + } + + // Check if the replacement is in a deleted part. + bool HasConflict = false; + for (const auto &Part : It->second) { + if (HasConflict = + hasConflict(Repl, Replacement(Repl.getFilePath(), Part.first, + Part.second, ""))) + break; + } + if (!HasConflict) + Repl_C_y[Repl.getFilePath().str()].push_back(Repl); + } + + // Shift Repl_C_y with Repl_A(ModifiedParts) + std::map> Repl_D; + for (const auto &Item : Repl_C_y) { + const auto &FilePath = Item.first; + const auto &Repls = Item.second; + // Check if the file has modified parts. + const auto &It = ModifiedParts.find(FilePath); + if (It == ModifiedParts.end()) { + Repl_D[FilePath] = Repls; + continue; + } + Repl_D[FilePath] = calculateUpdatedRanges(It->second, Repls); + } + + // Group Repl_B by file + const auto Repl_B_by_file = groupReplcementsByFile(Repl_B); + // Merge Repl_D and Repl_B + // Convert the repls to a map then merge line by line + std::map> Result; + for (const auto &Pair : Repl_B_by_file) { + std::map ReplBInLines = + convertReplcementsLineString(Pair.second); + if (Repl_D.find(Pair.first) != Repl_D.end()) { + // Merge Repl_D and Repl_B + std::map> ReplDInLines = + convertReplcementsLineString(Repl_D[Pair.first]); + Result[Pair.first] = + mergeMapsByLine(ReplBInLines, ReplDInLines, Pair.first); + } else { + // No Repl_D for this file, just add Repl_B. + Result[Pair.first] = Pair.second; + } + } + + return Result; +} +} // namespace clang::dpct diff --git a/clang/lib/DPCT/IncMigration/ReMigration.h b/clang/lib/DPCT/IncMigration/ReMigration.h new file mode 100644 index 000000000000..d49b1b5e8167 --- /dev/null +++ b/clang/lib/DPCT/IncMigration/ReMigration.h @@ -0,0 +1,273 @@ +//===----------------------- ReMigration.h --------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef __DPCT_REMIGRATION_H__ +#define __DPCT_REMIGRATION_H__ + +#include "clang/Tooling/Core/Replacement.h" +#include "clang/Tooling/ReplacementsYaml.h" + +extern std::optional< + std::function> + getLineStringHook; +extern std::optional< + std::function> + getLineNumberHook; +extern std::optional< + std::function> + getLineBeginOffsetHook; + +namespace clang { +namespace dpct { +class Hunk { +public: + enum HunkType : unsigned { + ModifyFile = 0, + AddFile, + DeleteFile, + MoveFile, + Unspecified + }; + +private: + HunkType HT = Unspecified; + +public: + Hunk(HunkType HT) : HT(HT) {} + HunkType getHunkType() const { return HT; } + virtual ~Hunk() = default; +}; + +class ModifyFileHunk : public Hunk, public tooling::Replacement { +public: + ModifyFileHunk() : Hunk(ModifyFile), Replacement() {} + ModifyFileHunk(const std::string &FilePath, unsigned Offset, unsigned Length, + const std::string &ReplacementText) + : Hunk(ModifyFile), + Replacement(FilePath, Offset, Length, ReplacementText) {} + ModifyFileHunk(const Replacement &R) : Hunk(ModifyFile), Replacement(R) {} +}; + +class AddFileHunk : public Hunk { + std::string NewFilePath; + +public: + AddFileHunk() : Hunk(AddFile) {} + AddFileHunk(std::string NewFilePath); + const std::string &getNewFilePath() const { return NewFilePath; } +}; + +class DeleteFileHunk : public Hunk { + std::string OldFilePath; + +public: + DeleteFileHunk() : Hunk(DeleteFile) {} + DeleteFileHunk(std::string OldFilePath); + const std::string &getOldFilePath() const { return OldFilePath; } +}; + +class MoveFileHunk : public Hunk, public tooling::Replacement { + std::string NewFilePath; + +public: + MoveFileHunk() : Hunk(MoveFile), Replacement() {} + MoveFileHunk(const std::string &FilePath, unsigned Offset, unsigned Length, + const std::string &ReplacementText, + const std::string &NewFilePath) + : Hunk(MoveFile), Replacement(FilePath, Offset, Length, ReplacementText), + NewFilePath(std::move(NewFilePath)) {} + MoveFileHunk(const Replacement &R, const std::string &NewFilePath) + : Hunk(MoveFile), Replacement(R), NewFilePath(NewFilePath) {} + std::string getNewFilePath() const { return NewFilePath; } +}; + +struct GitDiffChanges { + std::vector ModifyFileHunks; + std::vector AddFileHunks; + std::vector DeleteFileHunks; + std::vector MoveFileHunks; +}; +GitDiffChanges &getUpstreamChanges(); +GitDiffChanges &getUserChanges(); +std::shared_ptr & +getLastMigration(); + +class TaggedReplacement : public clang::tooling::Replacement { +public: + TaggedReplacement() = default; + TaggedReplacement(const clang::tooling::Replacement &R, + bool ContainsManualFix) + : clang::tooling::Replacement(R), ContainsManualFix(ContainsManualFix) {} + + TaggedReplacement(StringRef FilePath, unsigned Offset, unsigned Length, + StringRef ReplacementText, bool ContainsManualFix) + : TaggedReplacement(clang::tooling::Replacement(FilePath, Offset, Length, + ReplacementText), + ContainsManualFix) {} + bool containsManualFix() const { return ContainsManualFix; } + +private: + bool ContainsManualFix = true; +}; +std::vector +calculateUpdatedRanges(const clang::tooling::Replacements &Repls, + const std::vector &NewRepl); +std::map> +groupReplcementsByFile(const std::vector &Repls); +std::vector +splitReplInOrderToNotCrossLines(const std::vector &InRepls); +std::map> +convertReplcementsLineString(const std::vector &Repls); +std::vector +mergeMapsByLine(const std::map &MapA, + const std::map> &MapB, + const clang::tooling::UnifiedPath &FilePath); +std::vector mergeC1AndC2( + const std::vector &Repl_C1, + const GitDiffChanges &Repl_C2, + const std::map + &FileNameMap); + +std::map> +reMigrationMerge(const GitDiffChanges &Repl_A, + const std::vector &Repl_B, + const std::vector &Repl_C1, + const GitDiffChanges &Repl_C2, + const std::map &FileNameMap); +} // namespace dpct +} // namespace clang + +LLVM_YAML_IS_SEQUENCE_VECTOR(clang::dpct::ModifyFileHunk) +LLVM_YAML_IS_SEQUENCE_VECTOR(clang::dpct::AddFileHunk) +LLVM_YAML_IS_SEQUENCE_VECTOR(clang::dpct::DeleteFileHunk) +LLVM_YAML_IS_SEQUENCE_VECTOR(clang::dpct::MoveFileHunk) +LLVM_YAML_DECLARE_ENUM_TRAITS(clang::dpct::Hunk::HunkType) +namespace llvm { +namespace yaml { +inline void ScalarEnumerationTraits::enumeration( + IO &io, clang::dpct::Hunk::HunkType &value) { + io.enumCase(value, "ModifyFile", clang::dpct::Hunk::HunkType::ModifyFile); + io.enumCase(value, "AddFile", clang::dpct::Hunk::HunkType::AddFile); + io.enumCase(value, "DeleteFile", clang::dpct::Hunk::HunkType::DeleteFile); + io.enumCase(value, "MoveFile", clang::dpct::Hunk::HunkType::MoveFile); + io.enumCase(value, "Unspecified", clang::dpct::Hunk::HunkType::Unspecified); +} + +template <> struct MappingTraits { + struct NormalizedModifyFileHunk { + NormalizedModifyFileHunk(const IO &io) : Offset(0), Length(0) {} + NormalizedModifyFileHunk(const IO &io, clang::dpct::ModifyFileHunk &H) + : FilePath(H.getFilePath()), Offset(H.getOffset()), + Length(H.getLength()), ReplacementText(H.getReplacementText()) {} + + clang::dpct::ModifyFileHunk denormalize(const IO &) { + clang::dpct::ModifyFileHunk H(FilePath, Offset, Length, ReplacementText); + return H; + } + + std::string FilePath; + unsigned int Offset; + unsigned int Length; + std::string ReplacementText; + }; + + static void mapping(IO &Io, clang::dpct::ModifyFileHunk &H) { + MappingNormalization + Keys(Io, H); + Io.mapRequired("FilePath", Keys->FilePath); + Io.mapRequired("Offset", Keys->Offset); + Io.mapRequired("Length", Keys->Length); + Io.mapRequired("ReplacementText", Keys->ReplacementText); + } +}; + +template <> struct MappingTraits { + struct NormalizedAddFileHunk { + NormalizedAddFileHunk(const IO &io) {} + NormalizedAddFileHunk(const IO &io, clang::dpct::AddFileHunk &H) + : NewFilePath(H.getNewFilePath()) {} + + clang::dpct::AddFileHunk denormalize(const IO &io) { + clang::dpct::AddFileHunk H(NewFilePath); + return H; + } + + std::string NewFilePath; + }; + static void mapping(IO &Io, clang::dpct::AddFileHunk &H) { + MappingNormalization Keys( + Io, H); + Io.mapRequired("NewFilePath", Keys->NewFilePath); + } +}; + +template <> struct MappingTraits { + struct NormalizedDeleteFileHunk { + NormalizedDeleteFileHunk(const IO &io) {} + NormalizedDeleteFileHunk(const IO &io, clang::dpct::DeleteFileHunk &H) + : OldFilePath(H.getOldFilePath()) {} + + clang::dpct::DeleteFileHunk denormalize(const IO &io) { + clang::dpct::DeleteFileHunk H(OldFilePath); + return H; + } + + std::string OldFilePath; + }; + static void mapping(IO &Io, clang::dpct::DeleteFileHunk &H) { + MappingNormalization + Keys(Io, H); + Io.mapRequired("OldFilePath", Keys->OldFilePath); + } +}; + +template <> struct MappingTraits { + struct NormalizedMoveFileHunk { + NormalizedMoveFileHunk(const IO &io) : Offset(0), Length(0) {} + NormalizedMoveFileHunk(const IO &io, clang::dpct::MoveFileHunk &H) + : FilePath(H.getFilePath()), Offset(H.getOffset()), + Length(H.getLength()), ReplacementText(H.getReplacementText()), + NewFilePath(H.getNewFilePath()) {} + + clang::dpct::MoveFileHunk denormalize(const IO &io) { + clang::dpct::MoveFileHunk H(FilePath, Offset, Length, ReplacementText, + NewFilePath); + return H; + } + + std::string FilePath; + unsigned int Offset; + unsigned int Length; + std::string ReplacementText; + std::string NewFilePath; + }; + + static void mapping(IO &Io, clang::dpct::MoveFileHunk &H) { + MappingNormalization + Keys(Io, H); + Io.mapRequired("FilePath", Keys->FilePath); + Io.mapRequired("Offset", Keys->Offset); + Io.mapRequired("Length", Keys->Length); + Io.mapRequired("ReplacementText", Keys->ReplacementText); + Io.mapRequired("NewFilePath", Keys->NewFilePath); + } +}; + +template <> struct MappingTraits { + static void mapping(IO &Io, clang::dpct::GitDiffChanges &GDC) { + Io.mapOptional("ModifyFileHunks", GDC.ModifyFileHunks); + Io.mapOptional("AddFileHunks", GDC.AddFileHunks); + Io.mapOptional("DeleteFileHunks", GDC.DeleteFileHunks); + Io.mapOptional("MoveFileHunks", GDC.MoveFileHunks); + } +}; +} // namespace yaml +} // namespace llvm + +#endif // __DPCT_REMIGRATION_H__ diff --git a/clang/lib/DPCT/UserDefinedRules/PatternRewriter.cpp b/clang/lib/DPCT/UserDefinedRules/PatternRewriter.cpp index 911c8d86759a..4dea8b6799cf 100644 --- a/clang/lib/DPCT/UserDefinedRules/PatternRewriter.cpp +++ b/clang/lib/DPCT/UserDefinedRules/PatternRewriter.cpp @@ -25,6 +25,8 @@ #include #include +std::set MainSrcFilesHasCudaSyntex; + namespace clang { namespace dpct { @@ -464,7 +466,6 @@ static void applyExtenstionNameChange( std::string SrcFile = Input.substr(Pos, Next + ExtensionType.length() + 1 /*strlen of "."*/ - Pos); bool HasCudaSyntax = false; - std::set MainSrcFilesHasCudaSyntex; for (auto &Entry : DpctGlobalInfo::getMainSourceYamlTUR()->MainSourceFilesDigest) { if (Entry.HasCUDASyntax) diff --git a/clang/test/CMakeLists.txt b/clang/test/CMakeLists.txt index 20378b2b5e76..5752a6288757 100644 --- a/clang/test/CMakeLists.txt +++ b/clang/test/CMakeLists.txt @@ -95,6 +95,7 @@ list(APPEND CLANG_TEST_DEPS dpct-codepin-report-files scan-build-py pattern-rewriter-binary + gitdiff2yaml-binary ) if(CLANG_ENABLE_CIR) diff --git a/clang/test/dpct/autocomplete.c b/clang/test/dpct/autocomplete.c index 9a80e0d3401a..f2f86238a865 100644 --- a/clang/test/dpct/autocomplete.c +++ b/clang/test/dpct/autocomplete.c @@ -45,6 +45,7 @@ // DASH-NEXT: --output-verbosity= // DASH-NEXT: --process-all // DASH-NEXT: --query-api-mapping +// DASH-NEXT: --remigration // DASH-NEXT: --report-file-prefix // DASH-NEXT: --report-format= // DASH-NEXT: --report-only diff --git a/clang/test/dpct/help_option_check/lin/help_all.txt b/clang/test/dpct/help_option_check/lin/help_all.txt index 22186a896daa..c07e2662e0aa 100644 --- a/clang/test/dpct/help_option_check/lin/help_all.txt +++ b/clang/test/dpct/help_option_check/lin/help_all.txt @@ -106,6 +106,7 @@ All DPCT options to the --out-root directory. The --in-root option should be explicitly specified. Default: off. --query-api-mapping= - Query functionally compatible SYCL API to migrate CUDA API. + --remigration - Enable remigration. Default: off. --report-file-prefix= - Specify the prefix for the migration report file names. The full file name will have a suffix derived from the report-type, and an extension derived from the report-format. For example: .apis.csv or .stats.log. If this option is not diff --git a/clang/test/dpct/remigration/LastMigration.yaml b/clang/test/dpct/remigration/LastMigration.yaml new file mode 100644 index 000000000000..b06b3c4fb633 --- /dev/null +++ b/clang/test/dpct/remigration/LastMigration.yaml @@ -0,0 +1,369 @@ +--- +MainSourceFile: 'PATH_PLACEHOLDER/MainSrcFiles_placehold' +Replacements: + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 0 + Length: 26 + ReplacementText: "#include \n#include \n" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 45 + Length: 0 + ReplacementText: "/*\nDPCT1009:3: SYCL reports errors using exceptions and does not use error codes. Please replace the \"get_error_string_dummy(...)\" with a real error-handling function.\n*/\n" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 211 + Length: 11 + ReplacementText: 'dpct::err0' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 292 + Length: 325 + ReplacementText: '' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 708 + Length: 11 + ReplacementText: '' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 782 + Length: 0 + ReplacementText: "auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();\n " + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 792 + Length: 10 + ReplacementText: 'item_ct1.get_group(2)' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 805 + Length: 10 + ReplacementText: 'item_ct1.get_local_range(2)' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 818 + Length: 11 + ReplacementText: 'item_ct1.get_local_id(2)' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 931 + Length: 0 + ReplacementText: ' try ' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 932 + Length: 0 + ReplacementText: "\n dpct::device_ext &dev_ct1 = dpct::get_current_device();\n sycl::queue &q_ct1 = dev_ct1.in_order_queue();" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 975 + Length: 42 + ReplacementText: 'DPCT_CHECK_ERROR(input_d = sycl::malloc_device(size, q_ct1))' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1033 + Length: 43 + ReplacementText: 'DPCT_CHECK_ERROR(output_d = sycl::malloc_device(size, q_ct1))' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1099 + Length: 71 + ReplacementText: 'DPCT_CHECK_ERROR(q_ct1.memcpy(input_d, data, size * sizeof(float)).wait())' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1251 + Length: 0 + ReplacementText: " /*\n DPCT1049:0: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed.\n */\n" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1253 + Length: 64 + ReplacementText: "q_ct1.parallel_for(\n sycl::nd_range<3>(sycl::range<3>(1, 1, numBlocks) * sycl::range<3>(1, 1, blockSize), sycl::range<3>(1, 1, blockSize)), \n [=](sycl::nd_item<3> item_ct1) {\n scale_kernel(input_d, scale, output_d);\n });" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: true + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1317 + Length: 1 + ReplacementText: '' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1332 + Length: 23 + ReplacementText: 'DPCT_CHECK_ERROR(dev_ct1.queues_wait_and_throw())' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1358 + Length: 0 + ReplacementText: " /*\n DPCT1010:4: SYCL uses exceptions to report errors and does not use the error codes. The cudaGetLastError function call was replaced with 0. You need to rewrite this code.\n */\n" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1360 + Length: 11 + ReplacementText: 'dpct::err0' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1378 + Length: 18 + ReplacementText: '0' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1398 + Length: 0 + ReplacementText: " /*\n DPCT1000:2: Error handling if-stmt was detected but could not be rewritten.\n */\n" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1411 + Length: 11 + ReplacementText: '0' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1426 + Length: 0 + ReplacementText: " /*\n DPCT1009:5: SYCL reports errors using exceptions and does not use error codes. Please replace the \"get_error_string_dummy(...)\" with a real error-handling function.\n */\n" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1481 + Length: 23 + ReplacementText: 'dpct::get_error_string_dummy(err)' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1507 + Length: 0 + ReplacementText: " /*\n DPCT1001:1: The statement could not be removed.\n */\n" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1511 + Length: 17 + ReplacementText: 'dpct::dpct_free(input_d, q_ct1)' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1534 + Length: 18 + ReplacementText: 'dpct::dpct_free(output_d, q_ct1)' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1596 + Length: 98 + ReplacementText: 'DPCT_CHECK_ERROR(q_ct1.memcpy(output, output_d, size * sizeof(float)).wait())' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1710 + Length: 17 + ReplacementText: 'DPCT_CHECK_ERROR(dpct::dpct_free(input_d, q_ct1))' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1743 + Length: 18 + ReplacementText: 'DPCT_CHECK_ERROR(dpct::dpct_free(output_d, q_ct1))' + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false + - FilePath: 'PATH_PLACEHOLDER/test.cu' + Offset: 1765 + Length: 0 + ReplacementText: "\ncatch (sycl::exception const &exc) {\n std::cerr << exc.what() << \"Exception caught at file:\" << __FILE__ << \", line:\" << __LINE__ << std::endl;\n std::exit(1);\n}" + ConstantFlag: '' + ConstantOffset: 0 + InitStr: '' + NewHostVarName: '' + BlockLevelFormatFlag: false +MainSourceFilesDigest: + - MainSourceFile: 'PATH_PLACEHOLDER/test.cu' + Digest: 63232f272cda678b14a86a4243a28f63 + HasCUDASyntax: true +DpctVersion: 21.0.0 +SDKVersionMajor: '12' +SDKVersionMinor: '8' +MainHelperFileName: '' +USMLevel: '' +FeatureMap: {} +CompileTargets: {} +OptionMap: + AnalysisScopePath: + Value: '' + ValueVec: + - 'PATH_PLACEHOLDER' + Specified: false + AsyncHandler: + Value: 'false' + Specified: false + BuildScript: + Value: '0' + Specified: false + CodePinEnabled: + Value: 'false' + Specified: false + CommentsEnabled: + Value: 'false' + Specified: false + CompilationsDir: + Value: '' + Specified: false + CtadEnabled: + Value: 'false' + Specified: false + EnablepProfiling: + Value: 'false' + Specified: false + ExperimentalFlag: + Value: '0' + Specified: false + ExplicitNamespace: + Value: '20' + Specified: false + ExtensionDDFlag: + Value: '0' + Specified: false + ExtensionDEFlag: + Value: '4294967295' + Specified: false + HelperFuncPreferenceFlag: + Value: '0' + Specified: false + NDRangeDim: + Value: '3' + Specified: false + NoDRYPattern: + Value: 'false' + Specified: false + OptimizeMigration: + Value: 'false' + Specified: false + ProcessAll: + Value: 'false' + Specified: false + RuleFile: + Value: '' + Specified: false + SyclNamedLambda: + Value: 'false' + Specified: false + UseSYCLCompat: + Value: 'false' + Specified: false + UsmLevel: + Value: '1' + Specified: false +... diff --git a/clang/test/dpct/remigration/UpstreamChanges.yaml b/clang/test/dpct/remigration/UpstreamChanges.yaml new file mode 100644 index 000000000000..842c6b22accb --- /dev/null +++ b/clang/test/dpct/remigration/UpstreamChanges.yaml @@ -0,0 +1,19 @@ +--- +ModifyFileHunks: + - FilePath: test.cu + Offset: 26 + Length: 0 + ReplacementText: "#include \"cublas_v2.h\"\n" + - FilePath: test.cu + Offset: 708 + Length: 0 + ReplacementText: "#define CHECK_CUBLAS(call) \\\n do { \\\n cublasStatus_t err = call; \\\n if (err != CUBLAS_STATUS_SUCCESS) { \\\n fprintf(stderr, \"CUBLAS error in %s at line %d: %d\\n\", __FILE__, \\\n __LINE__, err); \\\n exit(EXIT_FAILURE); \\\n } \\\n } while (0)\n\n" + - FilePath: test.cu + Offset: 869 + Length: 0 + ReplacementText: "void scale_cublas(float *data, float scale, float *output, int size) {\n cublasHandle_t handle;\n CHECK_CUBLAS(cublasCreate(&handle));\n float alpha = scale;\n CHECK_CUBLAS(cublasSscal(handle, size, &alpha, data, 1));\n CHECK_CUDA(cudaDeviceSynchronize());\n CHECK_CUBLAS(cublasDestroy(handle));\n}\n\ntemplate\n" + - FilePath: test.cu + Offset: 1174 + Length: 404 + ReplacementText: " if (use_cublas) {\n scale_cublas(input_d, scale, output_d, size);\n } else {\n int blockSize = 256;\n int numBlocks = (size + blockSize - 1) / blockSize;\n scale_kernel<<>>(input_d, scale, output_d);\n CHECK_CUDA(cudaDeviceSynchronize());\n cudaError_t err = cudaGetLastError();\n if (err != cudaSuccess) {\n fprintf(stderr, \"CUDA kernel launch failed: %s\\n\",\n cudaGetErrorString(err));\n cudaFree(input_d);\n cudaFree(output_d);\n exit(EXIT_FAILURE);\n }\n" +... diff --git a/clang/test/dpct/remigration/UserChanges.yaml b/clang/test/dpct/remigration/UserChanges.yaml new file mode 100644 index 000000000000..b5adf495d2f6 --- /dev/null +++ b/clang/test/dpct/remigration/UserChanges.yaml @@ -0,0 +1,27 @@ +--- +ModifyFileHunks: + - FilePath: test.dp.cpp + Offset: 69 + Length: 508 + ReplacementText: '' + - FilePath: test.dp.cpp + Offset: 709 + Length: 92 + ReplacementText: " int idx = item_ct1.get_group(2) * item_ct1.get_local_range(2) +\n item_ct1.get_local_id(2);\n" + - FilePath: test.dp.cpp + Offset: 839 + Length: 69 + ReplacementText: "void scale(float *data, float scale, float *output, int size) try {\n" + - FilePath: test.dp.cpp + Offset: 1044 + Length: 264 + ReplacementText: " input_d = sycl::malloc_device(size, q_ct1);\n output_d = sycl::malloc_device(size, q_ct1);\n q_ct1.memcpy(input_d, data, size * sizeof(float)).wait();\n" + - FilePath: test.dp.cpp + Offset: 1386 + Length: 1252 + ReplacementText: " q_ct1.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, numBlocks) *\n sycl::range<3>(1, 1, blockSize),\n sycl::range<3>(1, 1, blockSize)),\n [=](sycl::nd_item<3> item_ct1) {\n scale_kernel(input_d, scale, output_d);\n });\n dev_ct1.queues_wait_and_throw();\n" + - FilePath: test.dp.cpp + Offset: 2639 + Length: 371 + ReplacementText: " q_ct1.memcpy(output, output_d, size * sizeof(float)).wait();\n dpct::dpct_free(input_d, q_ct1);\n dpct::dpct_free(output_d, q_ct1);\n} catch (sycl::exception const &exc) {\n std::cerr << exc.what() << \"Exception caught at file:\" << __FILE__\n << \", line:\" << __LINE__ << std::endl;\n" +... diff --git a/clang/test/dpct/remigration/expect.txt b/clang/test/dpct/remigration/expect.txt new file mode 100644 index 000000000000..049374105f14 --- /dev/null +++ b/clang/test/dpct/remigration/expect.txt @@ -0,0 +1,127 @@ +#include +#include +#include +#include + +<<<<<<< +/* +DPCT1009:1: SYCL reports errors using exceptions and does not use error codes. Please replace the "get_error_string_dummy(...)" with a real error-handling function. +*/ +#define CHECK_CUDA(call) \ +======= +void scale_kernel(float *data, float scale, float *output) { +>>>>>>> +<<<<<<< + dpct::err0 err = call; \ + \ +======= +>>>>>>> +<<<<<<< + int err = call; \ + if (err != 0) { \ +======= +>>>>>>> +<<<<<<< +void scale_kernel(float *data, float scale, float *output) { + auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>(); + int idx = item_ct1.get_group(2) * item_ct1.get_local_range(2) + item_ct1.get_local_id(2); +======= + auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>(); + int idx = item_ct1.get_group(2) * item_ct1.get_local_range(2) + + item_ct1.get_local_id(2); +>>>>>>> + output[idx] = data[idx] * scale; +} + +void scale_cublas(float *data, float scale, float *output, int size) try { + dpct::blas::descriptor_ptr handle; + CHECK_CUBLAS(DPCT_CHECK_ERROR(handle = new dpct::blas::descriptor())); + float alpha = scale; + CHECK_CUBLAS(DPCT_CHECK_ERROR(oneapi::mkl::blas::column_major::scal(handle->get_queue(), size, alpha, data, 1))); + CHECK_CUDA(DPCT_CHECK_ERROR(dpct::get_current_device().queues_wait_and_throw())); + CHECK_CUBLAS(DPCT_CHECK_ERROR(delete (handle))); +} +catch (sycl::exception const &exc) { + std::cerr << exc.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl; + std::exit(1); +} + +template +<<<<<<< +void scale(float *data, float scale, float *output, int size) try { + dpct::device_ext &dev_ct1 = dpct::get_current_device(); + sycl::queue &q_ct1 = dev_ct1.in_order_queue(); +======= +void scale(float *data, float scale, float *output, int size) try { + dpct::device_ext &dev_ct1 = dpct::get_current_device(); + sycl::queue &q_ct1 = dev_ct1.in_order_queue(); +>>>>>>> + float *input_d, *output_d; +<<<<<<< + CHECK_CUDA(DPCT_CHECK_ERROR(input_d = sycl::malloc_device(size, q_ct1))); + CHECK_CUDA(DPCT_CHECK_ERROR(output_d = sycl::malloc_device(size, q_ct1))); +======= + input_d = sycl::malloc_device(size, q_ct1); + output_d = sycl::malloc_device(size, q_ct1); + q_ct1.memcpy(input_d, data, size * sizeof(float)).wait(); +>>>>>>> +<<<<<<< + DPCT_CHECK_ERROR(q_ct1.memcpy(input_d, data, size * sizeof(float)).wait())); +======= +>>>>>>> + + if (use_cublas) { + scale_cublas(input_d, scale, output_d, size); + } else { + int blockSize = 256; + int numBlocks = (size + blockSize - 1) / blockSize; + /* + DPCT1049:0: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed. + */ + q_ct1.parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, numBlocks) * sycl::range<3>(1, 1, blockSize), sycl::range<3>(1, 1, blockSize)), + [=](sycl::nd_item<3> item_ct1) { + scale_kernel(input_d, scale, output_d); + }); + CHECK_CUDA(DPCT_CHECK_ERROR(dev_ct1.queues_wait_and_throw())); + /* + DPCT1010:4: SYCL uses exceptions to report errors and does not use the error codes. The cudaGetLastError function call was replaced with 0. You need to rewrite this code. + */ + dpct::err0 err = 0; + /* + DPCT1000:3: Error handling if-stmt was detected but could not be rewritten. + */ + if (err != 0) { + fprintf(stderr, "CUDA kernel launch failed: %s\n", + /* + DPCT1009:5: SYCL reports errors using exceptions and does not use error codes. Please replace the "get_error_string_dummy(...)" with a real error-handling function. + */ + dpct::get_error_string_dummy(err)); + /* + DPCT1001:2: The statement could not be removed. + */ + dpct::dpct_free(input_d, q_ct1); + dpct::dpct_free(output_d, q_ct1); + exit(EXIT_FAILURE); + } + } + +<<<<<<< + CHECK_CUDA(DPCT_CHECK_ERROR(q_ct1.memcpy(output, output_d, size * sizeof(float)).wait())); + CHECK_CUDA(DPCT_CHECK_ERROR(dpct::dpct_free(input_d, q_ct1))); + CHECK_CUDA(DPCT_CHECK_ERROR(dpct::dpct_free(output_d, q_ct1))); +} +catch (sycl::exception const &exc) { + std::cerr << exc.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl; + std::exit(1); +} +======= + q_ct1.memcpy(output, output_d, size * sizeof(float)).wait(); + dpct::dpct_free(input_d, q_ct1); + dpct::dpct_free(output_d, q_ct1); +} catch (sycl::exception const &exc) { + std::cerr << exc.what() << "Exception caught at file:" << __FILE__ + << ", line:" << __LINE__ << std::endl; + std::exit(1); +} +>>>>>>> diff --git a/clang/test/dpct/remigration/src.txt b/clang/test/dpct/remigration/src.txt new file mode 100644 index 000000000000..d8e0595dd594 --- /dev/null +++ b/clang/test/dpct/remigration/src.txt @@ -0,0 +1,68 @@ +#include "cuda_runtime.h" +#include "cublas_v2.h" +#include + +#define CHECK_CUDA(call) \ + do { \ + cudaError_t err = call; \ + if (err != cudaSuccess) { \ + fprintf(stderr, "CUDA error in %s at line %d: %s\n", __FILE__, __LINE__, \ + cudaGetErrorString(err)); \ + exit(EXIT_FAILURE); \ + } \ + } while (0) + +#define CHECK_CUBLAS(call) \ + do { \ + cublasStatus_t err = call; \ + if (err != CUBLAS_STATUS_SUCCESS) { \ + fprintf(stderr, "CUBLAS error in %s at line %d: %d\n", __FILE__, \ + __LINE__, err); \ + exit(EXIT_FAILURE); \ + } \ + } while (0) + +__global__ void scale_kernel(float *data, float scale, float *output) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + output[idx] = data[idx] * scale; +} + +void scale_cublas(float *data, float scale, float *output, int size) { + cublasHandle_t handle; + CHECK_CUBLAS(cublasCreate(&handle)); + float alpha = scale; + CHECK_CUBLAS(cublasSscal(handle, size, &alpha, data, 1)); + CHECK_CUDA(cudaDeviceSynchronize()); + CHECK_CUBLAS(cublasDestroy(handle)); +} + +template +void scale(float *data, float scale, float *output, int size) { + float *input_d, *output_d; + CHECK_CUDA(cudaMalloc(&input_d, size * sizeof(float))); + CHECK_CUDA(cudaMalloc(&output_d, size * sizeof(float))); + CHECK_CUDA( + cudaMemcpy(input_d, data, size * sizeof(float), cudaMemcpyHostToDevice)); + + if (use_cublas) { + scale_cublas(input_d, scale, output_d, size); + } else { + int blockSize = 256; + int numBlocks = (size + blockSize - 1) / blockSize; + scale_kernel<<>>(input_d, scale, output_d); + CHECK_CUDA(cudaDeviceSynchronize()); + cudaError_t err = cudaGetLastError(); + if (err != cudaSuccess) { + fprintf(stderr, "CUDA kernel launch failed: %s\n", + cudaGetErrorString(err)); + cudaFree(input_d); + cudaFree(output_d); + exit(EXIT_FAILURE); + } + } + + CHECK_CUDA(cudaMemcpy(output, output_d, size * sizeof(float), + cudaMemcpyDeviceToHost)); + CHECK_CUDA(cudaFree(input_d)); + CHECK_CUDA(cudaFree(output_d)); +} diff --git a/clang/test/dpct/remigration/test.cpp b/clang/test/dpct/remigration/test.cpp new file mode 100644 index 000000000000..772ce93bdab1 --- /dev/null +++ b/clang/test/dpct/remigration/test.cpp @@ -0,0 +1,12 @@ +// UNSUPPORTED: system-windows + +// RUN: cd %T +// RUN: mkdir out +// RUN: cp %S/LastMigration.yaml out/MainSourceFiles.yaml +// RUN: sed -i "s|PATH_PLACEHOLDER|$(pwd)|g" out/MainSourceFiles.yaml +// RUN: cp %S/UpstreamChanges.yaml . +// RUN: cp %S/UserChanges.yaml . +// RUN: cp %S/src.txt test.cu + +// RUN: dpct --out-root out ./test.cu --format-range=none --remigration +// RUN: diff %S/expect.txt out/test.dp.cpp diff --git a/clang/tools/CMakeLists.txt b/clang/tools/CMakeLists.txt index 9e4e4fcd8257..54c5619d1f00 100644 --- a/clang/tools/CMakeLists.txt +++ b/clang/tools/CMakeLists.txt @@ -58,6 +58,7 @@ add_clang_subdirectory(libclang) add_clang_subdirectory(dpct) add_clang_subdirectory(pattern-rewriter) +add_clang_subdirectory(gitdiff2yaml) add_clang_subdirectory(scan-build-py) add_clang_subdirectory(amdgpu-arch) add_clang_subdirectory(nvptx-arch) diff --git a/clang/tools/gitdiff2yaml/CMakeLists.txt b/clang/tools/gitdiff2yaml/CMakeLists.txt new file mode 100644 index 000000000000..2509b41bb3ca --- /dev/null +++ b/clang/tools/gitdiff2yaml/CMakeLists.txt @@ -0,0 +1,14 @@ +set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} ") +set(LLVM_LINK_COMPONENTS + support +) + +add_clang_tool(gitdiff2yaml-binary + src/main.cpp + src/gitdiff2yaml.cpp +) +set_target_properties(gitdiff2yaml-binary PROPERTIES OUTPUT_NAME "gitdiff2yaml") +add_llvm_install_targets("install-dpct-gitdiff2yaml" + DEPENDS gitdiff2yaml-binary + COMPONENT gitdiff2yaml-binary +) diff --git a/clang/tools/gitdiff2yaml/src/gitdiff2yaml.cpp b/clang/tools/gitdiff2yaml/src/gitdiff2yaml.cpp new file mode 100644 index 000000000000..b7fe1e65ca6f --- /dev/null +++ b/clang/tools/gitdiff2yaml/src/gitdiff2yaml.cpp @@ -0,0 +1,435 @@ +//===--- gitdiff2yaml.cpp -------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "gitdiff2yaml.h" + +#include "llvm/Support/YAMLTraits.h" +#include "llvm/Support/raw_os_ostream.h" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { +bool startsWith(const std::string &Str, const std::string &Prefix) { + return Str.size() >= Prefix.size() && + Str.compare(0, Prefix.size(), Prefix) == 0; +} + +struct HunkContext { + unsigned OldCurrentLine = 0; + bool InHunk = false; + bool FastForward = false; + std::string CurrentNewFilePath; + std::string CurrentOldFilePath; +}; + +bool parseHunkHeader(const std::string &Line, HunkContext &HC) { + const std::string HunkHeaderPrefix = "@@ -"; + if (!startsWith(Line, HunkHeaderPrefix)) + return false; + + // E.g., + // @@ -0,0 +1,3 @@ + // ^ + // |-- OldEnd + // E.g., + // @@ -24 +24 @@ + // ^ + // |-- OldEnd + size_t OldEnd = Line.find(' ', HunkHeaderPrefix.size()); + std::string OldPart = + Line.substr(HunkHeaderPrefix.size(), OldEnd - HunkHeaderPrefix.size()); + size_t Comma = OldPart.find(','); + if (Comma == std::string::npos) { + HC.OldCurrentLine = std::stoi(OldPart); + } else { + HC.OldCurrentLine = std::stoi(OldPart.substr(0, Comma)); + } + return true; +} + +// vec[0] is -1, which is a placeholder. +// vec[i] /*i is the line number (1-based)*/ is the offset at the line +// beginning. +// Assumption: the line ending in the file is '\n'. +std::vector calculateOldOffset(const std::string &OldFileContent) { + std::vector Ret; + Ret.push_back(-1); // Placeholder for Ret[0]. + std::istringstream OldFileStream(OldFileContent); + std::string Line; + unsigned Offset = 0; + + while (std::getline(OldFileStream, Line)) { + Ret.push_back(Offset); + Offset += Line.size() + + 1; // std::getline does not include the newline character. And we + // assume the line ending is '\n'. So add 1 here. + } + + return Ret; +} + +struct ModifyHunk { + ModifyHunk() = default; + ModifyHunk(std::string FilePath, unsigned Offset, unsigned Length, + std::string ReplacementText) + : FilePath(FilePath), Offset(Offset), Length(Length), + ReplacementText(ReplacementText) {}; + std::string FilePath; + unsigned Offset = 0; + unsigned Length = 0; + std::string ReplacementText; +}; + +struct AddHunk { + AddHunk() = default; + AddHunk(std::string NewFilePath) : NewFilePath(NewFilePath) {}; + std::string NewFilePath; +}; + +struct DeleteHunk { + DeleteHunk() = default; + DeleteHunk(std::string OldFilePath) : OldFilePath(OldFilePath) {}; + std::string OldFilePath; +}; + +struct MoveHunk { + MoveHunk() = default; + MoveHunk(std::string FilePath, unsigned Offset, unsigned Length, + std::string ReplacementText, std::string NewFilePath) + : FilePath(FilePath), Offset(Offset), Length(Length), + ReplacementText(ReplacementText), NewFilePath(NewFilePath) {}; + std::string FilePath; + unsigned Offset = 0; + unsigned Length = 0; + std::string ReplacementText; + std::string NewFilePath; +}; + +struct GitDiffChanges { + std::vector ModifyHunks; + std::vector AddHunks; + std::vector DeleteHunks; + std::vector MoveHunks; +}; +} // namespace + +LLVM_YAML_IS_SEQUENCE_VECTOR(ModifyHunk) +LLVM_YAML_IS_SEQUENCE_VECTOR(AddHunk) +LLVM_YAML_IS_SEQUENCE_VECTOR(DeleteHunk) +LLVM_YAML_IS_SEQUENCE_VECTOR(MoveHunk) +namespace llvm { +namespace yaml { +template <> struct MappingTraits { + struct NormalizedModifyFileHunk { + NormalizedModifyFileHunk(const IO &io) : Offset(0), Length(0) {} + NormalizedModifyFileHunk(const IO &io, ModifyHunk &H) + : FilePath(H.FilePath), Offset(H.Offset), Length(H.Length), + ReplacementText(H.ReplacementText) {} + ModifyHunk denormalize(const IO &) { + ModifyHunk H(FilePath, Offset, Length, ReplacementText); + return H; + } + std::string FilePath; + unsigned int Offset; + unsigned int Length; + std::string ReplacementText; + }; + + static void mapping(IO &Io, ModifyHunk &H) { + MappingNormalization Keys(Io, H); + Io.mapRequired("FilePath", Keys->FilePath); + Io.mapRequired("Offset", Keys->Offset); + Io.mapRequired("Length", Keys->Length); + Io.mapRequired("ReplacementText", Keys->ReplacementText); + } +}; + +template <> struct MappingTraits { + struct NormalizedAddFileHunk { + NormalizedAddFileHunk(const IO &io) {} + NormalizedAddFileHunk(const IO &io, AddHunk &H) + : NewFilePath(H.NewFilePath) {} + AddHunk denormalize(const IO &io) { + AddHunk H(NewFilePath); + return H; + } + std::string NewFilePath; + }; + static void mapping(IO &Io, AddHunk &H) { + MappingNormalization Keys(Io, H); + Io.mapRequired("NewFilePath", Keys->NewFilePath); + } +}; + +template <> struct MappingTraits { + struct NormalizedDeleteFileHunk { + NormalizedDeleteFileHunk(const IO &io) {} + NormalizedDeleteFileHunk(const IO &io, DeleteHunk &H) + : OldFilePath(H.OldFilePath) {} + DeleteHunk denormalize(const IO &io) { + DeleteHunk H(OldFilePath); + return H; + } + std::string OldFilePath; + }; + static void mapping(IO &Io, DeleteHunk &H) { + MappingNormalization Keys(Io, H); + Io.mapRequired("OldFilePath", Keys->OldFilePath); + } +}; + +template <> struct MappingTraits { + struct NormalizedMoveFileHunk { + NormalizedMoveFileHunk(const IO &io) : Offset(0), Length(0) {} + NormalizedMoveFileHunk(const IO &io, MoveHunk &H) + : FilePath(H.FilePath), Offset(H.Offset), Length(H.Length), + ReplacementText(H.ReplacementText), NewFilePath(H.NewFilePath) {} + MoveHunk denormalize(const IO &io) { + MoveHunk H(FilePath, Offset, Length, ReplacementText, NewFilePath); + return H; + } + std::string FilePath; + unsigned int Offset; + unsigned int Length; + std::string ReplacementText; + std::string NewFilePath; + }; + static void mapping(IO &Io, MoveHunk &H) { + MappingNormalization Keys(Io, H); + Io.mapRequired("FilePath", Keys->FilePath); + Io.mapRequired("Offset", Keys->Offset); + Io.mapRequired("Length", Keys->Length); + Io.mapRequired("ReplacementText", Keys->ReplacementText); + Io.mapRequired("NewFilePath", Keys->NewFilePath); + } +}; + +template <> struct MappingTraits { + static void mapping(IO &Io, GitDiffChanges &GDC) { + Io.mapOptional("ModifyFileHunks", GDC.ModifyHunks); + Io.mapOptional("AddFileHunks", GDC.AddHunks); + Io.mapOptional("DeleteFileHunks", GDC.DeleteHunks); + Io.mapOptional("MoveFileHunks", GDC.MoveHunks); + } +}; +} // namespace yaml +} // namespace llvm + +void printYaml(std::ostream &stream, const std::vector &Repls) { + GitDiffChanges GDC; + + for (const auto &R : Repls) { + if (R.OldFilePath == "/dev/null" && R.NewFilePath != "/dev/null") { + // Add replacement + AddHunk AH(R.NewFilePath); + GDC.AddHunks.push_back(AH); + continue; + } + if (R.OldFilePath != "/dev/null" && R.NewFilePath == "/dev/null") { + // Delete replacement + DeleteHunk DH(R.OldFilePath); + GDC.DeleteHunks.push_back(DH); + continue; + } + if (R.OldFilePath == R.NewFilePath && R.OldFilePath != "/dev/null") { + // Modify replacement + ModifyHunk MH(R.OldFilePath, R.Offset, R.Length, R.ReplacementText); + GDC.ModifyHunks.push_back(MH); + continue; + } + if (R.OldFilePath != R.NewFilePath) { + // Move replacement + MoveHunk MH(R.OldFilePath, R.Offset, R.Length, R.ReplacementText, + R.NewFilePath); + GDC.MoveHunks.push_back(MH); + continue; + } + throw std::runtime_error("Invalid replacement: " + R.OldFilePath + " -> " + + R.NewFilePath); + } + + llvm::raw_os_ostream raw_os(stream); + llvm::yaml::Output yout(raw_os); + yout << GDC; +} + +std::vector parseDiff(const std::string &diffOutput, + const std::string &RepoRoot) { + std::vector replacements; + std::istringstream iss(diffOutput); + std::string line; + + HunkContext HC; + std::vector CurrentOldFileOffset; + + std::optional< + std::pair> + DeleteInfo; + std::optional> + AddInfo; + + auto addRepl = [&]() { + Replacement R; + if (DeleteInfo.has_value() && AddInfo.has_value()) { + // replace-replacement + R.OldFilePath = HC.CurrentOldFilePath; + R.NewFilePath = HC.CurrentNewFilePath; + R.Length = DeleteInfo->second; + R.Offset = CurrentOldFileOffset[DeleteInfo->first]; + R.ReplacementText = AddInfo->second; + DeleteInfo.reset(); + AddInfo.reset(); + } else if (DeleteInfo.has_value()) { + // delete-replacement + R.OldFilePath = HC.CurrentOldFilePath; + R.NewFilePath = HC.CurrentNewFilePath; + R.Length = DeleteInfo->second; + R.Offset = CurrentOldFileOffset[DeleteInfo->first]; + R.ReplacementText = ""; + DeleteInfo.reset(); + } else if (AddInfo.has_value()) { + // insert-replacement + R.OldFilePath = HC.CurrentOldFilePath; + R.NewFilePath = HC.CurrentNewFilePath; + R.Length = 0; + R.Offset = CurrentOldFileOffset[AddInfo->first]; + R.ReplacementText = AddInfo->second; + AddInfo.reset(); + } + replacements.push_back(R); + }; + + // Don't use std::getline as condition of the while loop, because it will + // return false if the last line only containing EOF. + while (iss.good()) { + std::getline(iss, line); + if (startsWith(line, "diff --git")) { + if (HC.InHunk) + addRepl(); + HC.InHunk = false; + HC.FastForward = false; + continue; + } + if (HC.FastForward) + continue; + + if (startsWith(line, "---")) { + if (HC.InHunk) + addRepl(); + HC.InHunk = false; + HC.CurrentOldFilePath = + line.substr(4) == "/dev/null" ? "/dev/null" : line.substr(6); + if (HC.CurrentOldFilePath != "/dev/null") { + std::ifstream FileStream(RepoRoot + "/" + HC.CurrentOldFilePath); + if (!FileStream.is_open()) { + throw std::runtime_error("Failed to open file: " + RepoRoot + "/" + + HC.CurrentOldFilePath); + } + std::stringstream Buffer; + Buffer << FileStream.rdbuf(); + CurrentOldFileOffset = calculateOldOffset(Buffer.str()); + } + continue; + } + if (startsWith(line, "+++")) { + if (HC.InHunk) + addRepl(); + HC.InHunk = false; + HC.CurrentNewFilePath = + line.substr(4) == "/dev/null" ? "/dev/null" : line.substr(6); + if (HC.CurrentOldFilePath == "/dev/null" || + HC.CurrentNewFilePath == "/dev/null") { + HC.FastForward = true; + Replacement R; + R.OldFilePath = HC.CurrentOldFilePath; + R.NewFilePath = HC.CurrentNewFilePath; + replacements.emplace_back(R); + } + continue; + } + + if (parseHunkHeader(line, HC)) { + if (HC.InHunk) + addRepl(); + // Hunk start + HC.InHunk = true; + continue; + } + + // parse hunk body + // 1. Assume the line ending in the file is '\n'. + // 2. The pair (---, +++) should occur only once in one hunk, since + // --unified=0 + // 3. We use a variable to save the delete (-) operation. The continuous + // delete operations are treated as one operation. + // 4. After the delete operation, if the next line is one or more '+' + // operations, we make them as a replace-replacement. If the next line is a + // context line, the delete operation is a delete-replacement. Then clear + // the variable. + // 5. If we meet insertions ('+') when the variable is empty, we treat it as + // an insert-replacement. + if (HC.InHunk) { + switch (line[0]) { + case '-': { + if (!DeleteInfo.has_value()) { + auto Item = std::pair( + HC.OldCurrentLine, + line.length()); // +1 for the newline character, -1 for the + // '-' in the line beginng + DeleteInfo = Item; + } else { + DeleteInfo->second += + (line.length()); // +1 for the newline character, -1 for the + // '-' in the line beginng + } + HC.OldCurrentLine++; + break; + } + case '+': { + if (!AddInfo.has_value()) { + auto Item = std::pair( + HC.OldCurrentLine, line.substr(1) + LineEnd); + AddInfo = Item; + } else { + AddInfo->second += (line.substr(1) + LineEnd); + } + break; + } + } + continue; + } + } + if (HC.InHunk) + addRepl(); + + return replacements; +} + +std::string execGitCommand(const std::string &CMD) { + std::array Buffer; + std::unique_ptr Pipe(popen(CMD.c_str(), "r"), pclose); + if (!Pipe) { + throw std::runtime_error("popen() failed!"); + } + + std::string Result; + while (fgets(Buffer.data(), Buffer.size(), Pipe.get()) != nullptr) { + Result += Buffer.data(); + } + return Result; +} diff --git a/clang/tools/gitdiff2yaml/src/gitdiff2yaml.h b/clang/tools/gitdiff2yaml/src/gitdiff2yaml.h new file mode 100644 index 000000000000..b22f930b87bb --- /dev/null +++ b/clang/tools/gitdiff2yaml/src/gitdiff2yaml.h @@ -0,0 +1,30 @@ +//===---------------------- gitdiff2yaml.h ----------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef __GITDIFF2YAML_H__ +#define __GITDIFF2YAML_H__ + +#include +#include + +const std::string LineEnd = "\n"; + +struct Replacement { + std::string NewFilePath; + std::string OldFilePath; + unsigned Offset = 0; + unsigned Length = 0; + std::string ReplacementText; +}; + +std::string execGitCommand(const std::string &CMD); +std::vector parseDiff(const std::string &diffOutput, + const std::string &RepoRoot); +void printYaml(std::ostream &stream, const std::vector &Repls); + +#endif // __GITDIFF2YAML_H__ diff --git a/clang/tools/gitdiff2yaml/src/main.cpp b/clang/tools/gitdiff2yaml/src/main.cpp new file mode 100644 index 000000000000..2f8d1bc25870 --- /dev/null +++ b/clang/tools/gitdiff2yaml/src/main.cpp @@ -0,0 +1,125 @@ +//===------------------------- main.cpp -------------------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// Usage: +// $ cd /path/to/your/git/repo +// $ gitdiff2yaml -c -o +// This will output the clang replacements in YAML format. +// Limitation: +// (1) The workspace and the staging area should be clean before running +// this tool. +// (2) The line ending in the file should be '\n'. +//===----------------------------------------------------------------------===// + +#include "gitdiff2yaml.h" +#include "llvm/Support/CommandLine.h" +#include "llvm/Support/Path.h" + +#include +#include +#include +#include +#include + +static constexpr auto Description = R"--( + +This gitdiff2yaml tool converts the output of `git diff` into a YAML file +containing a list of replacements. +)--"; + +static constexpr auto Examples = R"--( +EXAMPLES: + +Output to the terminal: + + gitdiff2yaml -c + +Output to a file: + + gitdiff2yaml -c -o Changes.yaml + +)--"; + +static llvm::cl::OptionCategory &getG2YCategory() { + static llvm::cl::OptionCategory G2YCategory("gitdiff2yaml tool options"); + return G2YCategory; +} + +int main(int argc, char *argv[]) { + llvm::cl::opt OutputFilename( + "o", llvm::cl::desc("[optional] Specify output filename"), + llvm::cl::value_desc("filename"), llvm::cl::cat(getG2YCategory()), + llvm::cl::Optional); + + llvm::cl::opt OldCommitID( + "c", llvm::cl::desc("[required] Specify the old commit ID"), + llvm::cl::value_desc("commit_id"), llvm::cl::cat(getG2YCategory()), + llvm::cl::Required); + + llvm::cl::extrahelp MoreHelp(Examples); + + llvm::cl::HideUnrelatedOptions(getG2YCategory()); + + llvm::cl::ParseCommandLineOptions(argc, argv, Description); + + std::string RepoRoot = execGitCommand("git rev-parse --show-toplevel"); + RepoRoot = RepoRoot.substr(0, RepoRoot.size() - 1); // Remove the last '\n' + + std::string NewCommitID = execGitCommand("git log -1 --format=\"%H\""); + std::string DiffOutput = + execGitCommand("git diff --diff-algorithm=minimal --unified=0 " + + OldCommitID.getValue()); + + execGitCommand("git reset --hard " + OldCommitID.getValue()); + std::vector Repls = parseDiff(DiffOutput, RepoRoot); + + // Erase emtpy replacements + Repls.erase(std::remove_if(Repls.begin(), Repls.end(), + [](Replacement x) { + return (x.NewFilePath == "" && + x.OldFilePath == "" && x.Offset == 0 && + x.Length == 0 && + x.ReplacementText == ""); + }), + Repls.end()); + + // Erase unrelated replacements + Repls.erase(std::remove_if( + Repls.begin(), Repls.end(), + [](Replacement x) { + if (x.NewFilePath == "/dev/null") + return false; + llvm::StringRef PathRef = x.NewFilePath; + std::string Ext = + llvm::sys::path::extension(PathRef).substr(1).lower(); + if (Ext == "cu" || Ext == "cuh" || Ext == "cpp" || + Ext == "hpp" || Ext == "cxx" || Ext == "hxx" || + Ext == "cc" || Ext == "hh" || Ext == "c" || + Ext == "h") { + return false; + } + return true; + }), + Repls.end()); + + if (!OutputFilename.empty()) { + std::ofstream OutFile(OutputFilename.getValue()); + if (!OutFile.is_open()) { + std::cerr << "Failed to open output file: " << OutputFilename.getValue() + << std::endl; + return 1; + } + printYaml(OutFile, Repls); + OutFile.close(); + } else { + printYaml(std::cout, Repls); + } + + execGitCommand("git reset --hard " + NewCommitID); + + return 0; +} diff --git a/clang/unittests/CMakeLists.txt b/clang/unittests/CMakeLists.txt index aef28f914b64..531e1893cb2d 100644 --- a/clang/unittests/CMakeLists.txt +++ b/clang/unittests/CMakeLists.txt @@ -75,39 +75,43 @@ function(add_clang_unittest test_name) set_property(GLOBAL APPEND PROPERTY CLANG_UNITTEST_LLVM_COMPONENTS ${ARG_LLVM_COMPONENTS}) endfunction() -add_subdirectory(Basic) -add_subdirectory(Lex) -add_subdirectory(Parse) -add_subdirectory(Driver) -if(CLANG_ENABLE_STATIC_ANALYZER) - add_subdirectory(Analysis) - add_subdirectory(StaticAnalyzer) -endif() -add_subdirectory(ASTMatchers) -add_subdirectory(AST) -add_subdirectory(CrossTU) -add_subdirectory(Tooling) -add_subdirectory(Format) -add_subdirectory(Frontend) -add_subdirectory(Rewrite) -add_subdirectory(Sema) -add_subdirectory(CodeGen) -if(HAVE_CLANG_REPL_SUPPORT) - add_subdirectory(Interpreter) -endif() -# FIXME: libclang unit tests are disabled on Windows due -# to failures, mostly in libclang.VirtualFileOverlay_*. -if(NOT WIN32 AND CLANG_TOOL_LIBCLANG_BUILD) - add_subdirectory(libclang) -endif() -add_subdirectory(DirectoryWatcher) -add_subdirectory(Index) -add_subdirectory(InstallAPI) -add_subdirectory(Serialization) -add_subdirectory(Support) -if (CLANG_ENABLE_CIR) - add_subdirectory(CIR) -endif() +# SYCLomatic_CUSTOMIZATION begin +add_subdirectory(DPCT) +# SYCLomatic_CUSTOMIZATION else +# add_subdirectory(Basic) +# add_subdirectory(Lex) +# add_subdirectory(Parse) +# add_subdirectory(Driver) +# if(CLANG_ENABLE_STATIC_ANALYZER) +# add_subdirectory(Analysis) +# add_subdirectory(StaticAnalyzer) +# endif() +# add_subdirectory(ASTMatchers) +# add_subdirectory(AST) +# add_subdirectory(CrossTU) +# add_subdirectory(Tooling) +# add_subdirectory(Format) +# add_subdirectory(Frontend) +# add_subdirectory(Rewrite) +# add_subdirectory(Sema) +# add_subdirectory(CodeGen) +# if(HAVE_CLANG_REPL_SUPPORT) +# add_subdirectory(Interpreter) +# endif() +# # FIXME: libclang unit tests are disabled on Windows due +# # to failures, mostly in libclang.VirtualFileOverlay_*. +# if(NOT WIN32 AND CLANG_TOOL_LIBCLANG_BUILD) +# add_subdirectory(libclang) +# endif() +# add_subdirectory(DirectoryWatcher) +# add_subdirectory(Index) +# add_subdirectory(InstallAPI) +# add_subdirectory(Serialization) +# add_subdirectory(Support) +# if (CLANG_ENABLE_CIR) +# add_subdirectory(CIR) +# endif() +# SYCLomatic_CUSTOMIZATION end # If we're doing a single merged clang unit test binary, add that target after # all the previous subdirectories have been processed. diff --git a/clang/unittests/DPCT/CMakeLists.txt b/clang/unittests/DPCT/CMakeLists.txt new file mode 100644 index 000000000000..1d36997625d4 --- /dev/null +++ b/clang/unittests/DPCT/CMakeLists.txt @@ -0,0 +1,12 @@ +add_clang_unittest(DpctTests + ReMigrationTest.cpp + CLANG_LIBS + DPCT + clangBasic + clangFrontend + clangRewrite + clangTooling + clangToolingCore + LLVM_COMPONENTS + Support + ) diff --git a/clang/unittests/DPCT/ReMigrationTest.cpp b/clang/unittests/DPCT/ReMigrationTest.cpp new file mode 100644 index 000000000000..28c4aa6ccab1 --- /dev/null +++ b/clang/unittests/DPCT/ReMigrationTest.cpp @@ -0,0 +1,405 @@ +#include "../../lib/DPCT/IncMigration/ReMigration.h" +#include "clang/Tooling/Core/Replacement.h" +#include "clang/Tooling/Core/UnifiedPath.h" +#include "gtest/gtest.h" +#include + +using namespace llvm; +using namespace clang::tooling; +using namespace clang::dpct; + +class ReMigrationTest1 : public ::testing::Test { +protected: + void SetUp() override {} + void TearDown() override {} +}; + +TEST_F(ReMigrationTest1, calculateUpdatedRanges) { + // clang-format off + // Old base: +/* +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +*/ + // Old migrated: +/* +aaa zzz ccc +ppp aaa bb ccc +aaa yyy ccc +aaa bb ccc +aaxxx bb ccc +aaa bb ccc +qqq aaa bb ccc +*/ + // New base: +/* +aaa bb ccc +Hi bb Everyone +aaa bb ccc +aaa bb Everyone +Hi Hello Everyone +aaa Hi bb ccc +aaa bb ccc +*/ + // clang-format on + + std::vector NewRepl; + Replacements Repls; + Replacements Expected; + + NewRepl.push_back({"file1.cpp", 4, 2, "zzz", true}); + NewRepl.push_back({"file1.cpp", 11, 0, "ppp ", true}); + NewRepl.push_back({"file1.cpp", 26, 2, "yyy", true}); + NewRepl.push_back({"file1.cpp", 46, 1, "xxx", true}); + NewRepl.push_back({"file1.cpp", 65, 1, "\nqqq", true}); + + llvm::cantFail( + Repls.add(Replacement("file1.cpp", 11, 11, "Hi bb Everyone\n"))); + llvm::cantFail( + Repls.add(Replacement("file1.cpp", 33, 11, "aaa bb Everyone\n"))); + llvm::cantFail( + Repls.add(Replacement("file1.cpp", 44, 11, "Hi Hello Everyone\n"))); + llvm::cantFail( + Repls.add(Replacement("file1.cpp", 55, 11, "aaa Hi bb ccc\n"))); + + llvm::cantFail(Expected.add(Replacement("file1.cpp", 4, 2, "zzz"))); + llvm::cantFail(Expected.add(Replacement("file1.cpp", 30, 2, "yyy"))); + + std::vector Result = + calculateUpdatedRanges(Repls, NewRepl); + ASSERT_EQ(Expected.size(), Result.size()); + size_t Num = Expected.size(); + auto ExpectedIt = Expected.begin(); + auto ResultIt = Result.begin(); + for (size_t i = 0; i < Num; ++i) { + EXPECT_EQ(ExpectedIt->getFilePath(), ResultIt->getFilePath()); + EXPECT_EQ(ExpectedIt->getOffset(), ResultIt->getOffset()); + EXPECT_EQ(ExpectedIt->getLength(), ResultIt->getLength()); + EXPECT_EQ(ExpectedIt->getReplacementText(), ResultIt->getReplacementText()); + ExpectedIt++; + ResultIt++; + } +} + +TEST_F(ReMigrationTest1, groupReplcementsByFile) { + std::string FilePath1 = + clang::tooling::UnifiedPath("file1.cpp").getCanonicalPath().str(); + std::string FilePath2 = + clang::tooling::UnifiedPath("file2.cpp").getCanonicalPath().str(); + std::string FilePath3 = + clang::tooling::UnifiedPath("file3.cpp").getCanonicalPath().str(); + + std::vector Repls = {{FilePath1, 0, 0, ""}, + {FilePath2, 3, 0, ""}, + {FilePath1, 1, 0, ""}, + {FilePath3, 4, 0, ""}, + {FilePath2, 2, 0, ""}}; + std::map> Expected = { + {FilePath1, + { + {FilePath1, 0, 0, ""}, + {FilePath1, 1, 0, ""}, + }}, + {FilePath2, + { + {FilePath2, 2, 0, ""}, + {FilePath2, 3, 0, ""}, + }}, + {FilePath3, {{FilePath3, 4, 0, ""}}}}; + + auto Result = groupReplcementsByFile(Repls); + EXPECT_EQ(Expected.size(), Result.size()); + + size_t Num = Expected.size(); + auto ExpectedIt = Expected.begin(); + auto ResultIt = Result.begin(); + for (size_t i = 0; i < Num; ++i) { + EXPECT_EQ(ExpectedIt->first, ResultIt->first); + EXPECT_EQ(ExpectedIt->second.size(), ResultIt->second.size()); + size_t SubNum = ExpectedIt->second.size(); + std::sort(ExpectedIt->second.begin(), ExpectedIt->second.end()); + std::sort(ResultIt->second.begin(), ResultIt->second.end()); + for (size_t j = 0; j < SubNum; ++j) { + EXPECT_EQ(ExpectedIt->second[j].getFilePath(), + ResultIt->second[j].getFilePath()); + EXPECT_EQ(ExpectedIt->second[j].getOffset(), + ResultIt->second[j].getOffset()); + EXPECT_EQ(ExpectedIt->second[j].getLength(), + ResultIt->second[j].getLength()); + EXPECT_EQ(ExpectedIt->second[j].getReplacementText(), + ResultIt->second[j].getReplacementText()); + } + ExpectedIt++; + ResultIt++; + } +} + +class ReMigrationTest2 : public ::testing::Test { +protected: + void SetUp() override {} + void TearDown() override {} + // clang-format off +/* +aaabbbbccc +dddeeeefff +ggghhhhiii +zzzzzzzzzz +yyyyyyyyyy +*/ + // clang-format on + static StringRef getLineStringUnittest(clang::tooling::UnifiedPath FilePath, + unsigned LineNumber) { + static std::vector LineString = { + "aaabbbbccc\n", "dddeeeefff\n", "ggghhhhiii\n", "zzzzzzzzzz\n", + "yyyyyyyyyy\n"}; + return StringRef(LineString[LineNumber - 1]); + } + static unsigned getLineNumberUnittest(clang::tooling::UnifiedPath FilePath, + unsigned Offset) { + static std::vector LineOffsets = {0, 11, 22, 33, 44}; + auto Iter = + std::upper_bound(LineOffsets.begin(), LineOffsets.end(), Offset); + if (Iter == LineOffsets.end()) + return LineOffsets.size(); + return std::distance(LineOffsets.begin(), Iter); + } + static unsigned + getLineBeginOffsetUnittest(clang::tooling::UnifiedPath FilePath, + unsigned LineNumber) { + static std::unordered_map LineOffsets = { + {1, 0}, {2, 11}, {3, 22}, {4, 33}, {5, 44}}; + return LineOffsets[LineNumber]; + } +}; + +TEST_F(ReMigrationTest2, splitReplInOrderToNotCrossLines) { + // Example: + // + // Original repl: + // (ccc\ndddeeeefff\nggg) =>（jjj\nkkk） + // (zz\nyy) =>（xx) + // + // Splitted repls: + // (ccc\n) =>（jjj\nkkkhhhhiii\n） + // (dddeeeefff\n) => "" + // (ggghhhhiii\n) => "" + // (zz\n) => (xxyyyyyyyy\n) + // (yyyyyyyyyy) => "" + + getLineStringHook = this->getLineStringUnittest; + getLineNumberHook = this->getLineNumberUnittest; + getLineBeginOffsetHook = this->getLineBeginOffsetUnittest; + std::vector Repls = { + {"file1.cpp", 7, 18, "jjj\nkkk", true}, {"file1.cpp", 41, 5, "xx", true}}; + std::vector Expected = { + Replacement("file1.cpp", 7, 4, "jjj\nkkkhhhhiii\n"), + Replacement("file1.cpp", 11, 11, ""), + Replacement("file1.cpp", 22, 11, ""), + Replacement("file1.cpp", 41, 3, "xxyyyyyyyy\n"), + Replacement("file1.cpp", 44, 11, "")}; + auto Temp = splitReplInOrderToNotCrossLines(Repls); + std::vector Result; + // drop the tag info + for (const auto &R : Temp) { + Result.push_back({R.getFilePath(), R.getOffset(), R.getLength(), + R.getReplacementText()}); + } + std::sort(Result.begin(), Result.end()); + EXPECT_EQ(Expected, Result); +} + +class ReMigrationTest3 : public ::testing::Test { +protected: + void SetUp() override {} + void TearDown() override {} + // clang-format off +/* +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +aaa bb ccc +*/ + // clang-format on + static StringRef getLineStringUnittest(clang::tooling::UnifiedPath FilePath, + unsigned LineNumber) { + static std::string S = "aaa bb ccc\n"; + return StringRef(S); + } + static unsigned getLineNumberUnittest(clang::tooling::UnifiedPath FilePath, + unsigned Offset) { + static std::vector LineOffsets = {0, 11, 22, 33, 44, + 55, 66, 77, 88, 99}; + auto Iter = + std::upper_bound(LineOffsets.begin(), LineOffsets.end(), Offset); + if (Iter == LineOffsets.end()) + return LineOffsets.size(); + return std::distance(LineOffsets.begin(), Iter); + } + static unsigned + getLineBeginOffsetUnittest(clang::tooling::UnifiedPath FilePath, + unsigned LineNumber) { + static std::unordered_map LineOffsets = { + {1, 0}, {2, 11}, {3, 22}, {4, 33}, {5, 44}, + {6, 55}, {7, 66}, {8, 77}, {9, 88}, {10, 99}}; + return LineOffsets[LineNumber]; + } +}; + +TEST_F(ReMigrationTest3, convertReplcementsLineString) { + // clang-format off + // After appling repls: +/* +aaa zzz ccc +ppp aaa bb ccc +aaa bb ccc +aaa yyy ccc +aaa bb ccc +aaa bb ccq +qqqaa bb ccc +aaa bb ddd + +eee bb ccc +*/ + // clang-format on + getLineStringHook = this->getLineStringUnittest; + getLineNumberHook = this->getLineNumberUnittest; + getLineBeginOffsetHook = this->getLineBeginOffsetUnittest; + std::vector Repls = { + {"file1.cpp", 4, 2, "zzz", true}, + {"file1.cpp", 64, 3, "q\nqqq", true}, + {"file1.cpp", 33, 11, "aaa yyy ccc\n", true}, + {"file1.cpp", 11, 0, "ppp ", true}, + {"file1.cpp", 84, 18, "ddd\neee", true}}; + std::map Expected = {{1, "aaa zzz ccc\n"}, + {2, "ppp aaa bb ccc\n"}, + {4, "aaa yyy ccc\n"}, + {6, "aaa bb ccq\nqqqaa bb ccc\n"}, + {7, ""}, + {8, "aaa bb ddd\neee bb ccc\n"}, + {9, ""}, + {10, ""}}; + auto Temp = convertReplcementsLineString(Repls); + std::map Result; + // drop the tag info + for (const auto &R : Temp) { + Result[R.first] = R.second.first; + } + EXPECT_EQ(Expected, Result); +} + +TEST_F(ReMigrationTest3, mergeMapsByLine) { + getLineBeginOffsetHook = this->getLineBeginOffsetUnittest; + getLineStringHook = this->getLineStringUnittest; + std::map MapA = {{1, "zzzzz\n"}, {3, "xxxx1\n"}, + {4, "wwww1\n"}, {5, "vvvvv\n"}, + {7, "ppppp\n"}, {9, "iiijjjkkk\n"}}; + std::map> MapB = { + {2, {"yyyyy\n", true}}, + {3, {"xxxx2\n", true}}, + {4, {"wwww2\n", true}}, + {7, {"", true}}, + {8, {"qqqqq", true}}}; + UnifiedPath FilePath("test.cu"); + auto Result = mergeMapsByLine(MapA, MapB, FilePath); + std::sort(Result.begin(), Result.end()); + + std::vector Expected = { + Replacement("test.cu", 0, 11, "zzzzz\n"), + Replacement("test.cu", 11, 11, "yyyyy\n"), + Replacement("test.cu", 22, 22, + "<<<<<<<\nxxxx1\nwwww1\n=======\nxxxx2\nwwww2\n>>>>>>>\n"), + Replacement("test.cu", 44, 11, "vvvvv\n"), + Replacement("test.cu", 66, 11, "<<<<<<<\nppppp\n=======\n>>>>>>>\n"), + Replacement("test.cu", 77, 11, "qqqqq"), + Replacement("test.cu", 88, 11, "iiijjjkkk\n"), + }; + + ASSERT_EQ(Expected.size(), Result.size()); + size_t Num = Expected.size(); + auto ExpectedIt = Expected.begin(); + auto ResultIt = Result.begin(); + for (size_t i = 0; i < Num; ++i) { + EXPECT_EQ(ExpectedIt->getFilePath(), ResultIt->getFilePath()); + EXPECT_EQ(ExpectedIt->getOffset(), ResultIt->getOffset()); + EXPECT_EQ(ExpectedIt->getLength(), ResultIt->getLength()); + EXPECT_EQ(ExpectedIt->getReplacementText(), ResultIt->getReplacementText()); + ExpectedIt++; + ResultIt++; + } + + EXPECT_EQ(Expected, Result); +} + +TEST_F(ReMigrationTest1, mergeC1AndC2) { + // clang-format off + // original file: +/* +0123456789 +0123456789 +0123456789 +0123456789 +0123456789 +*/ + // migrated file: +/* +0123456789aaa +bbb12345678ccc789 +0123456789 +*/ + // updated file: +/* +0123zzz456789aaa +bbyyy2345678xxxc78www0123456789 +*/ + // clang-format on + std::vector Repl_C1 = { + Replacement("file1.cu", 10, 0, "aaa"), + Replacement("file1.cu", 11, 1, "bbb"), + Replacement("file1.cu", 20, 20, "ccc"), + }; + GitDiffChanges Repl_C2; + Repl_C2.ModifyFileHunks = { + Replacement("file1.dp.cpp", 30, 2, "www"), + Replacement("file1.dp.cpp", 25, 2, "xxx"), + Replacement("file1.dp.cpp", 16, 2, "yyy"), + Replacement("file1.dp.cpp", 4, 0, "zzz"), + }; + const std::map FileNameMap = { + {UnifiedPath("file1.dp.cpp").getCanonicalPath().str(), + UnifiedPath("file1.cu").getCanonicalPath().str()}}; + + std::vector Result = + mergeC1AndC2(Repl_C1, Repl_C2, FileNameMap); + std::vector Expected = {{"file1.cu", 4, 0, "zzz", true}, + {"file1.cu", 10, 0, "aaa", false}, + {"file1.cu", 11, 2, "bbyyy", true}, + {"file1.cu", 20, 20, "xxxc", true}, + {"file1.cu", 42, 2, "www", true}}; + std::sort(Result.begin(), Result.end()); + std::sort(Expected.begin(), Expected.end()); + + ASSERT_EQ(Expected.size(), Result.size()); + size_t Num = Expected.size(); + auto ExpectedIt = Expected.begin(); + auto ResultIt = Result.begin(); + for (size_t i = 0; i < Num; ++i) { + EXPECT_EQ(ExpectedIt->getFilePath(), ResultIt->getFilePath()); + EXPECT_EQ(ExpectedIt->getOffset(), ResultIt->getOffset()); + EXPECT_EQ(ExpectedIt->getLength(), ResultIt->getLength()); + EXPECT_EQ(ExpectedIt->getReplacementText(), ResultIt->getReplacementText()); + EXPECT_EQ(ExpectedIt->containsManualFix(), ResultIt->containsManualFix()); + ExpectedIt++; + ResultIt++; + } +}