Skip to content

[SYCLomatic] Fix migration issue for free-function-queries #2804

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.

Sign up for GitHub

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

Merged
ziranzha merged 6 commits into from
Apr 30, 2025
Merged

[SYCLomatic] Fix migration issue for free-function-queries #2804

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
1ece36c
[SYCLomatic] Fix migration issue for free-function-queries
zhiweij1 Apr 24, 2025
aec62ae
Refine
zhiweij1 Apr 28, 2025
3f415dc
Merge remote-tracking branch 'origin/SYCLomatic' into free_func
zhiweij1 Apr 28, 2025
22acd3b
Fix
zhiweij1 Apr 28, 2025
085ca92
Fix
zhiweij1 Apr 29, 2025
df11eb8
Merge remote-tracking branch 'origin/SYCLomatic' into free_func
zhiweij1 Apr 29, 2025
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
31 changes: 26 additions & 5 deletions clang/lib/DPCT/AnalysisInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7512,11 +7512,32 @@ void FreeQueriesInfo::printImmediateText(llvm::raw_ostream &OS, const Node *S,
return Info->printImmediateText(OS, S->getBeginLoc(), K);
}

#ifdef DPCT_DEBUG_BUILD
llvm::errs() << "Can not get FreeQueriesInfo for this FunctionDecl\n";
assert(0);
#endif // DPCT_DEBUG_BUILD

auto DFI = DeviceFunctionDecl::LinkRedecls(FD);
if (!DFI)
return;
auto Index = DpctGlobalInfo::getCudaKernelDimDFIIndexThenInc();
DpctGlobalInfo::insertCudaKernelDimDFIMap(Index, DFI);
switch (K) {
case FreeQueriesKind::NdItem: {
OS << MapNames::getClNamespace()
<< "ext::oneapi::this_work_item::get_nd_item<{{NEEDREPLACEG" +
std::to_string(Index) + "}}>()";
break;
}
case FreeQueriesKind::Group: {
OS << MapNames::getClNamespace()
<< "ext::oneapi::this_work_item::get_work_group<{{NEEDREPLACEG" +
std::to_string(Index) + "}}>()";
break;
}
case FreeQueriesKind::SubGroup: {
OS << MapNames::getClNamespace()
<< "ext::oneapi::this_work_item::get_sub_group()";
break;
}
default:
llvm_unreachable("Unexpected FreeQueriesKind");
}
Comment thread
zhiweij1 marked this conversation as resolved.
Outdated
} else {
if (auto DFI = DeviceFunctionDecl::LinkRedecls(FD))
DFI->setItem();
Expand Down
13 changes: 13 additions & 0 deletions clang/test/dpct/with_this_nd_item.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,3 +87,16 @@ int main() {

test3<<<32,32>>>();
}

struct TEST4 {
// CHECK: TEST4(int a) : thread_idx(sycl::ext::oneapi::this_work_item::get_nd_item<1>().get_local_id(0)) {}
__device__ TEST4(int a) : thread_idx(threadIdx.x) {}
int thread_idx;
};

// CHECK: int test5(const int ct, int numLane = 0) {
// CHECK-NEXT: if (!numLane) numLane = sycl::ext::oneapi::this_work_item::get_sub_group().get_local_range().get(0);
__device__ int test5(const int ct, const int numLane = warpSize) {
int r = ct * numLane;
return r;
}