Skip to content

Commit 0d4fb62

Browse files
daiyaan-ahmed6daiyaan-ahmed6
committed
Migration for cudaGraphExecUpdateResultInfo
Signed-off-by: Daiyaan Ahmed <[email protected]>
1 parent 164819e commit 0d4fb62

File tree

7 files changed

+211
-39
lines changed

7 files changed

+211
-39
lines changed

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

Lines changed: 49 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -640,8 +640,14 @@ void MapNames::setExplicitNamespaceMap(
640640
DpctGlobalInfo::useExtGraph()
641641
? getClNamespace() + "ext::oneapi::experimental::node_type"
642642
: "cudaGraphNodeType")},
643-
{"cudaGraphExecUpdateResultInfo", std::make_shared<TypeNameRule>("int")},
644-
{"cudaGraphExecUpdateResult", std::make_shared<TypeNameRule>("int")},
643+
{"cudaGraphExecUpdateResultInfo",
644+
std::make_shared<TypeNameRule>(DpctGlobalInfo::useExtGraph()
645+
? "int"
646+
: "cudaGraphExecUpdateResultInfo")},
647+
{"cudaGraphExecUpdateResult",
648+
std::make_shared<TypeNameRule>(DpctGlobalInfo::useExtGraph()
649+
? "int"
650+
: "cudaGraphExecUpdateResultInfo")},
645651
{"CUmem_advise", std::make_shared<TypeNameRule>("int")},
646652
{"CUmemorytype",
647653
std::make_shared<TypeNameRule>(getClNamespace() + "usm::alloc")},
@@ -1148,6 +1154,47 @@ void MapNames::setExplicitNamespaceMap(
11481154
? getClNamespace() +
11491155
"ext::oneapi::experimental::node_type::empty"
11501156
: "cudaGraphNodeTypeEmpty")},
1157+
{"cudaGraphExecUpdateSuccess",
1158+
std::make_shared<EnumNameRule>(
1159+
DpctGlobalInfo::useExtGraph() ? "1" : "cudaGraphExecUpdateSuccess")},
1160+
{"cudaGraphExecUpdateError",
1161+
std::make_shared<EnumNameRule>(
1162+
DpctGlobalInfo::useExtGraph() ? "0" : "cudaGraphExecUpdateError")},
1163+
{"cudaGraphExecUpdateErrorTopologyChanged",
1164+
std::make_shared<EnumNameRule>(
1165+
DpctGlobalInfo::useExtGraph()
1166+
? "0"
1167+
: "cudaGraphExecUpdateErrorTopologyChanged")},
1168+
{"cudaGraphExecUpdateErrorNodeTypeChanged",
1169+
std::make_shared<EnumNameRule>(
1170+
DpctGlobalInfo::useExtGraph()
1171+
? "0"
1172+
: "cudaGraphExecUpdateErrorNodeTypeChanged")},
1173+
{"cudaGraphExecUpdateErrorFunctionChanged",
1174+
std::make_shared<EnumNameRule>(
1175+
DpctGlobalInfo::useExtGraph()
1176+
? "0"
1177+
: "cudaGraphExecUpdateErrorFunctionChanged")},
1178+
{"cudaGraphExecUpdateErrorParametersChanged",
1179+
std::make_shared<EnumNameRule>(
1180+
DpctGlobalInfo::useExtGraph()
1181+
? "0"
1182+
: "cudaGraphExecUpdateErrorParametersChanged")},
1183+
{"cudaGraphExecUpdateErrorNotSupported",
1184+
std::make_shared<EnumNameRule>(
1185+
DpctGlobalInfo::useExtGraph()
1186+
? "0"
1187+
: "cudaGraphExecUpdateErrorNotSupported")},
1188+
{"cudaGraphExecUpdateErrorUnsupportedFunctionChange",
1189+
std::make_shared<EnumNameRule>(
1190+
DpctGlobalInfo::useExtGraph()
1191+
? "0"
1192+
: "cudaGraphExecUpdateErrorUnsupportedFunctionChange")},
1193+
{"cudaGraphExecUpdateErrorAttributesChanged",
1194+
std::make_shared<EnumNameRule>(
1195+
DpctGlobalInfo::useExtGraph()
1196+
? "0"
1197+
: "cudaGraphExecUpdateErrorAttributesChanged")},
11511198
// enum CUmem_advise_enum
11521199
{"CU_MEM_ADVISE_SET_READ_MOSTLY", std::make_shared<EnumNameRule>("0")},
11531200
{"CU_MEM_ADVISE_UNSET_READ_MOSTLY", std::make_shared<EnumNameRule>("0")},

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

