[SYCL][Graph] Add testing for sycl_ext_oneapi_local_memory (#16379)

- Adds testing that verifies the interaction between
sycl_ext_oneapi_local_memory and sycl_ext_oneapi_graph.
- Reorder the extensions in the sycl graph spec to be listed in
alphabetical order.
- Explicitly state in the sycl graph spec that using
sycl_ext_oneapi_local_memory is supported.

Author: Fábio Mestre

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -2075,6 +2075,45 @@ recording mode, as opposed to throwing.
 This section defines the interaction of `sycl_ext_oneapi_graph` with other
 extensions.
 
+==== sycl_ext_codeplay_enqueue_native_command
+
+`ext_codeplay_enqueue_native_command`, defined in
+link:../experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc[sycl_ext_codeplay_enqueue_native_command]
+cannot be used in graph nodes. A synchronous exception will be thrown with error
+code `invalid` if a user tries to add them to a graph.
+
+Removing this restriction is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
+==== sycl_ext_intel_queue_index
+
+The compute index queue property defined by
+link:../supported/sycl_ext_intel_queue_index.asciidoc[sycl_ext_intel_queue_index]
+is ignored during queue recording.
+
+Using this information is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
+==== sycl_ext_oneapi_bindless_images
+
+The new handler methods, and queue shortcuts, defined by
+link:../experimental/sycl_ext_oneapi_bindless_images.asciidoc[sycl_ext_oneapi_bindless_images]
+cannot be used in graph nodes. A synchronous exception will be thrown with error
+code `invalid` if a user tries to add them to a graph.
+
+Removing this restriction is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
+==== sycl_ext_oneapi_device_global
+
+The new handler methods, and queue shortcuts, defined by
+link:../experimental/sycl_ext_oneapi_device_global.asciidoc[sycl_ext_oneapi_device_global].
+cannot be used in graph nodes. A synchronous exception will be thrown with error
+code `invalid` if a user tries to add them to a graph.
+
+Removing this restriction is something we may look at for future revisions of
+`sycl_ext_oneapi_graph`.
+
 ==== sycl_ext_oneapi_discard_queue_events
 
 When recording a `sycl::queue` which has been created with the
@@ -2108,37 +2147,25 @@ nodes that are recorded from multiple queues and/or added by the explicit API:
 * The only commands which have an implicit dependency on the barrier command
   are those recorded from the same queue the barrier command was submitted to.
 
-==== sycl_ext_oneapi_memcpy2d
-
-The new handler methods, and queue shortcuts, defined by
-link:../supported/sycl_ext_oneapi_memcpy2d.asciidoc[sycl_ext_oneapi_memcpy2d]
-cannot be used in graph nodes. A synchronous exception will be thrown with
-error code `invalid` if a user tries to add them to a graph.
-
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
-
-==== sycl_ext_oneapi_queue_priority
-
-The queue priority property defined by
-link:../supported/sycl_ext_oneapi_queue_priority.asciidoc[sycl_ext_oneapi_queue_priority]
-is ignored during queue recording.
+==== sycl_ext_oneapi_enqueue_functions
 
-==== sycl_ext_oneapi_queue_empty
+The command submission functions defined in
+link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions]
+can be used adding nodes to a graph when creating a graph from queue recording.
+New methods are also defined that enable submitting an executable graph,
+e.g. directly to a queue without returning an event.
 
-The `queue::ext_oneapi_empty()` query defined by the
-link:../supported/sycl_ext_oneapi_queue_empty.asciidoc[sycl_ext_oneapi_queue_empty]
-extension behaves as normal during queue recording and is not captured to the graph.
-Recorded commands are not counted as submitted for the purposes of this query.
+==== sycl_ext_oneapi_free_function_kernels
 
-==== sycl_ext_intel_queue_index
+`sycl_ext_oneapi_free_function_kernels`, defined in
+link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[sycl_ext_oneapi_free_function_kernels]
+can be used with SYCL Graphs.
 
