[SYCLomatic] Keep CUDASyntaxInfo cross two runs

clang/lib/DPCT/AnalysisInfo.cpp

@@ -20,6 +20,7 @@
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/OperationKinds.h"
 #include "clang/ASTMatchers/ASTMatchers.h"
+#include "clang/Tooling/Core/UnifiedPath.h"
 #include "clang/Tooling/Tooling.h"
 #include <algorithm>
 #include <deque>
@@ -829,6 +830,19 @@ void DpctFileInfo::setFirstIncludeOffset(unsigned Offset) {
     HasInclusionDirectiveSet.insert(std::move(MF));
   }
 }
+void DpctFileInfo::setHeaderInserted(HeaderType Header) {
+  DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath][Header] = true;
+}
+void DpctFileInfo::setMathHeaderInserted(bool B) {
+  DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath][HeaderType::HT_Math] = B;
+}
+void DpctFileInfo::setAlgorithmHeaderInserted(bool B) {
+  DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath]
+                                           [HeaderType::HT_Algorithm] = B;
+}
+void DpctFileInfo::setTimeHeaderInserted(bool B) {
+  DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath][HeaderType::HT_Time] = B;
+}
 void DpctFileInfo::concatHeader(llvm::raw_string_ostream &OS) {}
 template <class FirstT, class... Args>
 void DpctFileInfo::concatHeader(llvm::raw_string_ostream &OS, FirstT &&First,
@@ -862,9 +876,9 @@ void DpctFileInfo::insertHeader(HeaderType Type, unsigned Offset,
             Type, FirstIncludeOffset.at(MF));
       }
   }
-  if (HeaderInsertedBitMap[Type])
+  if (DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath][Type])
     return;
-  HeaderInsertedBitMap[Type] = true;
+  DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath][Type] = true;
   std::string ReplStr;
   llvm::raw_string_ostream OS(ReplStr);
   std::string MigratedMacroDefinitionStr;
