[SYCLomatic] Fix the migration when the file in the soft link.

clang/include/clang/Tooling/Core/UnifiedPath.h

@@ -25,23 +25,28 @@ class UnifiedPath {
   UnifiedPath() = default;
   UnifiedPath(const std::string &Path, const std::string &CWD = ".")
       : _Path(Path) {
+    makeAbsolute(CWD);
     makeCanonical(CWD);
   }
   UnifiedPath(const llvm::StringRef Path, const std::string &CWD = ".")
       : _Path(Path.str()) {
+    makeAbsolute(CWD);
     makeCanonical(CWD);
   }
   UnifiedPath(const llvm::Twine &Path, const std::string &CWD = ".")
       : _Path(Path.str()) {
+    makeAbsolute(CWD);
     makeCanonical(CWD);
   }
   UnifiedPath(const llvm::SmallVectorImpl<char> &Path,
               const std::string &CWD = ".") {
     _Path = std::string(Path.data(), Path.size());
+    makeAbsolute(CWD);
     makeCanonical(CWD);
   }
   UnifiedPath(const char *Path, const std::string &CWD = ".") {
     _Path = std::string(Path);
+    makeAbsolute(CWD);
     makeCanonical(CWD);
   }
   bool equalsTo(const std::string &RHS) {
@@ -61,16 +66,21 @@ class UnifiedPath {
   }
   llvm::StringRef getCanonicalPath() const noexcept { return _CanonicalPath; }
   llvm::StringRef getPath() const noexcept { return _Path; }
+  llvm::StringRef getAbsolutePath() const noexcept { return _AbsolutePath; }
   void setPath(const std::string &NewPath) {
     _Path = NewPath;
+    _AbsolutePath.clear();
     _CanonicalPath.clear();
     makeCanonical();
+    makeAbsolute();
   }
 
 private:
   void makeCanonical(const std::string &CWD = ".");
+  void makeAbsolute(const std::string &CWD = ".");
   std::string _Path;
   std::string _CanonicalPath;
+  std::string _AbsolutePath;
   static std::unordered_map<std::string, std::string> CanonicalPathCache;
 };
 bool operator==(const clang::tooling::UnifiedPath &LHS,

clang/lib/DPCT/AnalysisInfo.cpp

@@ -1451,11 +1451,11 @@ DpctGlobalInfo::getAbsolutePath(FileEntryRef File) {
   if (auto RealPath = File.getFileEntry().tryGetRealPathName();
       !RealPath.empty())
     return clang::tooling::UnifiedPath(RealPath);
-
   llvm::SmallString<512> FilePathAbs(File.getName());
   SM->getFileManager().makeAbsolutePath(FilePathAbs);
   return clang::tooling::UnifiedPath(FilePathAbs);
 }
+
 std::pair<clang::tooling::UnifiedPath, unsigned>
 DpctGlobalInfo::getLocInfo(SourceLocation Loc, bool *IsInvalid) {
   if (SM->isMacroArgExpansion(Loc)) {

clang/lib/DPCT/SaveNewFiles.cpp

@@ -36,6 +36,7 @@
 
 #include <algorithm>
 #include <cassert>
+#include <filesystem>
 #include <fstream>
 #include <queue>
 
@@ -150,6 +151,87 @@ bool rewriteDir(clang::tooling::UnifiedPath &FilePath,
   return true;
 }
 
+bool rewriteAbsoluteDir(clang::tooling::UnifiedPath &FilePath,
+                        const clang::tooling::UnifiedPath &InRoot,
+                        const clang::tooling::UnifiedPath &OutRoot) {
+#if defined(_WIN64)
+  std::string Filename = sys::path::filename(FilePath.getPath()).str();
+#endif
+
+  if (!isChildPath(InRoot, FilePath) || DpctGlobalInfo::isExcluded(FilePath)) {
+    // Skip rewriting file path if FilePath is not child of InRoot
+    // E.g,
+    //  FilePath : /path/to/inc/util.cuh
+    //    InRoot : /path/to/inroot
+    //   OutRoot : /path/to/outroot
+    //  AnalysisScope : /path/to
+    return false;
+  }
+  auto PathDiff = std::mismatch(path::begin(FilePath.getAbsolutePath()),
+                                path::end(FilePath.getAbsolutePath()),
+                                path::begin(InRoot.getAbsolutePath()));
+  SmallString<512> NewFilePath = OutRoot.getAbsolutePath();
+  path::append(NewFilePath, PathDiff.first,
+               path::end(FilePath.getAbsolutePath()));
+#if defined(_WIN64)
+  sys::path::remove_filename(NewFilePath);
+  sys::path::append(NewFilePath, Filename);
+#endif
+
+  FilePath = NewFilePath;
+  return true;
+}
+
+void createSymLink(const clang::tooling::UnifiedPath &FilePath,
+                   const clang::tooling::UnifiedPath &InRoot,
+                   const clang::tooling::UnifiedPath &OutRoot,
+                   bool RewriteFileName) {
+  if (llvm::sys::fs::exists(FilePath.getCanonicalPath())) {
+    auto SourcePath = FilePath;
+    while (
+        isChildPath(InRoot.getAbsolutePath(), SourcePath.getAbsolutePath())) {
+      auto AbsolutePath = SourcePath.getAbsolutePath();
+      if (llvm::sys::fs::is_symlink_file(AbsolutePath)) {
+        tooling::UnifiedPath SymbolPath =
+            std::filesystem::read_symlink(
+                std::filesystem::path(AbsolutePath.str()))
+                .string();
+        if (RewriteFileName) {
+          rewriteFileName(SymbolPath);
+          rewriteFileName(SourcePath);
+        }
+        // The code is to create symbol file and link the file to target file.
+        // The code will iterate to get the parent path of the
+        // absolute path of the file. Then create the target directory if the
+        // canonical path is not exists in the out root path.
+        auto ReplPath = SourcePath;
+        rewriteAbsoluteDir(ReplPath, InRoot, OutRoot);
+
+        auto AbsoluteParentPath = sys::path::parent_path(AbsolutePath);
+        if (llvm::sys::fs::is_symlink_file(AbsoluteParentPath)) {
+          createSymLink(AbsoluteParentPath, InRoot, OutRoot, false);
+        }
+        auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath);
+        rewriteAbsoluteDir(ParentPath, InRoot, OutRoot);
+        if (!llvm::sys::fs::exists(ParentPath.getAbsolutePath())) {
+          createDirectories(ParentPath.getAbsolutePath());
+        }
+        try {
+          std::filesystem::create_symlink(
+              std::filesystem::path(SymbolPath.getPath().str()),
+              std::filesystem::path(ReplPath.getAbsolutePath().str()));
+        } catch (const std::filesystem::filesystem_error &e) {
+          std::cerr << "Error creating symbolic link: " << e.what()
+                    << std::endl;
+        }
+      }
+      SourcePath = path::parent_path(AbsolutePath);
+      if (!isChildOrSamePath(OutRoot, SourcePath)) {
+        return;
+      }
+    }
+  }
+}
 void rewriteFileName(clang::tooling::UnifiedPath &FileName) {
   rewriteFileName(FileName, FileName);
 }
@@ -220,6 +302,26 @@ static bool checkOverwriteAndWarn(StringRef OutFilePath, StringRef InFilePath) {
 void processallOptionAction(clang::tooling::UnifiedPath &InRoot,
                             clang::tooling::UnifiedPath &OutRoot,
                             bool IsForSYCL) {
+  extern DpctOption<clang::dpct::opt, bool> ProcessAll;
+  if (ProcessAll) {
+    std::error_code EC;
+    for (fs::recursive_directory_iterator Iter(Twine(InRoot.getPath()), EC),
+         End;
+         Iter != End; Iter.increment(EC)) {
+          // Skip output directory if it is in the in-root directory.
+      if (Iter->type() == fs::file_type::symlink_file) {
+
+        tooling::UnifiedPath RootSource = Iter->path();
+        if (IncludeFileMap.find(Iter->path()) != IncludeFileMap.end()) {
+          if (IncludeFileMap[RootSource]) {
+            rewriteFileName(RootSource);
+          }
+        }
+        createSymLink(Iter->path(), InRoot, OutRoot,
+                      IncludeFileMap[Iter->path()]);
+      }
+    }
+  }
   for (const auto &File : FilesNotInCompilationDB) {
     if (IncludeFileMap.find(File) != IncludeFileMap.end()) {
       // Skip the files parsed by dpct parser.

clang/lib/Tooling/Core/UnifiedPath.cpp

@@ -82,6 +82,29 @@ void UnifiedPath::makeCanonical(const std::string &CWD) {
 #endif
   CanonicalPathCache.insert(std::pair(_Path, _CanonicalPath));
 }
+
+void UnifiedPath::makeAbsolute(const std::string &CWD) {
+  if (_Path.empty()) {
+    return;
+  }
+  llvm::SmallString<512> Path(_Path);
+  llvm::sys::fs::expand_tilde(Path, Path);
+  if (!llvm::sys::path::is_absolute(Path)) {
+    llvm::SmallString<512> TempPath;
+    if (CWD == ".") {
+      llvm::sys::fs::current_path(TempPath);
+    } else {
+      UnifiedPath UnifiedCWD(CWD);
+      TempPath = UnifiedCWD.getCanonicalPath();
+    }
+    llvm::sys::path::append(TempPath, llvm::sys::path::Style::native, Path);
+    Path = TempPath;
+  }
+
+  llvm::sys::path::remove_dots(Path, /* remove_dot_dot= */ true);
+  llvm::sys::path::native(Path);
+  _AbsolutePath = Path.str();
+}
 std::unordered_map<std::string, std::string> UnifiedPath::CanonicalPathCache;
 bool operator==(const clang::tooling::UnifiedPath &LHS,
                 const clang::tooling::UnifiedPath &RHS) {

clang/lib/Tooling/Tooling.cpp

@@ -1120,9 +1120,8 @@ int ClangTool::run(ToolAction *Action) {
   // if input file(s) is not specified in command line, and the process-all
   // option is given in the comomand line, dpct tries to migrate or copy all
   // files from -in-root to the output directory.
-  if(SourcePaths.size() == 0 && DoGetRunRound() == 0) {
+  if (SourcePaths.size() == 0 && DoGetRunRound() == 0) {
     std::vector<std::string> FilesNotProcessed;
-
     // To traverse all the files in the directory specified by
     // -in-root, collecting *.cu files not processed by the first loop of
     // calling processFiles() into FilesNotProcessed, and copies the rest

clang/test/dpct/soft_link/soft_link_dir/link/test

@@ -0,0 +1 @@
+../target

clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp

@@ -0,0 +1,6 @@
+// CHECK: #include <sycl/sycl.hpp>
+// CHECK: #include <dpct/dpct.hpp>
+#include <cuda.h>
+int test() {
+  return 0;
+}

clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu

@@ -0,0 +1,54 @@
+// RUN: cd %S/link && rm -rf test && ln -nfs ../target test
+// RUN: dpct  --in-root=%S --out-root=%T/out  %s --cuda-include-path="%cuda-path/include"  --process-all -- -I %S/link/test -x cuda --cuda-host-only 
+// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s
+// RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp
+// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST  %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/test %}
+
+#include <cuda.h>
+#include <stdio.h>
+#include "test/test.hpp"
+#define VECTOR_SIZE 256
+
+__global__ void VectorAddKernel(float* A, float* B, float* C)
+{
+    A[threadIdx.x] = threadIdx.x + 1.0f;
+    B[threadIdx.x] = threadIdx.x + 1.0f;
+    C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x];
+}
+
+
+
+int main()
+{
+  //      CHECK:    dpct::device_ext &dev_ct1 = dpct::get_current_device();
+  // CHECK-NEXT:    sycl::queue &q_ct1 = dev_ct1.in_order_queue();
+  float *d_A, *d_B, *d_C;
+
+
+  cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float));
+  cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float));
+  cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float));
+
+
+  VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C);
+
+
+  float Result[VECTOR_SIZE] = {};
+  cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost);
+
+  //      CHECK:  dpct::dpct_free(d_A, q_ct1);
+  // CHECK-NEXT:  dpct::dpct_free(d_B, q_ct1);
+  // CHECK-NEXT:  dpct::dpct_free(d_C, q_ct1);
+  cudaFree(d_A);
+  cudaFree(d_B);
+  cudaFree(d_C);
+
+  for (int i = 0; i < VECTOR_SIZE; i++) {
+    if (i % 16 == 0) {
+      printf("\n");
+    }
+    printf("%f ", Result[i]);
+  }
+
+    return 0;
+}

clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp

@@ -0,0 +1 @@
+../../target/test/test.hpp

clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp

@@ -0,0 +1,6 @@
+// CHECK: #include <sycl/sycl.hpp>
+// CHECK: #include <dpct/dpct.hpp>
+#include <cuda.h>
+int test() {
+  return 0;
+}

clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu

@@ -0,0 +1,54 @@
+// RUN: cd %S/link/test &&  rm test.hpp && ln -nfs  ../../target/test/test.hpp test.hpp
+// RUN: dpct  --in-root=%S --out-root=%T/out  %s --cuda-include-path="%cuda-path/include" --process-all -- -I %S/link -x cuda --cuda-host-only
+// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s
+// RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp
+// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST  %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %}
+
+#include <cuda.h>
+#include <stdio.h>
+#include "test/test.hpp"
+#define VECTOR_SIZE 256
+
+__global__ void VectorAddKernel(float* A, float* B, float* C)
+{
+    A[threadIdx.x] = threadIdx.x + 1.0f;
+    B[threadIdx.x] = threadIdx.x + 1.0f;
+    C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x];
+}
+
+
+
+int main()
+{
+  //      CHECK:    dpct::device_ext &dev_ct1 = dpct::get_current_device();
+  // CHECK-NEXT:    sycl::queue &q_ct1 = dev_ct1.in_order_queue();
+  float *d_A, *d_B, *d_C;
+
+
+  cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float));
+  cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float));
+  cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float));
+
+
+  VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C);
+
+
+  float Result[VECTOR_SIZE] = {};
+  cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost);
+
+  //      CHECK:  dpct::dpct_free(d_A, q_ct1);
+  // CHECK-NEXT:  dpct::dpct_free(d_B, q_ct1);
+  // CHECK-NEXT:  dpct::dpct_free(d_C, q_ct1);
+  cudaFree(d_A);
+  cudaFree(d_B);
+  cudaFree(d_C);
+
+  for (int i = 0; i < VECTOR_SIZE; i++) {
+    if (i % 16 == 0) {
+      printf("\n");
+    }
+    printf("%f ", Result[i]);
+  }
+
+    return 0;
+}

clang/test/dpct/soft_link/soft_link_file_rename/link/test/test.cuh

@@ -0,0 +1 @@
+../../target/test/test.cuh

clang/test/dpct/soft_link/soft_link_file_rename/target/test/test.cuh

@@ -0,0 +1,6 @@
+// CHECK: #include <sycl/sycl.hpp>
+// CHECK: #include <dpct/dpct.hpp>
+#include <cuda.h>
+int test() {
+  return 0;
+}