Skip to content

[SYCLomatic] Match the typeAliasDecl for Cub Type, match template scan and reduce for block and warp. #2887

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
zhimingwang36 merged 5 commits into from
Jun 13, 2025
Merged

[SYCLomatic] Match the typeAliasDecl for Cub Type, match template scan and reduce for block and warp. #2887

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
bb8285a
init
ShengchenJ Jun 10, 2025
0f2e2f9
[SYCLomatic] Match the AliasTypeDecl for Cub Type.
ShengchenJ Jun 11, 2025
497e8fe
up
ShengchenJ Jun 11, 2025
ce0babe
up
ShengchenJ Jun 11, 2025
3fa1d56
Merge remote-tracking branch 'origin/SYCLomatic' into fix_cub
ShengchenJ Jun 12, 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
18 changes: 13 additions & 5 deletions clang/lib/DPCT/RulesLangLib/CUBAPIMigration.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -946,10 +946,14 @@ void CubRule::processCubDeclStmt(const DeclStmt *DS) {
if (isTypeInAnalysisScope(ObjCanonicalType.getTypePtr())) {
continue;
} else if (ObjTypeStr.find("class cub::WarpScan") == 0 ||
ObjTypeStr.find("class cub::WarpReduce") == 0) {
ObjTypeStr.find("class cub::WarpReduce") == 0 ||
ObjTypeStr.find("WarpScan<") == 0 ||
ObjTypeStr.find("WarpReduce<") == 0) {
Repl = DpctGlobalInfo::getSubGroup(DRE);
} else if (ObjTypeStr.find("class cub::BlockScan") == 0 ||
ObjTypeStr.find("class cub::BlockReduce") == 0) {
ObjTypeStr.find("class cub::BlockReduce") == 0 ||
ObjTypeStr.find("BlockScan<") == 0 ||
ObjTypeStr.find("BlockReduce<") == 0) {
Repl = DpctGlobalInfo::getGroup(DRE);
} else {
continue;
Expand All @@ -975,9 +979,12 @@ void CubRule::processCubTypeDefOrUsing(const TypedefNameDecl *TD) {
std::string TypeName = TD->getNameAsString();
auto &Context = dpct::DpctGlobalInfo::getContext();
auto &SM = dpct::DpctGlobalInfo::getSourceManager();
auto MyMatcher = compoundStmt(forEachDescendant(
typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasName(TypeName))))))
.bind("typeLoc")));

auto MyMatcher = compoundStmt(
forEachDescendant(typeLoc(loc(qualType(hasDeclaration(
anyOf(typedefDecl(hasName(TypeName)),
typeAliasDecl(hasName(TypeName)))))))
.bind("typeLoc")));
auto MatcherScope = DpctGlobalInfo::findAncestor<CompoundStmt>(TD);
if (!MatcherScope)
return;
Expand Down Expand Up @@ -1033,6 +1040,7 @@ void CubRule::processCubTypeDefOrUsing(const TypedefNameDecl *TD) {
}
}
}

if (DeleteFlag) {
emplaceTransformation(new ReplaceDecl(TD, ""));
} else {
Expand Down
40 changes: 40 additions & 0 deletions clang/test/dpct/cub/type/fix_alias_bug.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,40 @@
// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
// RUN: dpct --format-range=none -in-root %S -out-root %T/type/fix_alias_bug %S/fix_alias_bug.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
// RUN: FileCheck --input-file %T/type/fix_alias_bug/fix_alias_bug.dp.cpp %s
// RUN: %if build_lit %{icpx -c -fsycl %T/type/fix_alias_bug/fix_alias_bug.dp.cpp -o %T/type/fix_alias_bug/fix_alias_bug.dp.o %}

#include <cub/cub.cuh>
#include <cuda.h>
template <bool> struct warp_reduce;

template <> struct warp_reduce<true> {
template <typename T, typename reducer_t>
inline T operator()(const T &value_, bool all, const reducer_t &r) {
// CHECK: using warp_reduce_t = sycl::sub_group;

using warp_reduce_t = cub::WarpReduce<T, 10>;

typename warp_reduce_t::TempStorage dummy_storage;
// CHECK: warp_reduce_t warp_reduce(sycl::ext::oneapi::this_work_item::get_sub_group());

warp_reduce_t warp_reduce(dummy_storage);
}
};

template <> struct warp_reduce<false> {
template <typename T, typename reducer_t>
inline T operator()(const T &value_, bool all, const reducer_t &r) {
// CHECK: using warp_reduce_t = sycl::sub_group;

using warp_reduce_t = cub::WarpReduce<T, 10>;

typename warp_reduce_t::TempStorage dummy_storage;
// CHECK: warp_reduce_t warp_reduce(sycl::ext::oneapi::this_work_item::get_sub_group());
warp_reduce_t warp_reduce(dummy_storage);
}
};
int main() {
int a, b;
warp_reduce<true>()(a, true, b);
}