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

sycl/doc/design/CommandGraph.md

@@ -210,7 +210,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
 
@@ -325,6 +325,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,13 +56,14 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
  include(FetchContent)
 
- set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
+ 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 3487672ceba0fd3575b5f3f15a832b100dc5fbad)
+ 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