Skip to content

[SYCLomatic] Fix the migration of __syncthreads in the un-enabled if constexpr branch #2795

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 2 commits into
base: SYCLomatic
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
35 changes: 34 additions & 1 deletion clang/lib/DPCT/RulesLang/RulesLang.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -7733,6 +7733,38 @@ void SyncThreadsMigrationRule::registerMatcher(MatchFinder &MF) {
this);
}

bool SyncThreadsMigrationRule::noCorrespondingCEInInstantiatedTemplates(
const FunctionTemplateDecl *FTD, const CallExpr *CE) {
const auto &SM = DpctGlobalInfo::getSourceManager();
std::string FuncName =
CE->getDirectCallee()->getNameInfo().getName().getAsString();
auto CEMatcher = ast_matchers::findAll(
ast_matchers::callExpr(callee(functionDecl(hasName(FuncName))))
.bind("call"));
SourceLocation CELocation = SM.getSpellingLoc(CE->getBeginLoc());
auto DecomposedCELocation = SM.getDecomposedLoc(CELocation);
for (const auto &Spec : FTD->specializations()) {
if (!(Spec->hasBody()))
continue;
auto MatchedResults = ast_matchers::match(CEMatcher, *(Spec->getBody()),
DpctGlobalInfo::getContext());
for (auto &Node : MatchedResults) {
if (const auto *MatchedCE = Node.getNodeAs<CallExpr>("call")) {
SourceLocation MatchedCELocation =
SM.getSpellingLoc(MatchedCE->getBeginLoc());
auto DecomposedMatchedCELocation =
SM.getDecomposedLoc(MatchedCELocation);
if ((DecomposedCELocation.first == DecomposedMatchedCELocation.first) &&
(DecomposedCELocation.second ==
DecomposedMatchedCELocation.second)) {
return false;
}
}
}
}
return true;
}

void SyncThreadsMigrationRule::runRule(const MatchFinder::MatchResult &Result) {
static std::map<std::string, bool> LocationResultMapForTemplate;
auto emplaceReplacement = [&](BarrierFenceSpaceAnalyzerResult Res,
Expand Down Expand Up @@ -7774,7 +7806,8 @@ void SyncThreadsMigrationRule::runRule(const MatchFinder::MatchResult &Result) {
BarrierFenceSpaceAnalyzer A;
const FunctionTemplateDecl *FTD = FD->getDescribedFunctionTemplate();
if (FTD) {
if (FTD->specializations().empty()) {
if (FTD->specializations().empty() ||
noCorrespondingCEInInstantiatedTemplates(FTD, CE)) {
emplaceReplacement(A.analyze(CE), CE);
}
} else {
Expand Down
2 changes: 2 additions & 0 deletions clang/lib/DPCT/RulesLang/RulesLang.h
Original file line number Diff line number Diff line change
Expand Up @@ -809,6 +809,8 @@ class SyncThreadsMigrationRule
public:
void registerMatcher(ast_matchers::MatchFinder &MF) override;
void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
bool noCorrespondingCEInInstantiatedTemplates(const FunctionTemplateDecl *FTD,
const CallExpr *CE);
};

/// Migrate Function Attributes to Sycl kernel info, defined in
Expand Down
17 changes: 17 additions & 0 deletions clang/test/dpct/syncthreads.cu
Original file line number Diff line number Diff line change
Expand Up @@ -461,3 +461,20 @@ __global__ void test21(float *ptr1, float *ptr2, int step1, int step2) {
idx2 += step2;
}
}

template <typename T, const bool B> __device__ void test_22_d() {
// CHECK: if constexpr (B) {
// CHECK-NEXT: /*
// CHECK-NEXT: DPCT1065:{{[0-9]+}}: Consider replacing sycl::nd_item::barrier() with sycl::nd_item::barrier(sycl::access::fence_space::local_space) for better performance if there is no access to global memory.
// CHECK-NEXT: */
// CHECK-NEXT: item_ct1.barrier();
// CHECK-NEXT: }
if constexpr (B) {
__syncthreads();
}
}

__global__ void test_22() {
test_22_d<int, false>();
test_22_d<float, false>();
}