Lines changed: 32 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -58,22 +58,25 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
5858

5959
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
6060
UseExtGraph,
61-
MEMBER_CALL_FACTORY_ENTRY("cudaGraphExecUpdate", ARG(0), true, "update",
62-
DEREF(1)),
61+
CALL_FACTORY_ENTRY("cudaGraphExecUpdate",
62+
CALL(MapNames::getDpctNamespace() +
63+
"experimental::update",
64+
ARG(0), ARG(1), ARG(2))),
6365
UNSUPPORT_FACTORY_ENTRY("cudaGraphExecUpdate",
6466
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
6567
ARG("cudaGraphExecUpdate"),
6668
ARG("--use-experimental-features=graph"))))
6769

68-
CONDITIONAL_FACTORY_ENTRY(
70+
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
6971
UseExtGraph,
70-
ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY("cudaGraphNodeGetType", DEREF(1),
71-
MEMBER_CALL(ARG(0), true,
72-
"get_type"))),
72+
CALL_FACTORY_ENTRY("cudaGraphNodeGetType",
73+
CALL(MapNames::getDpctNamespace() +
74+
"experimental::get_node_type",
75+
ARG(0), ARG(1))),
7376
UNSUPPORT_FACTORY_ENTRY("cudaGraphNodeGetType",
7477
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
7578
ARG("cudaGraphNodeGetType"),
76-
ARG("--use-experimental-features=graph")))
79+
ARG("--use-experimental-features=graph"))))
7780

7881
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
7982
UseExtGraph,
@@ -113,3 +116,25 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
113116
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
114117
ARG("cudaGraphAddKernelNode"),
115118
ARG("--use-experimental-features=graph"))))
119+
120+
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
121+
UseExtGraph,
122+
CALL_FACTORY_ENTRY("cudaGraphKernelNodeGetParams",
123+
CALL(MapNames::getDpctNamespace() +
124+
"experimental::kernel_node_get_params",
125+
ARG(0), ARG(1))),
126+
UNSUPPORT_FACTORY_ENTRY("cudaGraphKernelNodeGetParams",
127+
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
128+
ARG("cudaGraphKernelNodeGetParams"),
129+
ARG("--use-experimental-features=graph"))))
130+
131+
ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
132+
UseExtGraph,
133+
CALL_FACTORY_ENTRY("cudaGraphKernelNodeSetParams",
134+
CALL(MapNames::getDpctNamespace() +
135+
"experimental::kernel_node_set_params",
136+
ARG(0), ARG(1))),
137+
UNSUPPORT_FACTORY_ENTRY("cudaGraphKernelNodeSetParams",
138+
Diagnostics::TRY_EXPERIMENTAL_FEATURE,
139+
ARG("cudaGraphKernelNodeSetParams"),
140+
ARG("--use-experimental-features=graph"))))

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

Lines changed: 11 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1922,7 +1922,7 @@ void EnumConstantRule::registerMatcher(MatchFinder &MF) {
19221922
"cufftType", "cudaMemoryType", "CUctx_flags_enum",
19231923
"CUpointer_attribute_enum", "CUmemorytype_enum",
19241924
"cudaGraphicsMapFlags", "cudaGraphicsRegisterFlags",
1925-
"cudaGraphNodeType"))),
1925+
"cudaGraphNodeType", "cudaGraphExecUpdateResult"))),
19261926
matchesName("CUDNN_.*"), matchesName("CUSOLVER_.*")))))
19271927
.bind("EnumConstant"),
19281928
this);
@@ -2041,7 +2041,16 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) {
20412041
EnumName == "cudaGraphNodeTypeMemset" ||
20422042
EnumName == "cudaGraphNodeTypeHost" ||
20432043
EnumName == "cudaGraphNodeTypeGraph" ||
2044-
EnumName == "cudaGraphNodeTypeEmpty")) {
2044+
EnumName == "cudaGraphNodeTypeEmpty" ||
2045+
EnumName == "cudaGraphExecUpdateSuccess" ||
2046+
EnumName == "cudaGraphExecUpdateError" ||
2047+
EnumName == "cudaGraphExecUpdateErrorTopologyChanged" ||
2048+
EnumName == "cudaGraphExecUpdateErrorNodeTypeChanged" ||
2049+
EnumName == "cudaGraphExecUpdateErrorFunctionChanged" ||
2050+
EnumName == "cudaGraphExecUpdateErrorParametersChanged" ||
2051+
EnumName == "cudaGraphExecUpdateErrorNotSupported" ||
2052+
EnumName == "cudaGraphExecUpdateErrorUnsupportedFunctionChange" ||
2053+
EnumName == "cudaGraphExecUpdateErrorAttributesChanged")) {
20452054
report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
20462055
EnumName, "--use-experimental-features=graph");
20472056
return;

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