-The compute index queue property defined by
-link:../supported/sycl_ext_intel_queue_index.asciidoc[sycl_ext_intel_queue_index]
-is ignored during queue recording.
+==== sycl_ext_oneapi_kernel_compiler_spirv
 
-Using this information is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
+The kernels loaded using
+link:../experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc[sycl_ext_oneapi_kernel_compiler_spirv]
+behave as normal when used in graph nodes.
 
 ==== sycl_ext_oneapi_kernel_properties
 
@@ -2147,62 +2174,41 @@ link:../experimental/sycl_ext_oneapi_kernel_properties.asciidoc[sycl_ext_oneapi_
 can be used in graph nodes in the same way as they are used in normal queue
 submission. 
 
-==== sycl_ext_oneapi_prod
+==== sycl_ext_oneapi_local_memory
 
-The new `sycl::queue::ext_oneapi_prod()` method added by
-link:../proposed/sycl_ext_oneapi_prod.asciidoc[sycl_ext_oneapi_prod]
-behaves as normal during queue recording and is not captured to the graph.
-Recorded commands are not counted as submitted for the purposes of its operation.
+Allocating local memory inside a graph kernel node with `sycl::ext::oneapi::group_local_memory()` or
+`sycl::ext::oneapi::group_local_memory_for_overwrite()` is supported. These methods are defined by
+link:../supported/sycl_ext_oneapi_local_memory.asciidoc[sycl_ext_oneapi_local_memory.]
 
-==== sycl_ext_oneapi_device_global
-
-The new handler methods, and queue shortcuts, defined by
-link:../experimental/sycl_ext_oneapi_device_global.asciidoc[sycl_ext_oneapi_device_global].
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
-
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
-
-==== sycl_ext_oneapi_bindless_images
+==== sycl_ext_oneapi_memcpy2d
 
 The new handler methods, and queue shortcuts, defined by
-link:../experimental/sycl_ext_oneapi_bindless_images.asciidoc[sycl_ext_oneapi_bindless_images]
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
+link:../supported/sycl_ext_oneapi_memcpy2d.asciidoc[sycl_ext_oneapi_memcpy2d]
+cannot be used in graph nodes. A synchronous exception will be thrown with
+error code `invalid` if a user tries to add them to a graph.
 
 Removing this restriction is something we may look at for future revisions of
 `sycl_ext_oneapi_graph`.
 
-==== sycl_ext_oneapi_kernel_compiler_spirv
-
-The kernels loaded using
-link:../experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc[sycl_ext_oneapi_kernel_compiler_spirv]
-behave as normal when used in graph nodes.
-
-==== sycl_ext_codeplay_enqueue_native_command
-
-`ext_codeplay_enqueue_native_command`, defined in
-link:../experimental/sycl_ext_codeplay_enqueue_native_command.asciidoc[sycl_ext_codeplay_enqueue_native_command]
-cannot be used in graph nodes. A synchronous exception will be thrown with error
-code `invalid` if a user tries to add them to a graph.
+==== sycl_ext_oneapi_prod
 
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
+The new `sycl::queue::ext_oneapi_prod()` method added by
+link:../proposed/sycl_ext_oneapi_prod.asciidoc[sycl_ext_oneapi_prod]
+behaves as normal during queue recording and is not captured to the graph.
+Recorded commands are not counted as submitted for the purposes of its operation.
 
-==== sycl_ext_oneapi_enqueue_functions
+==== sycl_ext_oneapi_queue_empty
 
-The command submission functions defined in
-link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions]
-can be used adding nodes to a graph when creating a graph from queue recording.
-New methods are also defined that enable submitting an executable graph,
-e.g. directly to a queue without returning an event.
+The `queue::ext_oneapi_empty()` query defined by the
+link:../supported/sycl_ext_oneapi_queue_empty.asciidoc[sycl_ext_oneapi_queue_empty]
+extension behaves as normal during queue recording and is not captured to the graph.
+Recorded commands are not counted as submitted for the purposes of this query.
 
