[SYCL] Add basic tests for virtual functions

sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp

@@ -0,0 +1,113 @@
+// UNSUPPORTED: cuda, hip, acc
+// FIXME: replace unsupported with an aspect check once we have it
+//
+// RUN: %{build} -o %t.out -Xclang -fsycl-allow-virtual-functions %helper-includes
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+
+#include "helpers.hpp"
+
+#include <iostream>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void increment(int *) { /* do nothhing */
+  }
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void multiply(int *) { /* do nothhing */
+  }
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void substract(int *) { /* do nothhing */
+  }
+};
+
+class IncrementBy1 : public Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 1; }
+};
+
+class IncrementBy1AndSubstractBy2 : public IncrementBy1 {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void substract(int *Data) override { *Data -= 2; }
+};
+
+class MultiplyBy2 : public Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void multiply(int *Data) override { *Data *= 2; }
+};
+
+class MultiplyBy2AndIncrementBy8 : public MultiplyBy2 {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 8; }
+};
+
+class SubstractBy4 : public Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void substract(int *Data) override { *Data -= 4; }
+};
+
+class SubstractBy4AndMultiplyBy4 : public SubstractBy4 {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void multiply(int *Data) override { *Data *= 4; }
+};
+
+void applyOp(int *DataPtr, Base *ObjPtr) {
+  ObjPtr->increment(DataPtr);
+  ObjPtr->substract(DataPtr);
+  ObjPtr->multiply(DataPtr);
+}
+
+int main() try {
+  using storage_t = obj_storage_t<IncrementBy1, IncrementBy1AndSubstractBy2,
+                                  MultiplyBy2, MultiplyBy2AndIncrementBy8,
+                                  SubstractBy4, SubstractBy4AndMultiplyBy4>;
+  storage_t HostStorage;
+  sycl::buffer<storage_t> DeviceStorage(sycl::range{1});
+
+  auto asyncHandler = [](sycl::exception_list list) {
+    for (auto &e : list)
+      std::rethrow_exception(e);
+  };
+
+  sycl::queue q(asyncHandler);
+
+  constexpr oneapi::properties props{oneapi::calls_indirectly<>};
+  for (unsigned TestCase = 0; TestCase < 6; ++TestCase) {
+    int HostData = 42;
+    int Data = HostData;
+    sycl::buffer<int> DataStorage(&Data, sycl::range{1});
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only);
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
+      CGH.single_task(props, [=]() {
+        auto *Ptr = StorageAcc[0].construct</* ret type = */ Base>(TestCase);
+        applyOp(DataAcc.get_multi_ptr<sycl::access::decorated::no>().get(),
+                Ptr);
+      });
+    });
+
+    Base *Ptr = HostStorage.construct</* ret type = */ Base>(TestCase);
+    applyOp(&HostData, Ptr);
+
+    sycl::host_accessor HostAcc(DataStorage);
+    assert(HostAcc[0] == HostData);
+  }
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}

sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp

@@ -0,0 +1,93 @@
+// UNSUPPORTED: cuda, hip, acc
+// FIXME: replace unsupported with an aspect check once we have it
+//
+// RUN: %{build} -o %t.out -Xclang -fsycl-allow-virtual-functions %helper-includes
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+
+#include "helpers.hpp"
+
+#include <iostream>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class AbstractOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void applyOp(int *) = 0;
+};
+
+class IncrementOp : public AbstractOp {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void applyOp(int *Data) final override { increment(Data); }
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void increment(int *) = 0;
+};
+
+class IncrementBy1 : public IncrementOp {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 1; }
+};
+
+class IncrementBy2 : public IncrementOp {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 2; }
+};
+
+class IncrementBy4 : public IncrementOp {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 4; }
+};
+
+class IncrementBy8 : public IncrementOp {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 8; }
+};
+
+void applyOp(int *Data, AbstractOp *Obj) { Obj->applyOp(Data); }
+
+int main() try {
+  using storage_t =
+      obj_storage_t<IncrementBy1, IncrementBy2, IncrementBy4, IncrementBy8>;
+
+  storage_t HostStorage;
+  sycl::buffer<storage_t> DeviceStorage(sycl::range{1});
+
+  auto asyncHandler = [](sycl::exception_list list) {
+    for (auto &e : list)
+      std::rethrow_exception(e);
+  };
+
+  sycl::queue q(asyncHandler);
+
+  constexpr oneapi::properties props{oneapi::calls_indirectly<>};
+  for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
+    int HostData = 42;
+    int Data = HostData;
+    sycl::buffer<int> DataStorage(&Data, sycl::range{1});
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only);
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
+      CGH.single_task(props, [=]() {
+        auto *Ptr =
+            StorageAcc[0].construct</* ret type = */ AbstractOp>(TestCase);
+        applyOp(DataAcc.get_multi_ptr<sycl::access::decorated::no>().get(),
+                Ptr);
+      });
+    });
+
+    auto *Ptr = HostStorage.construct</* ret type = */ AbstractOp>(TestCase);
+    Ptr->applyOp(&HostData);
+
+    sycl::host_accessor HostAcc(DataStorage);
+    assert(HostAcc[0] == HostData);
+  }
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}

sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp

@@ -0,0 +1,78 @@
+// UNSUPPORTED: cuda, hip, acc
+// FIXME: replace unsupported with an aspect check once we have it
+//
+// RUN: %{build} -o %t.out -Xclang -fsycl-allow-virtual-functions %helper-includes
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+
+#include "helpers.hpp"
+
+#include <iostream>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class BaseIncrement {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void increment(int *Data) { *Data += 1; }
+};
+
+class IncrementBy2 : public BaseIncrement {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 2; }
+};
+
+class IncrementBy4 : public BaseIncrement {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 4; }
+};
+
+class IncrementBy8 : public BaseIncrement {
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 8; }
+};
+
+int main() try {
+  using storage_t =
+      obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4, IncrementBy8>;
+
+  storage_t HostStorage;
+  sycl::buffer<storage_t> DeviceStorage(sycl::range{1});
+
+  auto asyncHandler = [](sycl::exception_list list) {
+    for (auto &e : list)
+      std::rethrow_exception(e);
+  };
+
+  sycl::queue q(asyncHandler);
+
+  constexpr oneapi::properties props{oneapi::calls_indirectly<>};
+  for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
+    int HostData = 42;
+    int Data = HostData;
+    sycl::buffer<int> DataStorage(&Data, sycl::range{1});
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only);
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
+      CGH.single_task(props, [=]() {
+        auto *Ptr =
+            StorageAcc[0].construct</* ret type = */ BaseIncrement>(TestCase);
+        Ptr->increment(
+            DataAcc.get_multi_ptr<sycl::access::decorated::no>().get());
+      });
+    });
+
+    auto *Ptr = HostStorage.construct</* ret type = */ BaseIncrement>(TestCase);
+    Ptr->increment(&HostData);
+
+    sycl::host_accessor HostAcc(DataStorage);
+    assert(HostAcc[0] == HostData);
+  }
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}

sycl/test-e2e/VirtualFunctions/README.md

@@ -0,0 +1,7 @@
+# E2E tests for `sycl_ext_oneapi_virtual_functions` extension
+
+Note about naming convention and files organization for this folder: the tests,
+files and directories are named and organized in a way that resembles their
+description in the corresponding test plan document: link to be inserted here
+later, but for now look into
+[intel/llvm#10540](https://github.com/intel/llvm/pull/10540) PR.

sycl/test-e2e/VirtualFunctions/helpers.hpp

@@ -0,0 +1,53 @@
+#include <algorithm>
+#include <type_traits>
+
+// TODO: strictly speaking, selecting a max alignment here may not be always
+// valid, but for test cases that we have now we expect alignment of all types
+// to be the same.
+// std::aligned_storage uses double under the hood which prevents us from
+// using it on some HW. Therefore we use a custom implementation.
+template <typename... T> struct aligned_storage {
+  static constexpr size_t Len = std::max({sizeof(T)...});
+  static constexpr size_t Align = std::max({alignof(T)...});
+
+  struct type {
+    alignas(Align) unsigned char data[Len];
+  };
+};
+
+// Helper data structure that automatically creates a right (in terms of size
+// and alignment) storage to accomodate a value of any of types T...
+template <typename... T> struct obj_storage_t {
+  static_assert(std::max({alignof(T)...}) == std::min({alignof(T)...}),
+                "Unsupported alignment of input types");
+  using type = typename aligned_storage<T...>::type;
+  static constexpr size_t size = std::max({sizeof(T)...});
+
+  type storage;
+
+  template <typename RetT> RetT *construct(const unsigned int TypeIndex) {
+    if (TypeIndex >= sizeof...(T)) {
+#ifndef __SYCL_DEVICE_ONLY__
+      assert(false && "Type index is invalid");
+#endif
+      return nullptr;
+    }
+
+    return constructHelper<RetT, T...>(TypeIndex, 0);
+  }
+
+private:
+  template <typename RetT> RetT *constructHelper(const int, const int) {
+    // Won't be ever called, but required to compile
+    return nullptr;
+  }
+
+  template <typename RetT, typename Type, typename... Rest>
+  RetT *constructHelper(const int TargetIndex, const int CurIndex) {
+    if (TargetIndex != CurIndex)
+      return constructHelper<RetT, Rest...>(TargetIndex, CurIndex + 1);
+
+    RetT *Ptr = new (reinterpret_cast<Type *>(&storage)) Type;
+    return Ptr;
+  }
+};

sycl/test-e2e/VirtualFunctions/lit.local.cfg

@@ -0,0 +1,6 @@
+import os
+
+# Tests are sharing some common header, but we don't won't to use relative
+# paths like "../../../helper.hpp" in them, so let's just register a
+# substitution to add directory with helper headers into include search path
+config.substitutions.append(("%helper-includes", "-I {}".format(os.path.dirname(os.path.abspath(__file__)))))