Lines changed: 35 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -58,7 +58,8 @@ void GraphRule::registerMatcher(MatchFinder &MF) {
5858
"cudaGraphInstantiate", "cudaGraphLaunch", "cudaGraphExecDestroy",
5959
"cudaGraphAddEmptyNode", "cudaGraphAddDependencies",
6060
"cudaGraphExecUpdate", "cudaGraphNodeGetType", "cudaGraphGetNodes",
61-
"cudaGraphGetRootNodes", "cudaGraphDestroy", "cudaGraphAddKernelNode");
61+
"cudaGraphGetRootNodes", "cudaGraphDestroy", "cudaGraphAddKernelNode",
62+
"cudaGraphKernelNodeGetParams", "cudaGraphKernelNodeSetParams");
6263
};
6364
MF.addMatcher(
6465
callExpr(callee(functionDecl(functionName()))).bind("FunctionCall"),
@@ -70,6 +71,12 @@ void GraphRule::registerMatcher(MatchFinder &MF) {
7071
recordType(hasDeclaration(recordDecl(typeName()))))))))
7172
.bind("Type"),
7273
this);
74+
75+
MF.addMatcher(memberExpr(hasObjectExpression(hasType(
76+
asString("cudaGraphExecUpdateResultInfo"))),
77+
member(hasName("result")))
78+
.bind("execUpdateResult"),
79+
this);
7380
}
7481

