Changed depends_on arg repl as well

oneapi-src · Feb 7, 2025 · 6a3891b · 6a3891b
1 parent d0948dc
commit 6a3891b
Show file tree

Hide file tree

Showing 2 changed files with 16 additions and 0 deletions.
diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp
@@ -6665,8 +6665,19 @@ KernelPrinter &KernelCallExpr::SubmitStmtsList::print(KernelPrinter &Printer) {
       Printer.line("cgh.depends_on(dpct::get_current_device().get_in_order_"
                    "queues_last_events());");
     } else {
+      Printer.line("#ifdef __INTEL_LLVM_COMPILER");
+      Printer.newLine();
       Printer.line("cgh.depends_on(dpct::get_default_queue().ext_oneapi_get_"
                    "last_event());");
+      Printer.newLine();
+      Printer.line("#else");
+      Printer.newLine();
+      Printer.line("auto e_opt = dpct::get_default_queue().ext_oneapi_get_last_"
+                    "event();");
+      Printer.newLine();
+      Printer.line("if (e_opt) cgh.depends_on(*e_opt);");
+      Printer.newLine();
+      Printer.line("#endif");
     }
     Printer.newLine();
   }

diff --git a/clang/test/dpct/kernel_implicit_sync.cu b/clang/test/dpct/kernel_implicit_sync.cu
@@ -27,7 +27,12 @@ int main() {
 
 // CHECK:  s1->submit(
 // CHECK:      [&](sycl::handler &cgh) {
+// CHECK:#ifdef __INTEL_LLVM_COMPILER
 // CHECK:        cgh.depends_on(dpct::get_default_queue().ext_oneapi_get_last_event());
+// CHECK:#else
+// CHECK:        auto e_opt = dpct::get_default_queue().ext_oneapi_get_last_event();
+// CHECK:        if (e_opt) cgh.depends_on(*e_opt);
+// CHECK:#endif
 // CHECK:        cgh.parallel_for(
 // CHECK:          sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
 // CHECK:          [=](sycl::nd_item<3> item_ct1) {