[SYCL][Graph][HIP] Add HIP backend support to SYCL-Graph

sycl/doc/design/CommandGraph.md

@@ -236,7 +236,7 @@ Implementation of UR command-buffers
 for each of the supported SYCL 2020 backends.
 
 Backends which are implemented currently are: [Level Zero](#level-zero),
-[CUDA](#cuda), and partial support for [OpenCL](#opencl).
+[CUDA](#cuda), [HIP](#hip) and partial support for [OpenCL](#opencl).
 
 ### Level Zero
 
@@ -351,6 +351,27 @@ An executable CUDA Graph, which contains all commands and synchronization
 information, is saved in the UR command-buffer to allow for efficient
 graph resubmission.
 
+### HIP
+
+The HIP backend offers a Graph managemenet API very similar to CUDA Graph
+feature for batching series of operations.
+The SYCL Graph HIP backend implementation is therefore very similar to that of CUDA.
+
+UR commands (e.g. kernels) are mapped as graph nodes using the
+[HIP Management API](https://docs.amd.com/projects/HIP/en/docs-5.5.0/doxygen/html/group___graph.html).
+Synchronization between commands (UR sync-points) is implemented
+using graph dependencies.
+Executable HIP Graphs can be submitted to a HIP stream
+in the same way as regular kernels.
+The HIP backend enables enqueuing events to wait for into a stream.
+It also allows signaling the completion of a submission with an event.
+Therefore, submitting a UR command-buffer consists only of submitting to a stream
+the executable HIP Graph that represent this series of operations.
+
+An executable HIP Graph, which contains all commands and synchronization
+information, is saved in the UR command-buffer to allow for efficient
+graph resubmission.
+
 ### OpenCL
 
 SYCL-Graph is only enabled for an OpenCL backend when the

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -56,14 +56,14 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 588615e90bfd2b889834120dfff172236c6b8aa8
-  # Merge: 4e69cc60 47084751
-  # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Thu Feb 22 16:10:13 2024 +0000
-  #     Merge pull request #1371 from pbalcer/l0-query-status-sync-deadlock
-  #     [L0] fix a deadlock in queue sync and event status query
-  set(UNIFIED_RUNTIME_TAG 588615e90bfd2b889834120dfff172236c6b8aa8)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/bensuo/unified-runtime.git")
+
+  # commit 3487672ceba0fd3575b5f3f15a832b100dc5fbad
+  # Author: Artur Gainullin <artur.gainullin@intel.com>
+  # Date:   Fri Feb 16 09:59:50 2024 -0800
+  #
+  #     [UR] Provide flexibility to replace unified-memory-framework repo and tag
+  set(UNIFIED_RUNTIME_TAG cmdbuf-support-hip)
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")

sycl/test-e2e/Graph/Explicit/enqueue_ordering.cpp

@@ -18,6 +18,10 @@ int main() {
     return 0;
   }
 
+  if (!Queue.get_device().has(sycl::aspect::usm_shared_allocations)) {
+    return 0;
+  }
+
   exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()};
 
   const size_t N = 10;

sycl/test-e2e/Graph/Explicit/work_group_size_prop.cpp

@@ -11,10 +11,9 @@
 // should. So the Sycl graph support cannot correctly catch the error and throw
 // the approriate exception for negative test. An issue has been reported
 // https://github.com/bashbaug/SimpleOpenCLSamples/issues/95
-// XFAIL: cuda
+// XFAIL: cuda, hip
 // UNSUPPORTED: opencl
 // Note: failing negative test with HIP in the original test
-// TODO: disable hip when HIP backend will be supported by Graph
 
 #define GRAPH_E2E_EXPLICIT
 

sycl/test-e2e/Graph/RecordReplay/work_group_size_prop.cpp

@@ -11,10 +11,9 @@
 // should. So the Sycl graph support cannot correctly catch the error and throw
 // the approriate exception for negative test. An issue has been reported
 // https://github.com/bashbaug/SimpleOpenCLSamples/issues/95
-// XFAIL: cuda
+// XFAIL: cuda, hip
 // UNSUPPORTED: opencl
 // Note: failing negative test with HIP in the original test
-// TODO: disable hip when HIP backend will be supported by Graph
 
 #define GRAPH_E2E_RECORD_REPLAY
 

sycl/test-e2e/Graph/device_query.cpp

@@ -20,7 +20,8 @@ int main() {
   auto Backend = Device.get_backend();
 
   if ((Backend == backend::ext_oneapi_level_zero) ||
-      (Backend == backend::ext_oneapi_cuda)) {
+      (Backend == backend::ext_oneapi_cuda) ||
+      (Backend == backend::ext_oneapi_hip)) {
     assert(SupportsGraphs == exp_ext::graph_support_level::native);
   } else if (Backend == backend::opencl) {
     // OpenCL backend support is conditional on the cl_khr_command_buffer