Skip to content

Commit 1243af4

Browse files
daiyaan-ahmed6daiyaan-ahmed6
committed
[SYCLomatic] Migration of cudaGraphKphKernelNodeParams, cudaGraphExecUpdateResult
Signed-off-by: Ahmed, Daiyaan <daiyaan.ahmed@intel.com>
1 parent c03712d commit 1243af4

14 files changed

Lines changed: 389 additions & 67 deletions

File tree

‎clang/lib/DPCT/ASTTraversal.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -158,6 +158,7 @@ REGISTER_RULE(TypeRemoveRule, PassKind::PK_Analysis)
158158
REGISTER_RULE(CompatWithClangRule, PassKind::PK_Migration)
159159
REGISTER_RULE(AssertRule, PassKind::PK_Migration)
160160
REGISTER_RULE(GraphRule, PassKind::PK_Migration)
161+
REGISTER_RULE(GraphAnalysisRule, PassKind::PK_Analysis)
161162
REGISTER_RULE(GraphicsInteropRule, PassKind::PK_Migration)
162163
REGISTER_RULE(RulesLangAddrSpaceConvRule, PassKind::PK_Migration)
163164

‎clang/lib/DPCT/AnalysisInfo.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2513,7 +2513,7 @@ unsigned DpctGlobalInfo::ExperimentalFlag = 0;
25132513
unsigned DpctGlobalInfo::HelperFuncPreferenceFlag = 0;
25142514
bool DpctGlobalInfo::AnalysisModeFlag = false;
25152515
bool DpctGlobalInfo::UseSYCLCompatFlag = false;
2516-
bool DpctGlobalInfo::CVersionCUDALaunchUsedFlag = false;
2516+
bool DpctGlobalInfo::UseWrapperRegisterFnPtrFlag = false;
25172517
unsigned int DpctGlobalInfo::ColorOption = 1;
25182518
std::unordered_map<int, std::shared_ptr<DeviceFunctionInfo>>
25192519
DpctGlobalInfo::CubPlaceholderIndexMap;