-==== sycl_ext_oneapi_free_function_kernels
+==== sycl_ext_oneapi_queue_priority
 
-`sycl_ext_oneapi_free_function_kernels`, defined in
-link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[sycl_ext_oneapi_free_function_kernels]
-can be used with SYCL Graphs.
+The queue priority property defined by
+link:../supported/sycl_ext_oneapi_queue_priority.asciidoc[sycl_ext_oneapi_queue_priority]
+is ignored during queue recording.
 
 ==== sycl_ext_oneapi_work_group_memory
 

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

@@ -0,0 +1,10 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// Extra run to check for immediate-command-list in Level Zero
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+#define GRAPH_E2E_EXPLICIT
+
+#include "../Inputs/compile_time_local_memory.cpp"

sycl/test-e2e/Graph/Inputs/compile_time_local_memory.cpp

@@ -0,0 +1,73 @@
+// Tests adding kernel nodes with local memory that is allocated using
+// the sycl_ext_oneapi_local_memory extension.
+
+#include "../graph_common.hpp"
+#include <sycl/ext/oneapi/group_local_memory.hpp>
+
+int main() {
+  queue Queue{};
+
+  using T = int;
+  constexpr size_t LocalSize = 128;
+
+  std::vector<T> HostData(Size);
+  std::iota(HostData.begin(), HostData.end(), 10);
+
+  exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()};
+
+  T *PtrA = malloc_device<T>(Size, Queue);
+
+  Queue.copy(HostData.data(), PtrA, Size);
+  Queue.wait_and_throw();
+
+  auto NodeA = add_node(Graph, Queue, [&](handler &CGH) {
+    CGH.parallel_for(nd_range({Size}, {LocalSize}), [=](nd_item<1> Item) {
+      multi_ptr<size_t[LocalSize], access::address_space::local_space>
+          LocalMem = sycl::ext::oneapi::group_local_memory<size_t[LocalSize]>(
+              Item.get_group());
+      *LocalMem[Item.get_local_linear_id()] = Item.get_global_linear_id() * 2;
+      PtrA[Item.get_global_linear_id()] +=
+          *LocalMem[Item.get_local_linear_id()];
+    });
+  });
+
+  add_node(
+      Graph, Queue,
+      [&](handler &CGH) {
+        depends_on_helper(CGH, NodeA);
+        CGH.parallel_for(nd_range({Size}, {LocalSize}), [=](nd_item<1> Item) {
+          multi_ptr<size_t[LocalSize], access::address_space::local_space>
+              LocalMem = sycl::ext::oneapi::group_local_memory_for_overwrite<
+                  size_t[LocalSize]>(Item.get_group());
+          *LocalMem[Item.get_local_linear_id()] =
+              Item.get_global_linear_id() + 4;
+          PtrA[Item.get_global_linear_id()] *=
+              *LocalMem[Item.get_local_linear_id()];
+        });
+      },
+      NodeA);
+
+  auto GraphExec = Graph.finalize();
+
+  for (unsigned n = 0; n < Iterations; n++) {
+    Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); });
+  }
+
+  Queue.wait_and_throw();
+
+  Queue.copy(PtrA, HostData.data(), Size);
+  Queue.wait_and_throw();
+
+  free(PtrA, Queue);
+
+  for (size_t i = 0; i < Size; i++) {
+    T Ref = 10 + i;
+    for (size_t iter = 0; iter < Iterations; ++iter) {
+      Ref += (i * 2);
+      Ref *= (i + 4);
+    }
+    assert(check_value(i, Ref, HostData[i], "PtrA"));
+  }
+
+  return 0;
+}

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

@@ -0,0 +1,10 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// Extra run to check for immediate-command-list in Level Zero
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+#define GRAPH_E2E_RECORD_REPLAY
+
+#include "../Inputs/compile_time_local_memory.cpp"