[SYCL] sycl_khr_free_function_commands extension

sycl/include/sycl/sycl.hpp

@@ -74,6 +74,7 @@
 #include <sycl/ext/intel/experimental/task_sequence_properties.hpp>
 #include <sycl/ext/intel/experimental/usm_properties.hpp>
 #include <sycl/ext/intel/usm_pointers.hpp>
+#include <sycl/ext/khr/free_function_commands.hpp>
 #include <sycl/ext/oneapi/bfloat16.hpp>
 #include <sycl/ext/oneapi/bindless_images.hpp>
 #include <sycl/ext/oneapi/device_global/device_global.hpp>

sycl/source/feature_test.hpp.in

@@ -111,9 +111,10 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_MEMORY 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1
-#define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1
+#define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1 
 #define SYCL_EXT_ONEAPI_NUM_COMPUTE_UNITS 1
 #define SYCL_EXT_ONEAPI_DEVICE_IMAGE_BACKEND_CONTENT 1
+#define SYCL_KHR_FREE_FUNCTION_COMMANDS 1
 // In progress yet
 #define SYCL_EXT_ONEAPI_ATOMIC16 0
 

sycl/test-e2e/FreeFunctionCommands/helpers.hpp

@@ -0,0 +1,19 @@
+#include <iostream>
+#include <string>
+
+static int Check(int *Data, int Expected, size_t Index,
+                 const std::string &TestName) {
+  if (Data[Index] == Expected)
+    return 0;
+  std::cout << "Failed " << TestName << " at index " << Index << " : "
+            << Data[Index] << " != " << Expected << std::endl;
+  return 1;
+}
+
+static int Check(int &Data, int Expected, const std::string &TestName) {
+  if (Data == Expected)
+    return 0;
+  std::cout << "Failed " << TestName << " : " << Data << " != " << Expected
+            << std::endl;
+  return 1;
+}

sycl/test-e2e/FreeFunctionCommands/launch.cpp

@@ -0,0 +1,78 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// This test checks whether the free function command launch is valid.
+#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/khr/free_function_commands.hpp>
+#include <sycl/usm.hpp>
+
+#include <array>
+
+#include "helpers.hpp"
+
+int main() {
+
+  int Failed = 0;
+
+  sycl::queue Queue;
+  constexpr size_t Dim = 8;
+
+  {
+    constexpr size_t N = Dim * Dim * Dim;
+
+    std::array<int, N> Numbers = {0};
+    {
+      sycl::buffer<int, 1> MemBuffer{Numbers.data(), sycl::range<1>{N}};
+      sycl::khr::submit(Queue, [&](sycl::handler &Handler) {
+        sycl::accessor MemAcc{MemBuffer, Handler, sycl::write_only};
+        sycl::khr::launch(Handler, sycl::range<3>(Dim, Dim, Dim),
+                          [=](sycl::item<3> Item) {
+                            size_t Index = Item.get_linear_id();
+                            MemAcc[Index] = 901;
+                          });
+      });
+    }
+    for (size_t i = 0; i < N; ++i)
+      Failed += Check(Numbers.data(), 901, i, "launch_with_buffer");
+  }
+
+  {
+    constexpr size_t N = Dim * Dim;
+    int *Numbers = sycl::malloc_shared<int>(N, Queue);
+
+    sycl::event Event =
+        sycl::khr::submit_tracked(Queue, [&](sycl::handler &Handler) {
+          sycl::khr::launch(Handler, sycl::range<1>(N),
+                            [=](sycl::item<1> Item) { Numbers[Item] = 2; });
+        });
+
+    sycl::khr::submit(Queue, [&](sycl::handler &Handler) {
+      Handler.depends_on(Event);
+      sycl::khr::launch(Handler, sycl::range<3>(Dim / 4, Dim / 2, Dim),
+                        [=](sycl::item<3> Item) {
+                          size_t Index = Item.get_linear_id();
+                          Numbers[Index] = Numbers[Index] + 900;
+                        });
+    });
+
+    Queue.wait();
+    for (size_t i = 0; i < N; ++i)
+      Failed += Check(Numbers, 902, i, "launch_with_usm");
+
+    sycl::khr::launch(Queue, sycl::range<2>(Dim, Dim), [=](sycl::item<2> Item) {
+      size_t Index = Item.get_linear_id();
+      Numbers[Index] = 903;
+    });
+
+    Queue.wait();
+    for (size_t i = 0; i < N; ++i)
+      Failed += Check(Numbers, 903, i, "launch_shortcut_with_usm");
+
+    sycl::free(Numbers, Queue);
+  }
+
+  return Failed;
+}