‎clang/lib/DPCT/AnalysisInfo.h‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1358,8 +1358,10 @@ class DpctGlobalInfo {
13581358
static bool useNoQueueDevice() {
13591359
return getHelperFuncPreference(HelperFuncPreference::NoQueueDevice);
13601360
}
1361-
static void setCVersionCUDALaunchUsed() { CVersionCUDALaunchUsedFlag = true; }
1362-
static bool isCVersionCUDALaunchUsed() { return CVersionCUDALaunchUsedFlag; }
1361+
static void setUseWrapperRegisterFnPtr() {
1362+
UseWrapperRegisterFnPtrFlag = true;
1363+
}
1364+
static bool useWrapperRegisterFnPtr() { return UseWrapperRegisterFnPtrFlag; }
13631365
static void setUseSYCLCompat(bool Flag = true) { UseSYCLCompatFlag = Flag; }
13641366
static bool useSYCLCompat() { return UseSYCLCompatFlag; }
13651367
static bool useEnqueueBarrier() {
@@ -1689,7 +1691,7 @@ class DpctGlobalInfo {
16891691
static unsigned HelperFuncPreferenceFlag;
16901692
static bool AnalysisModeFlag;
16911693
static bool UseSYCLCompatFlag;
1692-
static bool CVersionCUDALaunchUsedFlag;
1694+
static bool UseWrapperRegisterFnPtrFlag;
16931695
static unsigned int ColorOption;
16941696
static std::unordered_map<int, std::shared_ptr<DeviceFunctionInfo>>
16951697
CubPlaceholderIndexMap;

‎clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -507,6 +507,15 @@ TYPE_REWRITE_ENTRY(
507507
WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
508508
STR("--use-experimental-features=graph"))))
509509

510+
TYPE_REWRITE_ENTRY(
511+
"cudaKernelNodeParams",
512+
TYPE_CONDITIONAL_FACTORY(
513+
checkEnableGraphForType(),
514+
TYPE_FACTORY(STR(MapNames::getDpctNamespace() +
515+
"experimental::kernel_node_params")),
516+
WARNING_FACTORY(Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
517+
STR("--use-experimental-features=graph"))))
518+
510519
// Graphics Interop Handle
511520
TYPE_REWRITE_ENTRY(
512521
"cudaGraphicsResource",

‎clang/lib/DPCT/RuleInfra/MapNames.cpp‎

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -643,6 +643,14 @@ void MapNames::setExplicitNamespaceMap(
643643
DpctGlobalInfo::useExtGraph()
644644
? getClNamespace() + "ext::oneapi::experimental::node_type"
645645
: "cudaGraphNodeType")},
646+
{"cudaGraphExecUpdateResultInfo",
647+
std::make_shared<TypeNameRule>(DpctGlobalInfo::useExtGraph()
648+
? "int"
649+
: "cudaGraphExecUpdateResultInfo")},
650+
{"cudaGraphExecUpdateResult",
651+
std::make_shared<TypeNameRule>(DpctGlobalInfo::useExtGraph()
652+
? "int"
653+
: "cudaGraphExecUpdateResult")},
646654
{"CUmem_advise", std::make_shared<TypeNameRule>("int")},
647655
{"CUmemorytype",
648656
std::make_shared<TypeNameRule>(getClNamespace() + "usm::alloc")},
@@ -1154,6 +1162,47 @@ void MapNames::setExplicitNamespaceMap(
11541162
? getClNamespace() +
11551163
"ext::oneapi::experimental::node_type::empty"
11561164
: "cudaGraphNodeTypeEmpty")},
1165+
{"cudaGraphExecUpdateSuccess",
1166+
std::make_shared<EnumNameRule>(
1167+
DpctGlobalInfo::useExtGraph() ? "1" : "cudaGraphExecUpdateSuccess")},
1168+
{"cudaGraphExecUpdateError",
1169+
std::make_shared<EnumNameRule>(
1170+
DpctGlobalInfo::useExtGraph() ? "0" : "cudaGraphExecUpdateError")},
1171+
{"cudaGraphExecUpdateErrorTopologyChanged",
1172+
std::make_shared<EnumNameRule>(
1173+
DpctGlobalInfo::useExtGraph()
1174+
? "0"
1175+
: "cudaGraphExecUpdateErrorTopologyChanged")},
1176+
{"cudaGraphExecUpdateErrorNodeTypeChanged",
1177+
std::make_shared<EnumNameRule>(
1178+
DpctGlobalInfo::useExtGraph()
1179+
? "0"
1180+
: "cudaGraphExecUpdateErrorNodeTypeChanged")},
1181+
{"cudaGraphExecUpdateErrorFunctionChanged",
1182+
std::make_shared<EnumNameRule>(
1183+
DpctGlobalInfo::useExtGraph()
1184+
? "0"
1185+
: "cudaGraphExecUpdateErrorFunctionChanged")},
1186+
{"cudaGraphExecUpdateErrorParametersChanged",
1187+
std::make_shared<EnumNameRule>(
1188+
DpctGlobalInfo::useExtGraph()
1189+
? "0"
1190+
: "cudaGraphExecUpdateErrorParametersChanged")},
1191+
{"cudaGraphExecUpdateErrorNotSupported",
1192+
std::make_shared<EnumNameRule>(
1193+
DpctGlobalInfo::useExtGraph()
1194+
? "0"
1195+
: "cudaGraphExecUpdateErrorNotSupported")},
1196+
{"cudaGraphExecUpdateErrorUnsupportedFunctionChange",
1197+
std::make_shared<EnumNameRule>(
1198+
DpctGlobalInfo::useExtGraph()
1199+
? "0"
1200+
: "cudaGraphExecUpdateErrorUnsupportedFunctionChange")},
1201+
{"cudaGraphExecUpdateErrorAttributesChanged",
1202+
std::make_shared<EnumNameRule>(
1203+
DpctGlobalInfo::useExtGraph()
1204+
? "0"
1205+
: "cudaGraphExecUpdateErrorAttributesChanged")},
11571206
// enum CUmem_advise_enum
11581207
{"CU_MEM_ADVISE_SET_READ_MOSTLY", std::make_shared<EnumNameRule>("0")},
11591208
{"CU_MEM_ADVISE_UNSET_READ_MOSTLY", std::make_shared<EnumNameRule>("0")},

‎clang/lib/DPCT/RulesLang/APINamesGraph.inc‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -60,8 +60,10 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
6060

