[SYCL][Graph] Introduce native recording mode feature

sycl/include/sycl/detail/property_helper.hpp

@@ -44,8 +44,9 @@ enum DataLessPropKind {
   GraphDependOnAllLeaves = 19,
   GraphUpdatable = 20,
   GraphEnableProfiling = 21,
+  GraphEnableNativeRecording = 22,
   // Indicates the last known dataless property.
-  LastKnownDataLessPropKind = 21,
+  LastKnownDataLessPropKind = 22,
   // Exceeding 32 may cause ABI breaking change on some of OSes.
   DataLessPropKindSize = 32
 };

sycl/include/sycl/ext/oneapi/experimental/detail/properties/graph_properties.def

@@ -22,5 +22,9 @@ __SYCL_DATA_LESS_PROP(property::graph, updatable, GraphUpdatable)
 /// events returned by submissions of the executable graph
 __SYCL_DATA_LESS_PROP(property::graph, enable_profiling, GraphEnableProfiling)
 
+/// Property passed to command_graph constructor to enable native recording
+/// capability for improved performance when supported by the backend.
+__SYCL_DATA_LESS_PROP(property::graph, enable_native_recording, GraphEnableNativeRecording)
+
 #undef __SYCL_DATA_LESS_PROP
 #undef __SYCL_MANUALLY_DEFINED_PROP

sycl/include/sycl/ext/oneapi/experimental/graph/modifiable_graph.hpp

@@ -159,7 +159,7 @@ class __SYCL_EXPORT modifiable_command_graph
   /// Get a list of all root nodes (nodes without dependencies) in this graph.
   std::vector<node> get_root_nodes() const;
 
-  /// Query whether the graph contains no nodes.
+  /// Returns true if the graph contains no nodes.
   bool empty() const;
 
   /// Common Reference Semantics

sycl/source/detail/async_alloc.cpp

@@ -67,6 +67,14 @@ void *async_malloc(sycl::handler &h, sycl::usm::alloc kind, size_t size) {
         sycl::make_error_code(sycl::errc::feature_not_supported),
         "Only device backed asynchronous allocations are supported!");
 
+  // Allocations not supported in graph native recording mode
+  if (auto *Queue = h.impl->get_queue_or_null();
+      Queue && Queue->isNativeRecording()) {
+    throw sycl::exception(
+        sycl::make_error_code(sycl::errc::invalid),
+        "async_malloc is not supported in native recording mode.");
+  }
+
   detail::adapter_impl &Adapter = h.getContextImpl().getAdapter();
 
   // Get CG event dependencies for this allocation.
@@ -117,6 +125,14 @@ __SYCL_EXPORT void *async_malloc(const sycl::queue &q, sycl::usm::alloc kind,
 __SYCL_EXPORT void *async_malloc_from_pool(sycl::handler &h, size_t size,
                                            const memory_pool &pool) {
 
+  // Allocations not supported in graph native recording mode
+  if (auto *Queue = h.impl->get_queue_or_null();
+      Queue && Queue->isNativeRecording()) {
+    throw sycl::exception(
+        sycl::make_error_code(sycl::errc::invalid),
+        "async_malloc is not supported in native recording mode.");
+  }
+
   detail::adapter_impl &Adapter = h.getContextImpl().getAdapter();
   detail::memory_pool_impl &memPoolImpl = *detail::getSyclObjImpl(pool);
 
@@ -181,6 +197,14 @@ __SYCL_EXPORT void async_free(sycl::handler &h, void *ptr) {
     }
   }
 
+  // Not supported in graph native recording mode
+  if (auto *Queue = h.impl->get_queue_or_null();
+      Queue && Queue->isNativeRecording()) {
+    throw sycl::exception(
+        sycl::make_error_code(sycl::errc::invalid),
+        "async_free is not supported in native recording mode.");
+  }
+
   h.impl->MFreePtr = ptr;
   h.setType(detail::CGType::AsyncFree);
 }