From 9b163ab5861b11754716d5d71c78a64f33c42487 Mon Sep 17 00:00:00 2001
From: Tang Jiajun <jiajun.tang@intel.com>
Date: Wed, 29 May 2024 10:43:39 +0800
Subject: [PATCH] [SYCLomatic] Remove duplicated replacement of host_ct
 function. (#2017)

Signed-off-by: Tang, Jiajun jiajun.tang@intel.com
---
 clang/lib/DPCT/AnalysisInfo.cpp               | 20 ++++++++----------
 clang/lib/DPCT/AnalysisInfo.h                 |  6 ++----
 .../dpct/cuda_arch_test/math_in_cuda_arch.cu  | 21 +++++++++++++++++++
 3 files changed, 32 insertions(+), 15 deletions(-)
 create mode 100644 clang/test/dpct/cuda_arch_test/math_in_cuda_arch.cu

diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp
index 46570a6c66a6..e13d7925bb54 100644
--- a/clang/lib/DPCT/AnalysisInfo.cpp
+++ b/clang/lib/DPCT/AnalysisInfo.cpp
@@ -1850,10 +1850,8 @@ void DpctGlobalInfo::processCudaArchMacro() {
   }
 }
 
-void DpctGlobalInfo::generateHostCode(
-    std::multimap<unsigned int, std::shared_ptr<clang::dpct::ExtReplacement>>
-        &ProcessedReplList,
-    HostDeviceFuncLocInfo Info, unsigned ID) {
+void DpctGlobalInfo::generateHostCode(tooling::Replacements &ProcessedReplList,
+                                      HostDeviceFuncLocInfo Info, unsigned ID) {
   std::vector<std::shared_ptr<ExtReplacement>> ExtraRepl;
 
   unsigned int Pos, Len;
@@ -1861,13 +1859,12 @@ void DpctGlobalInfo::generateHostCode(
   StringRef SR(OriginText);
   RewriteBuffer RB;
   RB.Initialize(SR.begin(), SR.end());
-  for (auto &Element : ProcessedReplList) {
-    auto R = Element.second;
-    unsigned ROffset = R->getOffset();
+  for (const auto &R : ProcessedReplList) {
+    unsigned ROffset = R.getOffset();
     if (ROffset >= Info.FuncStartOffset && ROffset <= Info.FuncEndOffset) {
       Pos = ROffset - Info.FuncStartOffset;
-      Len = R->getLength();
-      RB.ReplaceText(Pos, Len, R->getReplacementText());
+      Len = R.getLength();
+      RB.ReplaceText(Pos, Len, R.getReplacementText());
     }
   }
   Pos = Info.FuncNameOffset - Info.FuncStartOffset;
@@ -1953,8 +1950,9 @@ void DpctGlobalInfo::postProcess() {
           if (LocInfo.Type == HDFuncInfoType::HDFI_Call) {
             continue;
           }
-          auto &ReplLists =
-              FileMap[LocInfo.FilePath]->getReplsSYCL()->getReplMap();
+          tooling::Replacements ReplLists;
+          FileMap[LocInfo.FilePath]->getReplsSYCL()->emplaceIntoReplSet(
+              ReplLists);
           generateHostCode(ReplLists, LocInfo, Info.PostFixId);
         }
       }
diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h
index fc091b99dd35..b2712ab89942 100644
--- a/clang/lib/DPCT/AnalysisInfo.h
+++ b/clang/lib/DPCT/AnalysisInfo.h
@@ -1106,10 +1106,8 @@ class DpctGlobalInfo {
   void buildKernelInfo();
   void buildReplacements();
   void processCudaArchMacro();
-  void generateHostCode(
-      std::multimap<unsigned int, std::shared_ptr<clang::dpct::ExtReplacement>>
-          &ProcessedReplList,
-      HostDeviceFuncLocInfo Info, unsigned ID);
+  void generateHostCode(tooling::Replacements &ProcessedReplList,
+                        HostDeviceFuncLocInfo Info, unsigned ID);
   void postProcess();
   void cacheFileRepl(clang::tooling::UnifiedPath FilePath,
                      std::pair<std::shared_ptr<ExtReplacements>,
diff --git a/clang/test/dpct/cuda_arch_test/math_in_cuda_arch.cu b/clang/test/dpct/cuda_arch_test/math_in_cuda_arch.cu
new file mode 100644
index 000000000000..894cb707eb33
--- /dev/null
+++ b/clang/test/dpct/cuda_arch_test/math_in_cuda_arch.cu
@@ -0,0 +1,21 @@
+// RUN: dpct --format-range=none -out-root %T/cuda_arch_test/math_in_cuda_arch %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++14
+// RUN: FileCheck --input-file %T/cuda_arch_test/math_in_cuda_arch/math_in_cuda_arch.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/cuda_arch_test/math_in_cuda_arch/math_in_cuda_arch.dp.cpp -o %T/cuda_arch_test/math_in_cuda_arch/math_in_cuda_arch.dp.o %}
+
+// CHECK: float f(float x) {
+// CHECK-EMPTY: 
+// CHECK-NEXT:   return expf(x);
+// CHECK-EMPTY: 
+// CHECK-NEXT: }
+// CHECK-NEXT: float f_host_ct0(float x) {
+// CHECK-EMPTY: 
+// CHECK-NEXT:   return sycl::exp(x);
+// CHECK-EMPTY: 
+// CHECK-NEXT: }
+__device__ __host__ float f(float x) {
+#if defined(__CUDA_ARCH__)
+  return ::expf(x);
+#else
+  return std::exp(x);
+#endif
+}