6161
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
6262
UseExtGraph,
63-
MEMBER_CALL_FACTORY_ENTRY("cudaGraphExecUpdate", ARG(0), true, "update",
64-
DEREF(1)),
63+
CALL_FACTORY_ENTRY("cudaGraphExecUpdate",
64+
CALL(MapNames::getDpctNamespace() +
65+
"experimental::update",
66+
ARG(0), ARG(1), ARG(2))),
6567
UNSUPPORT_FACTORY_ENTRY("cudaGraphExecUpdate",
6668
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
6769
ARG("cudaGraphExecUpdate"),

‎clang/lib/DPCT/RulesLang/MapNamesLang.cpp‎

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -362,5 +362,13 @@ const std::unordered_map<std::string, HelperFeatureEnum>
362362
{"sampler", HelperFeatureEnum::device_ext},
363363
};
364364

365+
// Graph kernel node params mapping
366+
MapNamesLang::MapTy GraphRule::KernelNodeParamNames{
367+
{"gridDim", "grid_dim"},
368+
{"blockDim", "block_dim"},
369+
{"kernelParams", "kernel_params"},
370+
{"sharedMemBytes", "shared_mem_bytes"},
371+
{"func", "func"}};
372+
365373
} // namespace dpct
366-
} // namespace clang
374+
} // namespace clang

‎clang/lib/DPCT/RulesLang/RulesLang.cpp‎

Lines changed: 36 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -348,23 +348,25 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
348348
"cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
349349
"cudaExternalSemaphoreHandleType", "CUstreamCallback",
350350
"cudaHostFn_t", "__nv_half2", "__nv_half", "cudaGraphNodeType",
351-
"CUsurfref", "CUdevice_P2PAttribute", "cudaIpcMemHandle_t"))))))
351+
"CUsurfref", "CUdevice_P2PAttribute", "cudaIpcMemHandle_t",
352+
"cudaGraphExecUpdateResultInfo"))))))
352353
.bind("cudaTypeDef"),
353354
this);
354355

