Skip to content

Commit 6ab7370

Browse files
TejaX-AlaghariTejaX-Alaghari
committed
Added launch_config objects to pass kernel properties
1 parent c725bae commit 6ab7370

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
@@ -5903,7 +5903,6 @@ void KernelCallExpr::printSubmit(KernelPrinter &Printer) {
59035903
Printer << "*" << getEvent() << " = ";
59045904
}
59055905

5906-
printStreamBase(Printer);
59075906
if (isDefaultStream()) {
59085907
SubmitStmts.DefaultStreamFlag = true;
59095908
}
@@ -5912,8 +5911,12 @@ void KernelCallExpr::printSubmit(KernelPrinter &Printer) {
59125911
SubmitStmts.ImplicitSyncFlag = true;
59135912
}
59145913
if (SubmitStmts.empty()) {
5914+
if (ExecutionConfig.Properties.empty()) {
5915+
printStreamBase(Printer);
5916+
}
59155917
printParallelFor(Printer, false);
59165918
} else {
5919+
printStreamBase(Printer);
59175920
(Printer << "submit(").newLine();
59185921
printSubmitLambda(Printer);
59195922
}
@@ -5945,12 +5948,20 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59455948
}
59465949
}
59475950
}
5951+
bool UseEnqueueFunctions = !ExecutionConfig.Properties.empty();
59485952
if (IsInSubmit) {
5949-
Printer.indent() << "cgh.";
5953+
Printer.indent();
5954+
if (!UseEnqueueFunctions) {
5955+
Printer << "cgh.";
5956+
}
59505957
}
59515958
if (!SubmitStmts.NdRangeList.empty() && DpctGlobalInfo::isCommentsEnabled())
59525959
Printer.line("// run the kernel within defined ND range");
5953-
Printer << "parallel_for";
5960+
if (UseEnqueueFunctions) {
5961+
Printer << MapNames::getExpNamespace() << "nd_launch";
5962+
} else {
5963+
Printer << "parallel_for";
5964+
}
59545965
if (DpctGlobalInfo::isSyclNamedLambda()) {
59555966
Printer << "<dpct_kernel_name<class " << getName() << "_"
59565967
<< LocInfo.LocHash;
@@ -5961,16 +5972,26 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59615972
}
59625973
(Printer << "(").newLine();
59635974
auto B = Printer.block();
5975+
std::unique_ptr<KernelPrinter::Block> LaunchConfigBlock;
5976+
if (UseEnqueueFunctions) {
5977+
(Printer.indent() << (IsInSubmit ? "cgh" : ExecutionConfig.Stream) << ",")
5978+
.newLine();
5979+
}
59645980
static std::string CanIgnoreRangeStr3D =
59655981
DpctGlobalInfo::getCtadClass(MapNames::getClNamespace() + "range", 3) +
59665982
"(1, 1, 1)";
59675983
static std::string CanIgnoreRangeStr1D =
59685984
DpctGlobalInfo::getCtadClass(MapNames::getClNamespace() + "range", 1) +
59695985
"(1)";
59705986
if (ExecutionConfig.NdRange != "") {
5987+
if (UseEnqueueFunctions) {
5988+
Printer.line(MapNames::getExpNamespace() + "launch_config(");
5989+
LaunchConfigBlock = std::move(Printer.block());
5990+
}
59715991
Printer.line(ExecutionConfig.NdRange + ",");
5972-
if (!ExecutionConfig.Properties.empty()) {
5973-
Printer << ExecutionConfig.Properties << ", ";
5992+
if (UseEnqueueFunctions) {
5993+
Printer.line(ExecutionConfig.Properties + "),");
5994+
LaunchConfigBlock.reset();
59745995
}
59755996
Printer.line("[=](", MapNames::getClNamespace(), "nd_item<3> ",
59765997
getItemName(), ")", ExecutionConfig.SubGroupSize, " {");
@@ -5980,6 +6001,10 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59806001
MemVarMap::getHeadWithoutPathCompression(
59816002
&(getFuncInfo()->getVarMap()))
59826003
->Dim == 1) {
6004+
if (UseEnqueueFunctions) {
6005+
Printer.line(MapNames::getExpNamespace() + "launch_config(");
6006+
LaunchConfigBlock = std::move(Printer.block());
6007+
}
59836008
DpctGlobalInfo::printCtadClass(Printer.indent(),
59846009
MapNames::getClNamespace() + "nd_range", 1)
59856010
<< "(";
@@ -5994,12 +6019,17 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
59946019
Printer << ", ";
59956020
Printer << ExecutionConfig.LocalSizeFor1D;
59966021
(Printer << "), ").newLine();
5997-
if (!ExecutionConfig.Properties.empty()) {
5998-
Printer << ExecutionConfig.Properties << ", ";
6022+
if (UseEnqueueFunctions) {
6023+
Printer.line(ExecutionConfig.Properties + "),");
6024+
LaunchConfigBlock.reset();
59996025
}
60006026
Printer.line("[=](" + MapNames::getClNamespace() + "nd_item<1> ",
60016027
getItemName(), ")", ExecutionConfig.SubGroupSize, " {");
60026028
} else {
6029+
if (UseEnqueueFunctions) {
6030+
Printer.line(MapNames::getExpNamespace() + "launch_config(");
6031+
LaunchConfigBlock = std::move(Printer.block());
6032+
}
60036033
Printer.indent();
60046034
Printer << MapNames::getClNamespace() + "nd_range<3>(";
60056035
if (ExecutionConfig.GroupSize == CanIgnoreRangeStr3D) {
@@ -6013,8 +6043,9 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) {
60136043
Printer << ", ";
60146044
Printer << ExecutionConfig.LocalSize;
60156045
(Printer << "), ").newLine();
6016-
if (!ExecutionConfig.Properties.empty()) {
6017-
Printer << ExecutionConfig.Properties << ", ";
6046+
if (UseEnqueueFunctions) {
6047+
Printer.line(ExecutionConfig.Properties + "),");
6048+
LaunchConfigBlock.reset();
60186049
}
60196050
Printer.line("[=](" + MapNames::getClNamespace() + "nd_item<3> ",
60206051
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)