Skip to content

Commit 69ee755

Browse files
authored
[SYCLomatic] Fix the issue that group utils header file not inserted but required (#2776)
Signed-off-by: intwanghao <hao3.wang@intel.com>
1 parent 74b957d commit 69ee755

File tree

3 files changed

+14
-6
lines changed

3 files changed

+14
-6
lines changed

clang/lib/DPCT/AnalysisInfo.cpp

Lines changed: 8 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6113,9 +6113,14 @@ void KernelCallExpr::addAccessorDecl() {
61136113
Tex.second->addDecl(OuterStmts.InitList, SubmitStmts.TextureList,
61146114
SubmitStmts.SamplerList, getQueueStr());
61156115
}
6116-
for (auto &Tmp : VM.getTempStorageMap()) {
6117-
Tmp.second->addAccessorDecl(SubmitStmts.AccessorList,
6118-
ExecutionConfig.LocalSize);
6116+
if (!VM.getTempStorageMap().empty()) {
6117+
DpctGlobalInfo::getInstance()
6118+
.insertFile(getFilePath())
6119+
->insertHeader(HT_DPCT_GROUP_Utils, RT_ForSYCLMigration);
6120+
for (auto &Tmp : VM.getTempStorageMap()) {
6121+
Tmp.second->addAccessorDecl(SubmitStmts.AccessorList,
6122+
ExecutionConfig.LocalSize);
6123+
}
61196124
}
61206125
}
61216126
void KernelCallExpr::buildInfo() {

clang/lib/DPCT/RulesLangLib/CUBAPIMigration.cpp

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -247,6 +247,9 @@ void CubMemberCallRule::runRule(
247247
if (const auto *BlockMC =
248248
getNodeAsType<CXXMemberCallExpr>(Result, "memberCall")) {
249249
EA.analyze(BlockMC);
250+
emplaceTransformation(EA.getReplacement());
251+
EA.applyAllSubExprRepl();
252+
250253
StringRef Name = BlockMC->getMethodDecl()->getName();
251254
bool isBlockRadixSort = Name == "Sort" || Name == "SortDescending" ||
252255
Name == "SortBlockedToStriped" ||
@@ -343,9 +346,9 @@ void CubMemberCallRule::runRule(
343346
}
344347
} else if (const auto *E2 = getNodeAsType<MemberExpr>(Result, "memberExpr")) {
345348
EA.analyze(E2);
349+
emplaceTransformation(EA.getReplacement());
350+
EA.applyAllSubExprRepl();
346351
}
347-
emplaceTransformation(EA.getReplacement());
348-
EA.applyAllSubExprRepl();
349352
}
350353

351354
void CubIntrinsicRule::registerMatcher(ast_matchers::MatchFinder &MF) {

clang/test/dpct/cub/blocklevel/blockload.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@
66
// RUN: %if build_lit %{icpx -c -fsycl %T/blocklevel/blockload/blockload.dp.cpp -o %T/blocklevel/blockload/blockload.dp.o %}
77

88
#include <cub/cub.cuh>
9-
9+
// CHECK: #include <dpct/group_utils.hpp>
1010
__global__ void DirectKernel(int *d_data, int valid_items, int default_value) {
1111
int thread_data[4] = {0};
1212
// CHECK: dpct::group::load_direct_blocked(item_ct1, d_data, thread_data);

0 commit comments

Comments
 (0)