SYCL Graphs and empty barriers

[al42and]

Describe the bug

Trying to record an empty barrier while recording a graph throws an exception. The same pattern works fine when not recording a graph.

By itself, this operation is meaningless. However, it might appear in more complex code, like the following:

initializeData(q1, options); // May or may not submit anything, depending on `options`
sycl::event ev = q1.ext_oneapi_submit_barrier();
computeData(q2, ev); // Compute stuff (some operations depend on `ev`)

I don't see anything in either https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_enqueue_barrier.asciidoc or https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc prohibiting such use, so I believe it should be allowed.

To reproduce

#include <sycl/sycl.hpp>
#include <iostream>

namespace syclex = sycl::ext::oneapi::experimental;

int main() {
  for (const auto &dev : sycl::device::get_devices()) {
    if (dev.has(sycl::aspect::ext_oneapi_graph)) {
      std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
                << std::endl;
      sycl::context ctx{dev};
      sycl::queue q1{ctx, dev, {sycl::property::queue::in_order()}};
      const std::vector<sycl::queue> queuesToRecord{q1};

      const sycl::property_list propList{
          syclex::property::graph::no_cycle_check()};
      syclex::command_graph<syclex::graph_state::modifiable> graph(ctx, dev,
                                                                   propList);

      for (bool useGraph : {false, true}) {
        std::cout << "Running " << (useGraph ? "with" : "without")
                  << " graphs..." << std::endl;
        if (useGraph) {
          bool result = graph.begin_recording(queuesToRecord);
          if (!result) {
            std::cout << "  Could not start the recording" << std::endl;
          }
        }
        sycl::event ev = q1.ext_oneapi_submit_barrier();
        if (useGraph) {
          graph.end_recording();
          auto instance = graph.finalize();
          q1.ext_oneapi_graph(instance).wait_and_throw();
        }
        std::cout << "  Ok!" << std::endl;
      }
    }
  }
}

$ clang++ -fsycl -O1 -g test_graph_event.cpp && ONEAPI_DEVICE_SELECTOR=level_zero:0 ./a.out
Running on Intel(R) Arc(TM) A770 Graphics
Running without graphs...
  Ok!
Running with graphs...
terminate called after throwing an instance of 'sycl::_V1::exception'
  what():  No event has been recorded for the specified graph node
Aborted (core dumped)

Environment

• Ubuntu 22.04 HWE
• Arc A770 (via L0), RTX 3060, RX6400

$ clang++ --version
clang version 19.0.0git (https://github.com/intel/llvm 76167854fb3edba3a575302ccde14392c671529b)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/aland/intel-sycl/llvm/build/install/bin

$ sycl-ls 
[opencl:gpu][opencl:0] Intel(R) OpenCL Graphics, Intel(R) Arc(TM) A770 Graphics OpenCL 3.0 NEO  [23.35.27191.9]
[opencl:gpu][opencl:1] Intel(R) OpenCL Graphics, Intel(R) UHD Graphics 770 OpenCL 3.0 NEO  [23.35.27191.9]
[opencl:cpu][opencl:2] Intel(R) OpenCL, 12th Gen Intel(R) Core(TM) i9-12900K OpenCL 3.0 (Build 0) [2023.16.12.0.12_195853.xmain-hotfix]
[level_zero:gpu][level_zero:0] Intel(R) Level-Zero, Intel(R) Arc(TM) A770 Graphics 1.3 [1.3.27191]
[level_zero:gpu][level_zero:1] Intel(R) Level-Zero, Intel(R) UHD Graphics 770 1.3 [1.3.27191]
[cuda:gpu][cuda:0] NVIDIA CUDA BACKEND, NVIDIA GeForce RTX 3060 8.6 [CUDA 12.4]
[hip:gpu][hip:0] AMD HIP BACKEND, AMD Radeon RX 6400 gfx1034 [HIP 60032.83]

Additional context

No response

[al42and]

The problem appears even in less pathological cases:

sycl::queue q1{ctx, dev, {sycl::property::queue::in_order()}};
sycl::queue q2{ctx, dev, {sycl::property::queue::in_order()}};
const std::vector<sycl::queue> queuesToRecord{q1, q2};
// TODO: Allocate buf_h and buf_d...
// TODO: Start recording the graph  
sycl::event ev = q1.copy(buf_h, buf_d, n);
// Now, we want all future operations in q2 to wait for the copy to complete...
q2.ext_oneapi_submit_barrier({ev}); // Ooops, throws!

[EwanC]

Thanks for the bug report, proposed fix - #13193

[al42and]

Thanks for the fix, @EwanC!