Skip to content

Commit 9b163ab

Browse files
authored
[SYCLomatic] Remove duplicated replacement of host_ct function. (#2017)
Signed-off-by: Tang, Jiajun [email protected]
1 parent d73a4a5 commit 9b163ab

File tree

3 files changed

+32
-15
lines changed

3 files changed

+32
-15
lines changed

clang/lib/DPCT/AnalysisInfo.cpp

Lines changed: 9 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1850,24 +1850,21 @@ void DpctGlobalInfo::processCudaArchMacro() {
18501850
}
18511851
}
18521852

1853-
void DpctGlobalInfo::generateHostCode(
1854-
std::multimap<unsigned int, std::shared_ptr<clang::dpct::ExtReplacement>>
1855-
&ProcessedReplList,
1856-
HostDeviceFuncLocInfo Info, unsigned ID) {
1853+
void DpctGlobalInfo::generateHostCode(tooling::Replacements &ProcessedReplList,
1854+
HostDeviceFuncLocInfo Info, unsigned ID) {
18571855
std::vector<std::shared_ptr<ExtReplacement>> ExtraRepl;
18581856

18591857
unsigned int Pos, Len;
18601858
std::string OriginText = Info.FuncContentCache;
18611859
StringRef SR(OriginText);
18621860
RewriteBuffer RB;
18631861
RB.Initialize(SR.begin(), SR.end());
1864-
for (auto &Element : ProcessedReplList) {
1865-
auto R = Element.second;
1866-
unsigned ROffset = R->getOffset();
1862+
for (const auto &R : ProcessedReplList) {
1863+
unsigned ROffset = R.getOffset();
18671864
if (ROffset >= Info.FuncStartOffset && ROffset <= Info.FuncEndOffset) {
18681865
Pos = ROffset - Info.FuncStartOffset;
1869-
Len = R->getLength();
1870-
RB.ReplaceText(Pos, Len, R->getReplacementText());
1866+
Len = R.getLength();
1867+
RB.ReplaceText(Pos, Len, R.getReplacementText());
18711868
}
18721869
}
18731870
Pos = Info.FuncNameOffset - Info.FuncStartOffset;
@@ -1953,8 +1950,9 @@ void DpctGlobalInfo::postProcess() {
19531950
if (LocInfo.Type == HDFuncInfoType::HDFI_Call) {
19541951
continue;
19551952
}
1956-
auto &ReplLists =
1957-
FileMap[LocInfo.FilePath]->getReplsSYCL()->getReplMap();
1953+
tooling::Replacements ReplLists;
1954+
FileMap[LocInfo.FilePath]->getReplsSYCL()->emplaceIntoReplSet(
1955+
ReplLists);
19581956
generateHostCode(ReplLists, LocInfo, Info.PostFixId);
19591957
}
19601958
}

clang/lib/DPCT/AnalysisInfo.h

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1106,10 +1106,8 @@ class DpctGlobalInfo {
11061106
void buildKernelInfo();
11071107
void buildReplacements();
11081108
void processCudaArchMacro();
1109-
void generateHostCode(
1110-
std::multimap<unsigned int, std::shared_ptr<clang::dpct::ExtReplacement>>
1111-
&ProcessedReplList,
1112-
HostDeviceFuncLocInfo Info, unsigned ID);
1109+
void generateHostCode(tooling::Replacements &ProcessedReplList,
1110+
HostDeviceFuncLocInfo Info, unsigned ID);
11131111
void postProcess();
11141112
void cacheFileRepl(clang::tooling::UnifiedPath FilePath,
11151113
std::pair<std::shared_ptr<ExtReplacements>,
Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
// 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
2+
// RUN: FileCheck --input-file %T/cuda_arch_test/math_in_cuda_arch/math_in_cuda_arch.dp.cpp --match-full-lines %s
3+
// 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 %}
4+
5+
// CHECK: float f(float x) {
6+
// CHECK-EMPTY:
7+
// CHECK-NEXT: return expf(x);
8+
// CHECK-EMPTY:
9+
// CHECK-NEXT: }
10+
// CHECK-NEXT: float f_host_ct0(float x) {
11+
// CHECK-EMPTY:
12+
// CHECK-NEXT: return sycl::exp(x);
13+
// CHECK-EMPTY:
14+
// CHECK-NEXT: }
15+
__device__ __host__ float f(float x) {
16+
#if defined(__CUDA_ARCH__)
17+
return ::expf(x);
18+
#else
19+
return std::exp(x);
20+
#endif
21+
}

0 commit comments

Comments
 (0)