Skip to content

Commit 51b410c

Browse files
TejaX-AlaghariTejaX-Alaghari
committed
Added launch_config objects to pass kernel properties
1 parent 2ab702b commit 51b410c

File tree

3 files changed

+60
-21
lines changed

3 files changed

+60
-21
lines changed

‎clang/lib/DPCT/AnalysisInfo.cpp‎

Lines changed: 40 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -5907,7 +5907,6 @@ void KernelCallExpr::printSubmit(KernelPrinter &Printer) {
59075907
Printer << "*" << getEvent() << " = ";
59085908
}
59095909

5910-
printStreamBase(Printer);
59115910
if (isDefaultStream()) {
59125911
SubmitStmts.DefaultStreamFlag = true;
59135912
}
@@ -5916,8 +5915,12 @@ void KernelCallExpr::printSubmit(KernelPrinter &Printer) {
59165915
SubmitStmts.ImplicitSyncFlag = true;
59175916
}
59185917
if (SubmitStmts.empty()) {
5918+
if (ExecutionConfig.Properties.empty()) {
5919+
printStreamBase(Printer);
5920+
}
59195921
printParallelFor(Printer, false);
59205922
} else {
5923+
printStreamBase(Printer);
59215924
(Printer << "submit(").newLine();
59225925
printSubmitLambda(Printer);
59235926
}
@@ -5949,12 +5952,20 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59495952
}
59505953
}
59515954
}
5955+
bool UseEnqueueFunctions = !ExecutionConfig.Properties.empty();
59525956
if (IsInSubmit) {
5953-
Printer.indent() << "cgh.";
5957+
Printer.indent();
5958+
if (!UseEnqueueFunctions) {
5959+
Printer << "cgh.";
5960+
}
59545961
}
59555962
if (!SubmitStmts.NdRangeList.empty() && DpctGlobalInfo::isCommentsEnabled())
59565963
Printer.line("// run the kernel within defined ND range");
5957-
Printer << "parallel_for";
5964+
if (UseEnqueueFunctions) {
5965+
Printer << MapNames::getExpNamespace() << "nd_launch";
5966+
} else {
5967+
Printer << "parallel_for";
5968+
}
59585969
if (DpctGlobalInfo::isSyclNamedLambda()) {
59595970
Printer << "<dpct_kernel_name<class " << getName() << "_"
59605971
<< LocInfo.LocHash;
@@ -5965,16 +5976,26 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59655976
}
59665977
(Printer << "(").newLine();
59675978
auto B = Printer.block();
5979+
std::unique_ptr<KernelPrinter::Block> LaunchConfigBlock;
5980+
if (UseEnqueueFunctions) {
5981+
(Printer.indent() << (IsInSubmit ? "cgh" : ExecutionConfig.Stream) << ",")
5982+
.newLine();
5983+
}
59685984
static std::string CanIgnoreRangeStr3D =
59695985
DpctGlobalInfo::getCtadClass(MapNames::getClNamespace() + "range", 3) +
59705986
"(1, 1, 1)";
59715987
static std::string CanIgnoreRangeStr1D =
59725988
DpctGlobalInfo::getCtadClass(MapNames::getClNamespace() + "range", 1) +
59735989
"(1)";
59745990
if (ExecutionConfig.NdRange != "") {
5991+
if (UseEnqueueFunctions) {
5992+
Printer.line(MapNames::getExpNamespace() + "launch_config(");
5993+
LaunchConfigBlock = std::move(Printer.block());
5994+
}
59755995
Printer.line(ExecutionConfig.NdRange + ",");
5976-
if (!ExecutionConfig.Properties.empty()) {
5977-
Printer << ExecutionConfig.Properties << ", ";
5996+
if (UseEnqueueFunctions) {
5997+
Printer.line(ExecutionConfig.Properties + "),");
5998+
LaunchConfigBlock.reset();
59785999
}
59796000
Printer.line("[=](", MapNames::getClNamespace(), "nd_item<3> ",
59806001
getItemName(), ")", ExecutionConfig.SubGroupSize, " {");
@@ -5984,6 +6005,10 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59846005
MemVarMap::getHeadWithoutPathCompression(
59856006
&(getFuncInfo()->getVarMap()))
59866007
->Dim == 1) {
6008+
if (UseEnqueueFunctions) {
6009+
Printer.line(MapNames::getExpNamespace() + "launch_config(");
6010+
LaunchConfigBlock = std::move(Printer.block());
6011+
}
59876012
DpctGlobalInfo::printCtadClass(Printer.indent(),
59886013
MapNames::getClNamespace() + "nd_range", 1)
59896014
<< "(";
@@ -5998,12 +6023,17 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59986023
Printer << ", ";
59996024
Printer << ExecutionConfig.LocalSizeFor1D;
60006025
(Printer << "), ").newLine();
6001-
if (!ExecutionConfig.Properties.empty()) {
6002-
Printer << ExecutionConfig.Properties << ", ";
6026+
if (UseEnqueueFunctions) {
6027+
Printer.line(ExecutionConfig.Properties + "),");
6028+
LaunchConfigBlock.reset();
60036029
}
60046030
Printer.line("[=](" + MapNames::getClNamespace() + "nd_item<1> ",
60056031
getItemName(), ")", ExecutionConfig.SubGroupSize, " {");
60066032
} else {
6033+
if (UseEnqueueFunctions) {
6034+
Printer.line(MapNames::getExpNamespace() + "launch_config(");
6035+
LaunchConfigBlock = std::move(Printer.block());
6036+
}
60076037
Printer.indent();
60086038
Printer << MapNames::getClNamespace() + "nd_range<3>(";
60096039
if (ExecutionConfig.GroupSize == CanIgnoreRangeStr3D) {
@@ -6017,8 +6047,9 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
60176047
Printer << ", ";
60186048
Printer << ExecutionConfig.LocalSize;
60196049
(Printer << "), ").newLine();
6020-
if (!ExecutionConfig.Properties.empty()) {
6021-
Printer << ExecutionConfig.Properties << ", ";
6050+
if (UseEnqueueFunctions) {
6051+
Printer.line(ExecutionConfig.Properties + "),");
6052+
LaunchConfigBlock.reset();
60226053
}
60236054
Printer.line("[=](" + MapNames::getClNamespace() + "nd_item<3> ",
60246055
getItemName(), ")", ExecutionConfig.SubGroupSize, " {");

‎clang/test/dpct/sync_api.cu‎

Lines changed: 14 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -109,9 +109,12 @@ int main() {
109109
// CHECK: {
110110
// CHECK-NEXT: auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
111111
// CHECK-EMPTY:
112-
// CHECK-NEXT: dpct::get_in_order_queue().parallel_for(
113-
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)),
114-
// CHECK-NEXT: exp_props, [=](sycl::nd_item<3> item_ct1) {
112+
// CHECK-NEXT: sycl::ext::oneapi::experimental::nd_launch(
113+
// CHECK-NEXT: dpct::get_in_order_queue(),
114+
// CHECK-NEXT: sycl::ext::oneapi::experimental::launch_config(
115+
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)),
116+
// CHECK-NEXT: exp_props),
117+
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
115118
// CHECK-NEXT: kernel();
116119
// CHECK-NEXT: });
117120
// CHECK-NEXT: }
@@ -193,12 +196,14 @@ int foo3() {
193196
// CHECK-NEXT: auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
194197
// CHECK-NEXT: dpct::has_capability_or_fail(dpct::get_in_order_queue().get_device(), {sycl::aspect::fp64});
195198
// CHECK-EMPTY:
196-
// CHECK-NEXT: dpct::get_in_order_queue().parallel_for(
197-
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
198-
// CHECK-NEXT: exp_props, [=](sycl::nd_item<3> item_ct1) {{\[\[}}sycl::reqd_sub_group_size(32){{\]\]}} {
199-
// CHECK-NEXT: foo2();
200-
// CHECK-NEXT: });
201-
// CHECK-NEXT: }
199+
// CHECK-NEXT: sycl::ext::oneapi::experimental::nd_launch(
200+
// CHECK-NEXT: dpct::get_in_order_queue(),
201+
// CHECK-NEXT: sycl::ext::oneapi::experimental::launch_config(
202+
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
203+
// CHECK-NEXT: exp_props),
204+
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {{\[\[}}sycl::reqd_sub_group_size(32){{\]\]}} {
205+
// CHECK-NEXT: foo2();
206+
// CHECK-NEXT: });
202207
foo2<<<1,1>>>();
203208
return 0;
204209
}

‎clang/test/dpct/sync_api_noneusm.cu‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -105,9 +105,12 @@ int main() {
105105
// CHECK: {
106106
// CHECK-NEXT: auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
107107
// CHECK-EMPTY:
108-
// CHECK-NEXT: dpct::get_out_of_order_queue().parallel_for(
109-
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)),
110-
// CHECK-NEXT: exp_props, [=](sycl::nd_item<3> item_ct1) {
108+
// CHECK-NEXT: sycl::ext::oneapi::experimental::nd_launch(
109+
// CHECK-NEXT: dpct::get_out_of_order_queue(),
110+
// CHECK-NEXT: sycl::ext::oneapi::experimental::launch_config(
111+
// CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)),
112+
// CHECK-NEXT: exp_props),
113+
// CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
111114
// CHECK-NEXT: kernel();
112115
// CHECK-NEXT: });
113116
// CHECK-NEXT: }

0 commit comments

Comments
 (0)