From 6ab73700b7d3fc86ceff90f232df9f19ccaa742d Mon Sep 17 00:00:00 2001 From: Teja Alaghari Date: Wed, 11 Jun 2025 05:01:14 +0800 Subject: [PATCH] Added launch_config objects to pass kernel properties --- clang/lib/DPCT/AnalysisInfo.cpp | 49 +++++++++++++++++++++++------ clang/test/dpct/sync_api.cu | 23 ++++++++------ clang/test/dpct/sync_api_noneusm.cu | 9 ++++-- 3 files changed, 60 insertions(+), 21 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 9639ea44278a..f183ce2e4a2d 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -5903,7 +5903,6 @@ void KernelCallExpr::printSubmit(KernelPrinter &Printer) { Printer << "*" << getEvent() << " = "; } - printStreamBase(Printer); if (isDefaultStream()) { SubmitStmts.DefaultStreamFlag = true; } @@ -5912,8 +5911,12 @@ void KernelCallExpr::printSubmit(KernelPrinter &Printer) { SubmitStmts.ImplicitSyncFlag = true; } if (SubmitStmts.empty()) { + if (ExecutionConfig.Properties.empty()) { + printStreamBase(Printer); + } printParallelFor(Printer, false); } else { + printStreamBase(Printer); (Printer << "submit(").newLine(); printSubmitLambda(Printer); } @@ -5945,12 +5948,20 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) { } } } + bool UseEnqueueFunctions = !ExecutionConfig.Properties.empty(); if (IsInSubmit) { - Printer.indent() << "cgh."; + Printer.indent(); + if (!UseEnqueueFunctions) { + Printer << "cgh."; + } } if (!SubmitStmts.NdRangeList.empty() && DpctGlobalInfo::isCommentsEnabled()) Printer.line("// run the kernel within defined ND range"); - Printer << "parallel_for"; + if (UseEnqueueFunctions) { + Printer << MapNames::getExpNamespace() << "nd_launch"; + } else { + Printer << "parallel_for"; + } if (DpctGlobalInfo::isSyclNamedLambda()) { Printer << " LaunchConfigBlock; + if (UseEnqueueFunctions) { + (Printer.indent() << (IsInSubmit ? "cgh" : ExecutionConfig.Stream) << ",") + .newLine(); + } static std::string CanIgnoreRangeStr3D = DpctGlobalInfo::getCtadClass(MapNames::getClNamespace() + "range", 3) + "(1, 1, 1)"; @@ -5968,9 +5984,14 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) { DpctGlobalInfo::getCtadClass(MapNames::getClNamespace() + "range", 1) + "(1)"; if (ExecutionConfig.NdRange != "") { + if (UseEnqueueFunctions) { + Printer.line(MapNames::getExpNamespace() + "launch_config("); + LaunchConfigBlock = std::move(Printer.block()); + } Printer.line(ExecutionConfig.NdRange + ","); - if (!ExecutionConfig.Properties.empty()) { - Printer << ExecutionConfig.Properties << ", "; + if (UseEnqueueFunctions) { + Printer.line(ExecutionConfig.Properties + "),"); + LaunchConfigBlock.reset(); } Printer.line("[=](", MapNames::getClNamespace(), "nd_item<3> ", getItemName(), ")", ExecutionConfig.SubGroupSize, " {"); @@ -5980,6 +6001,10 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) { MemVarMap::getHeadWithoutPathCompression( &(getFuncInfo()->getVarMap())) ->Dim == 1) { + if (UseEnqueueFunctions) { + Printer.line(MapNames::getExpNamespace() + "launch_config("); + LaunchConfigBlock = std::move(Printer.block()); + } DpctGlobalInfo::printCtadClass(Printer.indent(), MapNames::getClNamespace() + "nd_range", 1) << "("; @@ -5994,12 +6019,17 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) { Printer << ", "; Printer << ExecutionConfig.LocalSizeFor1D; (Printer << "), ").newLine(); - if (!ExecutionConfig.Properties.empty()) { - Printer << ExecutionConfig.Properties << ", "; + if (UseEnqueueFunctions) { + Printer.line(ExecutionConfig.Properties + "),"); + LaunchConfigBlock.reset(); } Printer.line("[=](" + MapNames::getClNamespace() + "nd_item<1> ", getItemName(), ")", ExecutionConfig.SubGroupSize, " {"); } else { + if (UseEnqueueFunctions) { + Printer.line(MapNames::getExpNamespace() + "launch_config("); + LaunchConfigBlock = std::move(Printer.block()); + } Printer.indent(); Printer << MapNames::getClNamespace() + "nd_range<3>("; if (ExecutionConfig.GroupSize == CanIgnoreRangeStr3D) { @@ -6013,8 +6043,9 @@ void KernelCallExpr::printParallelFor(KernelPrinter &Printer, bool IsInSubmit) { Printer << ", "; Printer << ExecutionConfig.LocalSize; (Printer << "), ").newLine(); - if (!ExecutionConfig.Properties.empty()) { - Printer << ExecutionConfig.Properties << ", "; + if (UseEnqueueFunctions) { + Printer.line(ExecutionConfig.Properties + "),"); + LaunchConfigBlock.reset(); } Printer.line("[=](" + MapNames::getClNamespace() + "nd_item<3> ", getItemName(), ")", ExecutionConfig.SubGroupSize, " {"); diff --git a/clang/test/dpct/sync_api.cu b/clang/test/dpct/sync_api.cu index 91e6a2d07283..feeb20fabde3 100644 --- a/clang/test/dpct/sync_api.cu +++ b/clang/test/dpct/sync_api.cu @@ -109,9 +109,12 @@ int main() { // CHECK: { // CHECK-NEXT: auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync}; // CHECK-EMPTY: - // CHECK-NEXT: dpct::get_in_order_queue().parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)), - // CHECK-NEXT: exp_props, [=](sycl::nd_item<3> item_ct1) { + // CHECK-NEXT: sycl::ext::oneapi::experimental::nd_launch( + // CHECK-NEXT: dpct::get_in_order_queue(), + // CHECK-NEXT: sycl::ext::oneapi::experimental::launch_config( + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)), + // CHECK-NEXT: exp_props), + // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { // CHECK-NEXT: kernel(); // CHECK-NEXT: }); // CHECK-NEXT: } @@ -193,12 +196,14 @@ int foo3() { // CHECK-NEXT: auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync}; // CHECK-NEXT: dpct::has_capability_or_fail(dpct::get_in_order_queue().get_device(), {sycl::aspect::fp64}); // CHECK-EMPTY: - // CHECK-NEXT: dpct::get_in_order_queue().parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), - // CHECK-NEXT: exp_props, [=](sycl::nd_item<3> item_ct1) {{\[\[}}sycl::reqd_sub_group_size(32){{\]\]}} { - // CHECK-NEXT: foo2(); - // CHECK-NEXT: }); - // CHECK-NEXT: } + // CHECK-NEXT: sycl::ext::oneapi::experimental::nd_launch( + // CHECK-NEXT: dpct::get_in_order_queue(), + // CHECK-NEXT: sycl::ext::oneapi::experimental::launch_config( + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), + // CHECK-NEXT: exp_props), + // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {{\[\[}}sycl::reqd_sub_group_size(32){{\]\]}} { + // CHECK-NEXT: foo2(); + // CHECK-NEXT: }); foo2<<<1,1>>>(); return 0; } diff --git a/clang/test/dpct/sync_api_noneusm.cu b/clang/test/dpct/sync_api_noneusm.cu index 39000c186baf..69dc2f85e449 100644 --- a/clang/test/dpct/sync_api_noneusm.cu +++ b/clang/test/dpct/sync_api_noneusm.cu @@ -105,9 +105,12 @@ int main() { // CHECK: { // CHECK-NEXT: auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync}; // CHECK-EMPTY: - // CHECK-NEXT: dpct::get_out_of_order_queue().parallel_for( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)), - // CHECK-NEXT: exp_props, [=](sycl::nd_item<3> item_ct1) { + // CHECK-NEXT: sycl::ext::oneapi::experimental::nd_launch( + // CHECK-NEXT: dpct::get_out_of_order_queue(), + // CHECK-NEXT: sycl::ext::oneapi::experimental::launch_config( + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 2)), + // CHECK-NEXT: exp_props), + // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { // CHECK-NEXT: kernel(); // CHECK-NEXT: }); // CHECK-NEXT: }