sycl/test-e2e/FreeFunctionCommands/launch_grouped.cpp

@@ -0,0 +1,75 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// This test checks whether the free function command launch grouped is valid.
+
+#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/khr/free_function_commands.hpp>
+#include <sycl/usm.hpp>
+
+#include "helpers.hpp"
+
+int main() {
+
+  int Failed = 0;
+  sycl::queue Queue;
+  constexpr size_t N = 1024;
+
+  {
+    std::array<int, N> Numbers = {0};
+    {
+      sycl::buffer<int> MemBuffer{Numbers.data(), N};
+      Queue.submit([&](sycl::handler &Handler) {
+        sycl::accessor MemAcc{MemBuffer, Handler, sycl::write_only};
+        sycl::khr::launch_grouped(Handler, sycl::range<3>(8, 8, N / 64),
+                                  sycl::range<3>(4, 4, 4),
+                                  [=](sycl::nd_item<3> Item) {
+                                    size_t Index = Item.get_global_linear_id();
+                                    MemAcc[Index] = 301;
+                                  });
+      });
+    }
+    for (size_t i = 0; i < N; ++i)
+      Failed +=
+          Check(Numbers.data(), 301, i, "launch_grouped_with_sycl_buffer");
+  }
+
+  {
+    int *Numbers = sycl::malloc_shared<int>(N, Queue);
+
+    sycl::event Event =
+        sycl::khr::submit_tracked(Queue, [&](sycl::handler &Handler) {
+          sycl::khr::launch_grouped(Handler, sycl::range<3>(4, 4, N / 16),
+                                    sycl::range<3>(4, 4, 4),
+                                    [=](sycl::nd_item<3> Item) {
+                                      Numbers[Item.get_global_linear_id()] = 2;
+                                    });
+        });
+
+    sycl::khr::submit(Queue, [&](sycl::handler &Handler) {
+      Handler.depends_on(Event);
+      sycl::khr::launch_grouped(Handler, sycl::range<1>(N), sycl::range<1>(8),
+                                [=](sycl::nd_item<1> Item) {
+                                  size_t Index = Item.get_global_linear_id();
+                                  Numbers[Index] = Numbers[Index] + 300;
+                                });
+    });
+    Queue.wait();
+    for (size_t i = 0; i < N; ++i)
+      Failed += Check(Numbers, 302, i, "launch_grouped_with_usm");
+
+    sycl::khr::launch_grouped(Queue, sycl::range<2>(8, N / 8),
+                              sycl::range<2>(8, 8), [=](sycl::nd_item<2> Item) {
+                                Numbers[Item.get_global_linear_id()] = 303;
+                              });
+    Queue.wait();
+    for (size_t i = 0; i < N; ++i)
+      Failed += Check(Numbers, 303, i, "launch_grouped_shortcut_with_usm");
+    sycl::free(Numbers, Queue);
+  }
+
+  return Failed;
+}

sycl/test-e2e/FreeFunctionCommands/launch_grouped_reduce.cpp

