[SYCLomatic] Fix missed header file in codpein cuda out-root folder. (#…

…2637)

Signed-off-by: Chen, Sheng S <[email protected]>

clang/lib/DPCT/FileGenerator/GenFiles.cpp

@@ -1087,97 +1087,88 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool,
         "continues.\n");
   }
 
-  // The necessary header files which have no replacements will be copied to
-  // "-out-root" directory.
-  for (const auto &Entry : IncludeFileMap) {
-    // Generated SYCL file in outroot. E.g., /path/to/outroot/a.dp.cpp
-    clang::tooling::UnifiedPath FilePath = Entry.first;
-    // Generated CUDA file in outroot_debug. E.g., /path/to/outroot_debug/a.cu
-    clang::tooling::UnifiedPath DebugFilePath = Entry.first;
-    // Original CUDA file in inroot. E.g., /path/to/inroot/a.cu
-    clang::tooling::UnifiedPath OriginalFilePath = Entry.first;
-    if (!Entry.second) {
-      bool IsExcluded = DpctGlobalInfo::isExcluded(FilePath);
-      if (IsExcluded) {
-        continue;
-      }
-      // Always migrate *.cuh files to *.dp.hpp files,
-      // Always migrate *.cu files to *.dp.cpp files.
-      SourceProcessType FileType = GetSourceFileType(FilePath);
-      SmallString<512> TempFilePath(FilePath.getCanonicalPath());
-      if (FileType & SPT_CudaHeader) {
-        path::replace_extension(TempFilePath,
-                                DpctGlobalInfo::getSYCLHeaderExtension());
-      } else if (FileType & SPT_CudaSource) {
-        path::replace_extension(TempFilePath,
-                                DpctGlobalInfo::getSYCLSourceExtension());
-      }
-      FilePath = TempFilePath;
-
-      if (!rewriteCanonicalDir(FilePath, InRoot, OutRoot)) {
-        continue;
-      }
-
-      if (dpct::DpctGlobalInfo::isCodePinEnabled() &&
-          !rewriteCanonicalDir(DebugFilePath, InRoot, CUDAMigratedOutRoot)) {
-        continue;
-      }
+  // Copy the necessary files (header files, src files without replacements) to
+  // OutRoot
+  auto copyNecessaryFiles = [&](clang::tooling::UnifiedPath OutRoot,
+                                bool IsCUDAMigrationOutRoot = false) {
+    for (const auto &Entry : IncludeFileMap) {
+      // Generated SYCL file in outroot. E.g., /path/to/outroot/a.dp.cpp
+      clang::tooling::UnifiedPath FilePath = Entry.first;
+      if (!Entry.second) {
+        bool IsExcluded = DpctGlobalInfo::isExcluded(FilePath);
+        if (IsExcluded) {
+          continue;
+        }
+        // Always migrate *.cuh files to *.dp.hpp files,
+        // Always migrate *.cu files to *.dp.cpp files.
+        SourceProcessType FileType = GetSourceFileType(FilePath);
+        SmallString<512> TempFilePath(FilePath.getCanonicalPath());
+        if ((FileType & SPT_CudaHeader) && !IsCUDAMigrationOutRoot) {
+          path::replace_extension(TempFilePath,
+                                  DpctGlobalInfo::getSYCLHeaderExtension());
+        } else if ((FileType & SPT_CudaSource) && !IsCUDAMigrationOutRoot) {
+          path::replace_extension(TempFilePath,
+                                  DpctGlobalInfo::getSYCLSourceExtension());
+        }
+        FilePath = TempFilePath;
 
-      // Check for another file with SYCL extension. For example
-      // In in-root we have
-      //  * src.cpp
-      //  * src.cu
-      // After migration we will end up replacing src.cpp with migrated src.cu
-      // when the --sycl-file-extension is cpp
-      // In such a case warn the user.
-      if (checkOverwriteAndWarn(FilePath.getCanonicalPath(),
-                                Entry.first.getCanonicalPath()))
-        continue;
+        if (!rewriteCanonicalDir(FilePath, InRoot, OutRoot)) {
+          continue;
+        }
 
-      // If the file needs no replacement and it already exist, don't
-      // make any changes
-      if (fs::exists(FilePath.getCanonicalPath())) {
-        // A header file with this name already exists.
-        llvm::errs() << "File '" << FilePath
-                     << "' already exists; skipping it.\n";
-        continue;
-      }
+        // Check for another file with SYCL extension. For example
+        // In in-root we have
+        //  * src.cpp
+        //  * src.cu
+        // After migration we will end up replacing src.cpp with migrated src.cu
+        // when the --sycl-file-extension is cpp
+        // In such a case warn the user.
+        if (checkOverwriteAndWarn(FilePath.getCanonicalPath(),
+                                  Entry.first.getCanonicalPath()))
+          continue;
 
-      createDirectories(path::parent_path(FilePath.getCanonicalPath()));
-      dpct::RawFDOStream Stream(FilePath.getCanonicalPath());
+        // If the file needs no replacement and it already exist, don't
+        // make any changes
+        if (fs::exists(FilePath.getCanonicalPath())) {
+          // A header file with this name already exists.
+          llvm::errs() << "File '" << FilePath
+                       << "' already exists; skipping it.\n";
+          continue;
+        }
 
-      llvm::Expected<FileEntryRef> Result =
-          Tool.getFiles().getFileRef(Entry.first.getCanonicalPath());
-      if (auto E = Result.takeError()) {
-        continue;
-      }
-      if (MapNames::PatternRewriters.empty()) {
-        Rewrite
-            .getEditBuffer(Sources.getOrCreateFileID(
-                *Result, clang::SrcMgr::C_User /*normal user code*/))
-            .write(Stream);
-      } else {
-        std::string OutputString;
-        llvm::raw_string_ostream RSW(OutputString);
-        Rewrite
-            .getEditBuffer(Sources.getOrCreateFileID(
-                *Result, clang::SrcMgr::C_User /*normal user code*/))
-            .write(RSW);
-        applyPatternRewriter(OutputString, Stream);
-      }
+        createDirectories(path::parent_path(FilePath.getCanonicalPath()));
+        dpct::RawFDOStream Stream(FilePath.getCanonicalPath());
 
-      // This will help us to detect the same output filename
-      // for two different input files
-      OutFilePath2InFilePath[FilePath.getCanonicalPath().str()] =
-          Entry.first.getCanonicalPath().str();
+        llvm::Expected<FileEntryRef> Result =
+            Tool.getFiles().getFileRef(Entry.first.getCanonicalPath());
+        if (auto E = Result.takeError()) {
+          continue;
+        }
+        if (MapNames::PatternRewriters.empty() || IsCUDAMigrationOutRoot) {
+          Rewrite
+              .getEditBuffer(Sources.getOrCreateFileID(
+                  *Result, clang::SrcMgr::C_User /*normal user code*/))
+              .write(Stream);
+        } else {
+          std::string OutputString;
+          llvm::raw_string_ostream RSW(OutputString);
+          Rewrite
+              .getEditBuffer(Sources.getOrCreateFileID(
+                  *Result, clang::SrcMgr::C_User /*normal user code*/))
+              .write(RSW);
+          applyPatternRewriter(OutputString, Stream);
+        }
 
-      if (dpct::DpctGlobalInfo::isCodePinEnabled()) {
-        // Copy non-replacement CUDA files into debug folder
-        fs::copy_file(OriginalFilePath.getCanonicalPath(),
-                      DebugFilePath.getCanonicalPath());
+        // This will help us to detect the same output filename
+        // for two different input files
+        OutFilePath2InFilePath[FilePath.getCanonicalPath().str()] =
+            Entry.first.getCanonicalPath().str();
       }
     }
-  }
+  };
+  copyNecessaryFiles(OutRoot);
+  if (dpct::DpctGlobalInfo::isCodePinEnabled())
+    copyNecessaryFiles(CUDAMigratedOutRoot, true);
 
   std::string ScriptFineName = "Makefile.dpct";
   if (!BuildScriptFile.empty())

clang/lib/Tooling/Tooling.cpp

@@ -720,7 +720,7 @@ static void injectResourceDir(CommandLineArguments &Args, const char *Argv0,
 // if return value is -1, means current input file \p File is not processed,
 // if return value < -1, report return value to upper caller,
 // other values are ignored.
-int ClangTool::processFiles(llvm::StringRef File,bool &ProcessingFailed,
+int ClangTool::processFiles(llvm::StringRef File, bool &ProcessingFailed,
                      bool &FileSkipped, int &StaticSymbol, ToolAction *Action) {
     //clear error# counter
     CurFileParseErrCnt=0;

clang/test/dpct/debug_test/copy_file_in_db/common/test1.cpp

@@ -0,0 +1,5 @@
+// RUN: echo
+// CHECK:  int foo1() {
+int foo1() {
+return 0;
+}

clang/test/dpct/debug_test/copy_file_in_db/compile_commands.json

@@ -0,0 +1,17 @@
+[
+    {
+        "command": "g++ common/test1.cpp",
+        "directory": "./",
+        "file": "./common/test1.cpp"
+    },
+    {
+        "command": "g++ test2.cpp",
+        "directory": "./",
+        "file": "./test2.cpp"
+    },
+    {
+        "command": "nvcc test3.cu",
+        "directory": "./",
+        "file": "./test3.cu"
+    }
+]

clang/test/dpct/debug_test/copy_file_in_db/test.h

@@ -0,0 +1,8 @@
+#include <cuda_runtime.h>
+
+// CHECK: /*
+// CHECK: DPCT1057:{{[0-9]+}}: Variable float_to_force was used in host code and device code. float_to_force type was updated to be used in SYCL device code and new float_to_force_host_ct1 was generated to be used in host code. You need to update the host code manually to use the new float_to_force_host_ct1.
+// CHECK: */
+// CHECK: static const float float_to_force_host_ct1 = (float)(1ll << 40);
+// CHECK: static dpct::constant_memory<const float, 0> float_to_force((float)(1ll << 40));
+static __constant__ const float float_to_force = (float)(1ll << 40);

clang/test/dpct/debug_test/copy_file_in_db/test2.cpp

@@ -0,0 +1,5 @@
+#include "test.h"
+// RUN: echo
+int foo2() {
+
+}

clang/test/dpct/debug_test/copy_file_in_db/test3.cu

@@ -0,0 +1,19 @@
+// RUN: mkdir -p %T/common
+// RUN: cd %T
+// RUN: cat %S/compile_commands.json > %T/compile_commands.json
+// RUN: cat %S/common/test1.cpp > %T/common/test1.cpp
+// RUN: cat %S/test2.cpp > %T/test2.cpp
+// RUN: cat %S/test3.cu > %T/test3.cu
+// RUN: cat %S/test.h > %T/test.h
+// RUN: dpct --in-root=%T --out-root=%T/out -p %T --format-range=none --cuda-include-path="%cuda-path/include" --enable-codepin
+// RUN: FileCheck %S/common/test1.cpp --match-full-lines --input-file %T/out_codepin_cuda/common/test1.cpp
+#include "test.h"
+
+__global__ void kernel(){
+    float a = float_to_force;
+}
+
+int main() {
+    kernel<<<1, 1>>>();
+    return 0;      
+}