7582
void GraphRule::runRule(const MatchFinder::MatchResult &Result) {
@@ -100,10 +107,16 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) {
100107
if (auto *FD = dyn_cast<FunctionDecl>(RHS_DRE->getDecl())) {
101108
std::string FuncName = FD->getNameAsString();
102109
std::string WrapperName = FuncName;
103-
std::string AccessOperator = VD->getType()->isPointerType() ? "->" : ".";
104-
std::string ReplacementStr = VarName + AccessOperator + "set_func("
105-
"(void*) dpct::wrapper_register(&" + WrapperName ;
106-
emplaceTransformation(new ReplaceToken(BO->getBeginLoc(), BO->getEndLoc(), std::move(ReplacementStr)));
110+
std::string AccessOperator =
111+
VD->getType()->isPointerType() ? "->" : ".";
112+
std::string ReplacementStr =
113+
VarName + AccessOperator +
114+
"set_func("
115+
"(void*) dpct::wrapper_register(&" +
116+
WrapperName;
117+
emplaceTransformation(
118+
new ReplaceToken(BO->getBeginLoc(), BO->getEndLoc(),
119+
std::move(ReplacementStr)));
107120
emplaceTransformation(new InsertAfterStmt(BO, ")"));
108121
return;
109122
}
@@ -124,6 +137,23 @@ void GraphRule::runRule(const MatchFinder::MatchResult &Result) {
124137
}
125138
return;
126139
}
140+
if (auto ME = getNodeAsType<MemberExpr>(Result, "execUpdateResult")) {
141+
auto MD = ME->getMemberDecl();
142+
const Expr *Base = ME->getBase();
143+
if (MD->getNameAsString() == "result") {
144+
if (auto *DRE = dyn_cast<DeclRefExpr>(Base)) {
145+
SourceLocation StartLoc = Base->getBeginLoc();
146+
SourceLocation EndLoc = ME->getEndLoc();
147+
const SourceManager &SM = *Result.SourceManager;
148+
EndLoc = Lexer::getLocForEndOfToken(EndLoc, 0, SM,
149+
Result.Context->getLangOpts());
150+
std::string VarNameStr = DRE->getNameInfo().getAsString();
151+
emplaceTransformation(
152+
new ReplaceToken(StartLoc, EndLoc, std::move(VarNameStr)));
153+
}
154+
}
155+
return;
156+
}
127157
const CallExpr *CE = getNodeAsType<CallExpr>(Result, "FunctionCall");
128158
if (!CE) {
129159
return;

‎clang/lib/DPCT/SrcAPI/APINames.inc‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -456,12 +456,12 @@ ENTRY(cudaGraphInstantiateWithFlags, cudaGraphInstantiateWithFlags, false, NO_FL
456456
ENTRY(cudaGraphInstantiateWithParams, cudaGraphInstantiateWithParams, false, NO_FLAG, P4, "comment")
457457
ENTRY(cudaGraphKernelNodeCopyAttributes, cudaGraphKernelNodeCopyAttributes, false, NO_FLAG, P4, "comment")
458458
ENTRY(cudaGraphKernelNodeGetAttribute, cudaGraphKernelNodeGetAttribute, false, NO_FLAG, P4, "comment")
459-
ENTRY(cudaGraphKernelNodeGetParams, cudaGraphKernelNodeGetParams, false, NO_FLAG, P4, "comment")
459+
ENTRY(cudaGraphKernelNodeGetParams, cudaGraphKernelNodeGetParams, true, NO_FLAG, P4, "Successful/DPCT1119")
460460
ENTRY(cudaGraphKernelNodeSetAttribute, cudaGraphKernelNodeSetAttribute, false, NO_FLAG, P4, "comment")
461461
ENTRY(cudaGraphKernelNodeSetEnabled, cudaGraphKernelNodeSetEnabled, false, NO_FLAG, P4, "comment")
462462
ENTRY(cudaGraphKernelNodeSetGridDim, cudaGraphKernelNodeSetGridDim, false, NO_FLAG, P4, "comment")
463463
ENTRY(cudaGraphKernelNodeSetParam, cudaGraphKernelNodeSetParam, false, NO_FLAG, P4, "comment")
464-
ENTRY(cudaGraphKernelNodeSetParams, cudaGraphKernelNodeSetParams, false, NO_FLAG, P4, "comment")
464+
ENTRY(cudaGraphKernelNodeSetParams, cudaGraphKernelNodeSetParams, true, NO_FLAG, P4, "Successful/DPCT1119")
465465
ENTRY(cudaGraphKernelNodeUpdatesApply, cudaGraphKernelNodeUpdatesApply, false, NO_FLAG, P4, "comment")
466466
ENTRY(cudaGraphLaunch, cudaGraphLaunch, true, NO_FLAG, P4, "Successful/DPCT1119")
467467
ENTRY(cudaGraphMemAllocNodeGetParams, cudaGraphMemAllocNodeGetParams, false, NO_FLAG, P4, "comment")

‎clang/runtime/dpct-rt/include/dpct/graph.hpp‎

Lines changed: 62 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,12 @@ struct kernel_node_params {
3333
unsigned int shared_mem_bytes{};
3434

3535
std::vector<dpct::experimental::node_ptr> dependencies{};
36+
kernel_node_params() = default;
37+
kernel_node_params(const kernel_node_params &other)
38+
: block_dim(other.block_dim), grid_dim(other.grid_dim),
39+
kernel_params(other.kernel_params), func(other.func),
40+
shared_mem_bytes(other.shared_mem_bytes),
41+
dependencies(other.dependencies) {}
3642

3743
public:
3844
void set_block_dim(const dpct::dim3 &block_dim) {
@@ -142,6 +148,7 @@ class graph_mgr {
142148
dpct::experimental::node_ptr *dependencies,
143149
std::size_t numberOfDependencies,
144150
dpct::experimental::kernel_node_params *params) {
151+
node_graph_params_map[*node] = std::make_pair(graph, params);
145152
for (std::size_t i = 0; i < numberOfDependencies; i++) {
146153
params->add_dependency(dependencies[i]);
147154
}
@@ -156,7 +163,6 @@ class graph_mgr {
156163
for (std::size_t i = 0; i < kernel_params_vector.size(); i++) {
157164
auto &node_kernel_params_pair = kernel_params_vector[i];
158165
auto node_params = node_kernel_params_pair.second;
159-
160166
const auto &dependency_ptrs = node_params->get_dependencies();
161167
std::vector<sycl::ext::oneapi::experimental::node> dependencies;
162168
dependencies.reserve(dependency_ptrs.size());
@@ -184,9 +190,12 @@ class graph_mgr {
184190
}
185191
node_kernel_params_pair.first = new_node;
186192
}
187-
auto final_graph = graph->finalize();
193+
execGraph = new sycl::ext::oneapi::experimental::command_graph<
194+
sycl::ext::oneapi::experimental::graph_state::executable>(
195+
graph->finalize(
196+
sycl::ext::oneapi::experimental::property::graph::updatable{}));
188197
queue->submit(
189-
[&](sycl::handler &cgh) { cgh.ext_oneapi_graph(final_graph); });
198+
[&](sycl::handler &cgh) { cgh.ext_oneapi_graph(*execGraph); });
190199
}
191200

192201
void instantiate(dpct::experimental::command_graph_exec_ptr *execGraph,
@@ -195,7 +204,31 @@ class graph_mgr {
195204
}
196205

197206
void kernel_node_get_params(dpct::experimental::node_ptr node,
198-
dpct::experimental::kernel_node_params *params) {}
207+
dpct::experimental::kernel_node_params *params) {
208+
auto it = node_graph_params_map.find(node);
209+
if (it == node_graph_params_map.end()) {
210+
return;
211+
}
212+
*params = *(it->second.second);
213+
}
214+
215+
void kernel_node_set_params(dpct::experimental::node_ptr node,
216+
dpct::experimental::kernel_node_params *params) {
217+
node_graph_params_map[node].second = params;
218+
}
219+
220+
void get_node_type(dpct::experimental::node_ptr node,
221+
sycl::ext::oneapi::experimental::node_type *nodeType) {
222+
if (node_graph_params_map.find(node) != node_graph_params_map.end()) {
223+
*nodeType = sycl::ext::oneapi::experimental::node_type::kernel;
224+
} else {
225+
if (node) {
226+
*nodeType = node->get_type();
227+
} else {
228+
*nodeType = sycl::ext::oneapi::experimental::node_type::empty;
229+
}
230+
}
231+
}
199232

200233
private:
201234
std::unordered_map<sycl::queue *, command_graph_ptr> queue_graph_map;
@@ -214,8 +247,9 @@ class graph_mgr {
214247
dpct::experimental::kernel_node_params *>>>
215248
graph_kernel_node_params_map;
216249
std::unordered_map<dpct::experimental::node_ptr,
217-
dpct::experimental::kernel_node_params>
218-
node_params_map;
250+
std::pair<dpct::experimental::command_graph_ptr,
251+
dpct::experimental::kernel_node_params *>>
252+
node_graph_params_map;
219253
};
220254
} // namespace detail
221255

@@ -326,11 +360,31 @@ static void launch(dpct::experimental::command_graph_exec_ptr execGraph,
326360

327361
static void
328362
kernel_node_get_params(dpct::experimental::node_ptr node,
329-
dpct::experimental::kernel_node_params *params) {}
363+
dpct::experimental::kernel_node_params *params) {
364+
detail::graph_mgr::instance().kernel_node_get_params(node, params);
365+
}
330366

331367
static void
332368
kernel_node_set_params(dpct::experimental::node_ptr node,
333-
dpct::experimental::kernel_node_params *params) {}
369+
dpct::experimental::kernel_node_params *params) {
370+
detail::graph_mgr::instance().kernel_node_set_params(node, params);
371+
}
372+
373+
static void
374+
get_node_type(dpct::experimental::node_ptr node,
375+
sycl::ext::oneapi::experimental::node_type *nodeType) {
376+
detail::graph_mgr::instance().get_node_type(node, nodeType);
377+
}
378+
379+
static void update(dpct::experimental::command_graph_exec_ptr graphExec,
380+
dpct::experimental::command_graph_ptr graph,
381+
int *updateResultInfo) {
382+
graphExec->update(*graph);
383+
if (!graphExec) {
384+
*updateResultInfo = 0;
385+
}
386+
*updateResultInfo = 1;
387+
}
334388

335389
} // namespace experimental
336390
} // namespace dpct

0 commit comments

Comments
 (0)