@@ -0,0 +1,81 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// This test checks whether the free function command launch grouped reduce is
+// valid.
+
+#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/khr/free_function_commands.hpp>
+#include <sycl/reduction.hpp>
+#include <sycl/usm.hpp>
+
+#include "helpers.hpp"
+
+int main() {
+
+  int Failed = 0;
+  sycl::queue Queue;
+
+  {
+    int Result = 0;
+    constexpr size_t Dim = 8;
+    constexpr size_t N = Dim * Dim * Dim - 1;
+    constexpr int ExpectedResult = N * (N + 1) / 2;
+    {
+      sycl::buffer<int> sumBuf{&Result, 1};
+      Queue.submit([&](sycl::handler &Handler) {
+        sycl::khr::launch_grouped_reduce(
+            Handler, sycl::range<3>(Dim, Dim, Dim), sycl::range<3>(8, 8, 8),
+            [=](sycl::nd_item<3> Item, auto &Sum) {
+              Sum += Item.get_local_linear_id();
+            },
+            sycl::reduction(sumBuf, Handler, sycl::plus<>()));
+      });
+    }
+    Failed +=
+        Check(Result, ExpectedResult, "launch_grouped_reduce_with_buffer");
+  }
+
+  {
+    int *Result = sycl::malloc_shared<int>(1, Queue);
+    Result[0] = 0;
+    constexpr size_t N = 1024;
+    constexpr int ExpectedResult = ((N - 1) * N) / 2;
+
+    sycl::event Event =
+        sycl::khr::submit_tracked(Queue, [&](sycl::handler &Handler) {
+          sycl::khr::launch_grouped_reduce(
+              Handler, sycl::range<3>(8, 8, N / 64), sycl::range<3>(8, 8, 8),
+              [=](sycl::nd_item<3> Item, auto &Sum) { Sum += 1; },
+              sycl::reduction(Result, sycl::plus<>()));
+        });
+
+    sycl::khr::submit(Queue, [&](sycl::handler &Handler) {
+      Handler.depends_on(Event);
+      sycl::khr::launch_grouped_reduce(
+          Handler, sycl::range<2>(16, N / 16), sycl::range<2>(8, 8),
+          [=](sycl::nd_item<2> Item, auto &Sum) { Sum += 1; },
+          sycl::reduction(Result, sycl::plus<>()));
+    });
+    Queue.wait();
+    Failed += Check(Result[0], N * 2, "launch_grouped_reduce_with_usm");
+
+    Result[0] = 0;
+    sycl::khr::launch_grouped_reduce(
+        Queue, sycl::range<1>(N), sycl::range<1>(8),
+        [=](sycl::nd_item<1> Item, auto &Sum) {
+          Sum += Item.get_global_linear_id();
+        },
+        sycl::reduction(Result, sycl::plus<>()));
+
+    Queue.wait();
+    Failed += Check(Result[0], ExpectedResult,
+                    "launch_grouped_reduce_shortcut_with_usm");
+    sycl::free(Result, Queue);
+  }
+
+  return Failed;
+}

sycl/test-e2e/FreeFunctionCommands/launch_reduce.cpp

@@ -0,0 +1,76 @@
+// REQUIRES: aspect-usm_shared_allocations
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// This test checks whether the free function command launch reduce is valid.
+#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/khr/free_function_commands.hpp>
+#include <sycl/reduction.hpp>
+#include <sycl/usm.hpp>
+
+#include "helpers.hpp"
+
+int main() {
+
+  int Failed = 0;
+  sycl::queue Queue;
+
+  {
+    int Result = 0;
+    constexpr size_t Dim = 8;
+    constexpr size_t N = Dim * Dim * Dim - 1;
+    constexpr int ExpectedResult = N * (N + 1) / 2;
+    {
+      sycl::buffer<int> sumBuf{&Result, 1};
+      Queue.submit([&](sycl::handler &Handler) {
+        sycl::khr::launch_reduce(
+            Handler, sycl::range<3>(Dim, Dim, Dim),
+            [=](sycl::item<3> Item, auto &Sum) { Sum += Item.get_linear_id(); },
+            sycl::reduction(sumBuf, Handler, sycl::plus<>()));
+      });
+    }
+    Failed += Check(Result, ExpectedResult, "launch_reduce_with_buffer");
+  }
+
+  {
+    int *Result = sycl::malloc_shared<int>(1, Queue);
+    Result[0] = 0;
+    constexpr size_t N = 1024;
+    constexpr int ExpectedResult = ((N - 1) * N) / 2;
+
+    sycl::event Event =
+        sycl::khr::submit_tracked(Queue, [&](sycl::handler &Handler) {
+          sycl::khr::launch_reduce(
+              Handler, sycl::range<1>(N),
+              [=](sycl::id<1> Id, auto &Sum) {
+                int NegativeId = -(int)Id;
+                Sum += NegativeId;
+              },
+              sycl::reduction(Result, sycl::plus<>()));
+        });
+
+    Queue.submit([&](sycl::handler &Handler) {
+      Handler.depends_on(Event);
+      sycl::khr::launch_reduce(
+          Handler, sycl::range<1>(N),
+          [=](sycl::item<1> Item, auto &Sum) { Sum += Item.get_linear_id(); },
+          sycl::reduction(Result, sycl::plus<>()));
+    });
+
+    Queue.wait();
+    Failed += Check(Result[0], 0, "launch_reduce_with_usm");
+
+    sycl::khr::launch_reduce(
+        Queue, sycl::range<1>(N), [=](sycl::id<1> Id, auto &Sum) { Sum += Id; },
+        sycl::reduction(Result, sycl::plus<>()));
+
+    Queue.wait();
+    Failed +=
+        Check(Result[0], ExpectedResult, "launch_reduce_shortcut_with_usm");
+    sycl::free(Result, Queue);
+  }
+
+  return Failed;
+}