Skip to content

[SYCLomatic] Support migration for cub::{StoreDirectBlocked, StoreDirectStriped} API #1305

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

Closed
wants to merge 23 commits into from
Closed

[SYCLomatic] Support migration for cub::{StoreDirectBlocked, StoreDirectStriped} API #1305

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
23 commits
Select commit Hold shift + click to select a range
d842dfe
cub::blockload migration
abhilash1910 Sep 15, 2023
c014557
add blockload + fix
abhilash1910 Sep 18, 2023
aadbb89
add block load sample
abhilash1910 Sep 18, 2023
3e83abb
add line
abhilash1910 Sep 18, 2023
c856837
add prototype kernel api
abhilash1910 Sep 19, 2023
8b66f8e
fix bugs
abhilash1910 Sep 20, 2023
8b5589a
fix bugs
abhilash1910 Sep 20, 2023
7841813
Merge branch 'oneapi-src:SYCLomatic' into blockload_support
abhilash1910 Sep 26, 2023
ea33c67
Merge branch 'oneapi-src:SYCLomatic' into blockload_support
abhilash1910 Oct 3, 2023
f190341
Merge branch 'oneapi-src:SYCLomatic' into blockload_support
abhilash1910 Jul 11, 2024
d0a0e8b
modify to store
abhilash1910 Jul 11, 2024
b336ff2
add store rewriter
abhilash1910 Jul 11, 2024
2550602
revert changes
abhilash1910 Jul 11, 2024
7efbc81
revert changes
abhilash1910 Jul 11, 2024
5f7d10b
fix issue
abhilash1910 Jul 11, 2024
3b0494f
remove manual parsing
abhilash1910 Jul 11, 2024
d2ad0e5
Update APINames_CUB.inc
abhilash1910 Jul 11, 2024
9993f85
fix test
abhilash1910 Jul 11, 2024
a064b46
format fix
abhilash1910 Aug 8, 2024
23735aa
add newline
abhilash1910 Aug 8, 2024
5fc94c6
fix include file
abhilash1910 Aug 13, 2024
fd9c641
fix blank line
abhilash1910 Aug 13, 2024
ee15a11
Merge branch 'SYCLomatic' into blockload_support
abhilash1910 Aug 22, 2024
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
2 changes: 1 addition & 1 deletion clang/lib/DPCT/APINames_CUB.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -107,7 +107,6 @@ ENTRY_MEMBER_FUNCTION(cub::BlockExchange, cub::BlockExchange, ScatterToStripedFl
ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, InitHistogram, InitHistogram, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Histogram, Histogram, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Composite, Composite, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockStore, cub::BlockStore, Store, Store, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, Sort, Sort, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, SortDescending, SortDescending, false, NO_FLAG, P4, "Comment")
Expand All @@ -123,6 +122,7 @@ ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Offset, Offset, fals
ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Rotate, Rotate, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Up, Up, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Down, Down, false, NO_FLAG, P4, "Comment")
ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, true, NO_FLAG, P4, "Comment")


Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let's remove extra blank line change.

// Device Level
Expand Down
16 changes: 10 additions & 6 deletions clang/lib/DPCT/CUBAPIMigration.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -558,14 +558,14 @@ void CubDeviceLevelRule::removeRedundantTempVar(const CallExpr *CE) {
void CubRule::registerMatcher(ast_matchers::MatchFinder &MF) {
MF.addMatcher(
typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName(
"WarpScan", "WarpReduce", "BlockScan", "BlockReduce"))))))
"WarpScan", "WarpReduce", "BlockScan", "BlockReduce", "BlockLoad"))))))
.bind("TypeLoc"),
this);

MF.addMatcher(
typedefDecl(
hasType(hasCanonicalType(qualType(hasDeclaration(namedDecl(hasAnyName(
"WarpScan", "WarpReduce", "BlockScan", "BlockReduce")))))))
"WarpScan", "WarpReduce", "BlockScan", "BlockReduce", "BlockLoad")))))))
.bind("TypeDefDecl"),
this);

Expand Down Expand Up @@ -684,7 +684,8 @@ void CubRule::processCubDeclStmt(const DeclStmt *DS) {
ObjTypeStr.find("class cub::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("class cub::BlockLoad") == 0) {
Repl = DpctGlobalInfo::getGroup(DRE);
} else {
continue;
Expand Down Expand Up @@ -749,7 +750,8 @@ void CubRule::processCubTypeDef(const TypedefDecl *TD) {
!(ObjTypeStr.find("class cub::WarpScan") == 0 ||
ObjTypeStr.find("class cub::WarpReduce") == 0 ||
ObjTypeStr.find("class cub::BlockScan") == 0 ||
ObjTypeStr.find("class cub::BlockReduce") == 0)) {
ObjTypeStr.find("class cub::BlockReduce") == 0 ||
ObjTypeStr.find("class cub::BlockLoad") == 0)) {
DeleteFlag = false;
break;
}
Expand Down Expand Up @@ -1304,7 +1306,8 @@ void CubRule::processCubMemberCall(const CXXMemberCallExpr *MC) {
ObjTypeStr.find("class cub::WarpReduce") == 0) {
processWarpLevelMemberCall(MC);
} else if (ObjTypeStr.find("class cub::BlockScan") == 0 ||
ObjTypeStr.find("class cub::BlockReduce") == 0) {
ObjTypeStr.find("class cub::BlockReduce") == 0 ||
ObjTypeStr.find("class cub::BlockLoad") == 0) {
processBlockLevelMemberCall(MC);
} else {
report(MC->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, ObjTypeStr);
Expand All @@ -1328,7 +1331,8 @@ void CubRule::processTypeLoc(const TypeLoc *TL) {
MapNames::getClNamespace() + "sub_group",
SM));
} else if (TypeName.find("class cub::BlockScan") == 0 ||
TypeName.find("class cub::BlockReduce") == 0) {
TypeName.find("class cub::BlockReduce") == 0 ||
ObjTypeStr.find("class cub::BlockLoad") == 0) {
auto DeviceFuncDecl = DpctGlobalInfo::findAncestor<FunctionDecl>(TL);
if (DeviceFuncDecl && (DeviceFuncDecl->hasAttr<CUDADeviceAttr>() ||
DeviceFuncDecl->hasAttr<CUDAGlobalAttr>())) {
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/DPCT/ExprAnalysis.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1170,7 +1170,7 @@ void ExprAnalysis::analyzeType(TypeLoc TL, const Expr *CSCE,
}
}
if (OS.str() != "cub::WarpScan" && OS.str() != "cub::WarpReduce" &&
OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan") {
OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan" && OS.str() != "cub::BlockLoad") {
SR.setEnd(TSTL.getTemplateNameLoc());
}
analyzeTemplateSpecializationType(TSTL);
Expand Down