355356
MF.addMatcher(
356-
typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName(
357-
"cooperative_groups::__v1::coalesced_group",
358-
"cooperative_groups::__v1::grid_group",
359-
"cooperative_groups::__v1::thread_block_tile", "cudaGraph_t",
360-
"cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource",
361-
"cudaGraphicsResource_t", "CUgraphicsResource",
362-
"cudaExternalMemory_t", "cudaExternalMemoryHandleDesc",
363-
"cudaExternalMemoryMipmappedArrayDesc",
364-
"cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t",
365-
"cudaExternalSemaphoreHandleDesc",
366-
"cudaExternalSemaphoreSignalParams",
367-
"cudaExternalSemaphoreWaitParams"))))))
357+
typeLoc(
358+
loc(qualType(hasDeclaration(namedDecl(hasAnyName(
359+
"cooperative_groups::__v1::coalesced_group",
360+
"cooperative_groups::__v1::grid_group",
361+
"cooperative_groups::__v1::thread_block_tile", "cudaGraph_t",
362+
"cudaGraphExec_t", "cudaGraphNode_t", "cudaGraphicsResource",
363+
"cudaGraphicsResource_t", "CUgraphicsResource",
364+
"cudaExternalMemory_t", "cudaExternalMemoryHandleDesc",
365+
"cudaExternalMemoryMipmappedArrayDesc",
366+
"cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t",
367+
"cudaExternalSemaphoreHandleDesc",
368+
"cudaExternalSemaphoreSignalParams",
369+
"cudaExternalSemaphoreWaitParams", "cudaKernelNodeParams"))))))
368370
.bind("cudaTypeDefEA"),
369371
this);
370372
MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl(
@@ -937,9 +939,11 @@ void TypeInDeclRule::runRule(const MatchFinder::MatchResult &Result) {
937939
}
938940

939941
if (CanonicalTypeStr == "cudaGraphExecUpdateResult") {
940-
report(TL->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false,
941-
CanonicalTypeStr);
942-
return;
942+
if (!DpctGlobalInfo::useExtGraph()) {
943+
report(TL->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
944+
"cudaGraphExecUpdateResult",
945+
"--use-experimental-features=graph");
946+
}
943947
}
944948

945949
if (CanonicalTypeStr == "cudaGraphicsRegisterFlags" ||
@@ -1941,7 +1945,8 @@ void EnumConstantRule::registerMatcher(MatchFinder &MF) {
19411945
"cufftType", "cudaMemoryType", "CUctx_flags_enum",
19421946
"CUpointer_attribute_enum", "CUmemorytype_enum",
19431947
"cudaGraphicsMapFlags", "cudaGraphicsRegisterFlags",
1944-
"cudaGraphNodeType", "CUdevice_P2PAttribute_enum"))),
1948+
"cudaGraphNodeType", "CUdevice_P2PAttribute_enum",
1949+
"cudaGraphExecUpdateResult"))),
19451950
matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*")))))
19461951
.bind("EnumConstant"),
19471952
this);
@@ -2061,7 +2066,16 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) {
20612066
EnumName == "cudaGraphNodeTypeMemset" ||
20622067
EnumName == "cudaGraphNodeTypeHost" ||
20632068
EnumName == "cudaGraphNodeTypeGraph" ||
2064-
EnumName == "cudaGraphNodeTypeEmpty")) {
2069+
EnumName == "cudaGraphNodeTypeEmpty" ||
2070+
EnumName == "cudaGraphExecUpdateSuccess" ||
2071+
EnumName == "cudaGraphExecUpdateError" ||
2072+
EnumName == "cudaGraphExecUpdateErrorTopologyChanged" ||
2073+
EnumName == "cudaGraphExecUpdateErrorNodeTypeChanged" ||
2074+
EnumName == "cudaGraphExecUpdateErrorFunctionChanged" ||
2075+
EnumName == "cudaGraphExecUpdateErrorParametersChanged" ||
2076+
EnumName == "cudaGraphExecUpdateErrorNotSupported" ||
2077+
EnumName == "cudaGraphExecUpdateErrorUnsupportedFunctionChange" ||
2078+
EnumName == "cudaGraphExecUpdateErrorAttributesChanged")) {
20652079
report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
20662080
EnumName, "--use-experimental-features=graph");
20672081
return;
@@ -4638,7 +4652,7 @@ void KernelCallRefRule::runRule(
46384652
(OuterFD->getTemplatedKind() ==
46394653
FunctionDecl::TemplatedKind::TK_FunctionTemplate)) {
46404654
std::string TypeRepl;
4641-
if (DpctGlobalInfo::isCVersionCUDALaunchUsed()) {
4655+
if (DpctGlobalInfo::useWrapperRegisterFnPtr()) {
46424656
if ((IsTemplateRelated &&
46434657
(!DRE->hasExplicitTemplateArgs() ||
46444658
(DRE->getNumTemplateArgs() <= TemplateParamNum))) ||
@@ -4647,7 +4661,7 @@ void KernelCallRefRule::runRule(
46474661
}
46484662
}
46494663
insertWrapperPostfix<DeclRefExpr>(
4650-
DRE, std::move(TypeRepl), DpctGlobalInfo::isCVersionCUDALaunchUsed());
4664+
DRE, std::move(TypeRepl), DpctGlobalInfo::useWrapperRegisterFnPtr());
46514665
}
46524666
}
46534667
if (auto ULE =
@@ -4684,7 +4698,7 @@ void KernelCallRefRule::runRule(
46844698
}
46854699
}
46864700
insertWrapperPostfix<UnresolvedLookupExpr>(
4687-
ULE, getTypeRepl(ULE), DpctGlobalInfo::isCVersionCUDALaunchUsed());
4701+
ULE, getTypeRepl(ULE), DpctGlobalInfo::useWrapperRegisterFnPtr());
46884702
}
46894703
}
46904704

@@ -4957,7 +4971,7 @@ void KernelCallRule::runRule(
49574971

49584972
if (!getAddressedRef(CalleeDRE)) {
49594973
if (IsFuncTypeErased) {
4960-
DpctGlobalInfo::setCVersionCUDALaunchUsed();
4974+
DpctGlobalInfo::setUseWrapperRegisterFnPtr();
49614975
}
49624976
std::string ReplStr;
49634977
llvm::raw_string_ostream OS(ReplStr);

‎clang/lib/DPCT/RulesLang/RulesLang.h‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -998,7 +998,17 @@ class CompatWithClangRule : public NamedMigrationRule<CompatWithClangRule> {
998998
void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
999999
};
10001000

1001+
class GraphAnalysisRule : public NamedMigrationRule<GraphAnalysisRule> {
1002+
public:
1003+
void registerMatcher(ast_matchers::MatchFinder &MF) override;
1004+
void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
1005+
};
1006+
10011007
class GraphRule : public NamedMigrationRule<GraphRule> {
1008+
static MapNames::MapTy KernelNodeParamNames;
1009+
const Expr *getAssignedBO(const Expr *E, ASTContext &Context);
1010+
const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context);
1011+
10021012
public:
10031013
void registerMatcher(ast_matchers::MatchFinder &MF) override;
10041014
void runRule(const ast_matchers::MatchFinder::MatchResult &Result);

0 commit comments

Comments
 (0)