@@ -898,10 +912,11 @@ void DpctFileInfo::insertHeader(HeaderType Type, unsigned Offset,
     concatHeader(OS, getHeaderSpelling(Type));
     if (DpctGlobalInfo::useSYCLCompat()) {
       concatHeader(OS, getHeaderSpelling(HT_COMPAT_SYCLcompat));
-      HeaderInsertedBitMap[HT_COMPAT_SYCLcompat] = true;
+      DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath]
+                                               [HT_COMPAT_SYCLcompat] = true;
     } else {
       concatHeader(OS, getHeaderSpelling(HT_DPCT_Dpct));
-      HeaderInsertedBitMap[HT_DPCT_Dpct] = true;
+      DpctGlobalInfo::getHeaderInsertedBitMap()[FilePath][HT_DPCT_Dpct] = true;
     }
     DpctGlobalInfo::printUsingNamespace(OS);
     if (DpctGlobalInfo::useNoQueueDevice()) {
@@ -1456,8 +1471,8 @@ std::string DpctGlobalInfo::getStringForRegexReplacement(StringRef MatchedStr) {
     return getStringForRegexDefaultQueueAndDevice(
         HelperFuncType::HFT_DefaultQueuePtr, Index);
   case 'E': {
-    auto &Vec = DpctGlobalInfo::getInstance().getCSourceFileInfo();
-    return Vec[Index]->hasCUDASyntax()
+    auto &Vec = DpctGlobalInfo::getInstance().getCSourceFileList();
+    return DpctGlobalInfo::hasCUDASyntax(Vec[Index])
                ? ("c" + DpctGlobalInfo::getSYCLSourceExtension())
                : "c";
   }
@@ -2493,7 +2508,7 @@ bool DpctGlobalInfo::CVersionCUDALaunchUsedFlag = false;
 unsigned int DpctGlobalInfo::ColorOption = 1;
 std::unordered_map<int, std::shared_ptr<DeviceFunctionInfo>>
     DpctGlobalInfo::CubPlaceholderIndexMap;
-std::vector<std::shared_ptr<DpctFileInfo>> DpctGlobalInfo::CSourceFileInfo;
+std::vector<tooling::UnifiedPath> DpctGlobalInfo::CSourceFileList;
 bool DpctGlobalInfo::OptimizeMigrationFlag = false;
 std::unordered_map<std::string, std::shared_ptr<PriorityReplInfo>>
     DpctGlobalInfo::PriorityReplInfoMap;
@@ -2520,6 +2535,8 @@ std::vector<std::pair<std::string, std::vector<std::string>>>
 std::unordered_set<std::string> DpctGlobalInfo::NeedParenAPISet = {};
 std::unordered_set<std::string>
     DpctGlobalInfo::CustomHelperFunctionAddtionalIncludes = {};
+std::unordered_map<clang::tooling::UnifiedPath, std::bitset<32>>
+    DpctGlobalInfo::HeaderInsertedBitMap = {};
 ///// class DpctNameGenerator /////
 void DpctNameGenerator::printName(const FunctionDecl *FD,
                                   llvm::raw_ostream &OS) {

clang/lib/DPCT/AnalysisInfo.h

@@ -37,6 +37,7 @@
 
 #include "clang/Format/Format.h"
 #include "clang/Frontend/CompilerInstance.h"
+#include "clang/Tooling/Core/UnifiedPath.h"
 
 llvm::StringRef getReplacedName(const clang::NamedDecl *D);
 void setGetReplacedNamePtr(llvm::StringRef (*Ptr)(const clang::NamedDecl *D));
@@ -404,18 +405,10 @@ class DpctFileInfo {
   void setFileEnterOffset(unsigned Offset);
   void setFirstIncludeOffset(unsigned Offset);
   void setLastIncludeOffset(unsigned Offset) { LastIncludeOffset = Offset; }
-  void setHeaderInserted(HeaderType Header) {
-    HeaderInsertedBitMap[Header] = true;
-  }
-  void setMathHeaderInserted(bool B = true) {
-    HeaderInsertedBitMap[HeaderType::HT_Math] = B;
-  }
-  void setAlgorithmHeaderInserted(bool B = true) {
-    HeaderInsertedBitMap[HeaderType::HT_Algorithm] = B;
-  }
-  void setTimeHeaderInserted(bool B = true) {
-    HeaderInsertedBitMap[HeaderType::HT_Time] = B;
-  }
+  void setHeaderInserted(HeaderType Header);
+  void setMathHeaderInserted(bool B = true);
+  void setAlgorithmHeaderInserted(bool B = true);
+  void setTimeHeaderInserted(bool B = true);
 
   void concatHeader(llvm::raw_string_ostream &OS);
   template <class FirstT, class... Args>
@@ -540,7 +533,6 @@ class DpctFileInfo {
 
   void setCCLVerValue(std::string Value) { CCLVerValue = Value; }
   std::string getCCLVerValue() { return CCLVerValue; }
-  bool hasCUDASyntax() { return HeaderInsertedBitMap[HeaderType::HT_SYCL]; }
 
   std::shared_ptr<tooling::TranslationUnitReplacements> PreviousTUReplFromYAML =
       nullptr;
@@ -602,7 +594,6 @@ class DpctFileInfo {
   std::set<std::shared_ptr<DpctFileInfo> /*MainFile*/> HasInclusionDirectiveSet;
   std::vector<std::string> InsertedHeaders;
   std::vector<std::string> InsertedHeadersCUDA;
-  std::bitset<32> HeaderInsertedBitMap;
   std::bitset<32> UsingInsertedBitMap;
   bool AddOneDplHeaders = false;
   std::vector<std::shared_ptr<ExtReplacement>> IncludeDirectiveInsertions;
@@ -1397,6 +1388,10 @@ class DpctGlobalInfo {
   getCustomHelperFunctionAddtionalIncludes() {
     return CustomHelperFunctionAddtionalIncludes;
   }
+  static std::unordered_map<clang::tooling::UnifiedPath, std::bitset<32>> &
+  getHeaderInsertedBitMap() {
+    return HeaderInsertedBitMap;
+  }
   std::shared_ptr<DpctFileInfo>
   insertFile(const clang::tooling::UnifiedPath &FilePath) {
     return insertObject(FileMap, FilePath);
@@ -1440,8 +1435,8 @@ class DpctGlobalInfo {
   getCubPlaceholderIndexMap() {
     return CubPlaceholderIndexMap;
   }
-  std::vector<std::shared_ptr<DpctFileInfo>> &getCSourceFileInfo() {
-    return CSourceFileInfo;
+  std::vector<tooling::UnifiedPath> &getCSourceFileList() {
+    return CSourceFileList;
   }
   static std::unordered_map<std::string, std::shared_ptr<PriorityReplInfo>> &
   getPriorityReplInfoMap() {
@@ -1500,6 +1495,9 @@ class DpctGlobalInfo {
   static void printUsingNamespace(llvm::raw_ostream &);
   // #tokens, name of the second token, SourceRange of a macro
   static std::tuple<unsigned int, std::string, SourceRange> LastMacroRecord;
+  static bool hasCUDASyntax(tooling::UnifiedPath FilePath) {
+    return HeaderInsertedBitMap[FilePath][HeaderType::HT_SYCL];
+  }
 
 private:
   DpctGlobalInfo();
@@ -1690,7 +1688,7 @@ class DpctGlobalInfo {
   static unsigned int ColorOption;
   static std::unordered_map<int, std::shared_ptr<DeviceFunctionInfo>>
       CubPlaceholderIndexMap;
-  static std::vector<std::shared_ptr<DpctFileInfo>> CSourceFileInfo;
+  static std::vector<tooling::UnifiedPath> CSourceFileList;
   static bool OptimizeMigrationFlag;
   static std::unordered_map<std::string, std::shared_ptr<PriorityReplInfo>>
       PriorityReplInfoMap;
@@ -1719,6 +1717,8 @@ class DpctGlobalInfo {
       CodePinDumpFuncDepsVec;
   static std::unordered_set<std::string> NeedParenAPISet;
   static std::unordered_set<std::string> CustomHelperFunctionAddtionalIncludes;
+  static std::unordered_map<clang::tooling::UnifiedPath, std::bitset<32>>
+      HeaderInsertedBitMap;
 };
 
 /// Generate mangle name of FunctionDecl as key of DeviceFunctionInfo.

clang/lib/DPCT/FileGenerator/GenFiles.cpp

@@ -213,12 +213,10 @@ void rewriteFileName(std::string &FileName, const std::string &FullPathName) {
                               DpctGlobalInfo::getSYCLSourceExtension());
     } else if (FileType & SPT_CppSource) {
       if (Extension == ".c") {
-        if (auto FileInfo = DpctGlobalInfo::getInstance().findFile(FileName)) {
-          if (FileInfo->hasCUDASyntax()) {
-            path::replace_extension(
-                CanonicalPathStr,
-                Extension + DpctGlobalInfo::getSYCLSourceExtension());
-          }
+        if (DpctGlobalInfo::hasCUDASyntax(FileName)) {
+          path::replace_extension(CanonicalPathStr,
+                                  Extension +
+                                      DpctGlobalInfo::getSYCLSourceExtension());
         }
       } else {
         path::replace_extension(CanonicalPathStr,
@@ -578,11 +576,8 @@ int writeReplacementsToFiles(
         auto Hash = llvm::sys::fs::md5_contents(Entry.first);
 
         bool HasCUDASyntax = false;
-        if (auto FileInfo =
-                dpct::DpctGlobalInfo::getInstance().findFile(Entry.first)) {
-          if (FileInfo->hasCUDASyntax()) {
-            HasCUDASyntax = true;
-          }
+        if (DpctGlobalInfo::hasCUDASyntax(Entry.first)) {
+          HasCUDASyntax = true;
         }
 
         bool IsMainSrcFileChanged = false;

clang/lib/DPCT/RulesInclude/InclusionHeaders.cpp

@@ -161,10 +161,10 @@ void IncludesCallbacks::InclusionDirective(
     SmallString<512> NewFileName(FileName.str());
 
     if (Extension == ".c") {
-      auto &Vec = DpctGlobalInfo::getInstance().getCSourceFileInfo();
+      auto &Vec = DpctGlobalInfo::getInstance().getCSourceFileList();
       path::replace_extension(
           NewFileName, "{{NEEDREPLACEE" + std::to_string(Vec.size()) + "}}");
-      Vec.push_back(DpctGlobalInfo::getInstance().insertFile(IncludedFile));
+      Vec.push_back(IncludedFile);
     } else {
       clang::tooling::UnifiedPath NewFilePath = FileName;
       rewriteFileName(NewFilePath, IncludedFile);

clang/test/dpct/CUDASyntaxInfo/a.cpp

@@ -0,0 +1,6 @@
+// UNSUPPORTED: system-windows
+// RUN: echo "empty command"
+
+#include "cuda_runtime.h"
+
+float2 f2;

clang/test/dpct/CUDASyntaxInfo/a.cu

@@ -0,0 +1,20 @@
+// UNSUPPORTED: system-windows
+// RUN: cd %S
+// RUN: dpct --out-root %T -p . --cuda-include-path="%cuda-path/include"
+// RUN: FileCheck %s --match-full-lines --input-file %T/MainSourceFiles.yaml
+// RUN: %if build_lit %{icpx -c -fsycl %T/a.cpp.dp.cpp -o %T/a.cpp.dp.o %}
+// RUN: %if build_lit %{icpx -c -fsycl %T/a.dp.cpp -o %T/a.dp.o %}
+
+// CHECK: HasCUDASyntax:   true
+// CHECK: HasCUDASyntax:   true
+
+__host__ __device__ int test(int axis) {
+  int a = 1;
+#if !defined(__CUDA_ARCH__)
+  cudaDeviceSynchronize();
+  return 0;
+#else
+  return axis == 0 ? threadIdx.x : axis == 1 ? threadIdx.y : threadIdx.z;
+#endif
+  return a;
+}

clang/test/dpct/CUDASyntaxInfo/compile_commands.json

@@ -0,0 +1,12 @@
+[
+    {
+        "command": "c++ -c a.cpp",
+        "directory": ".",
+        "file": "a.cpp"
+    },
+    {
+        "command": "nvcc -c a.cu",
+        "directory": ".",
+        "file": "a.cu"
+    }
+]