[SYCL] Add one more test for virtual functions (#14705)

Test plan can be found in #10540. This PR introduces a basic test case
for a scenario where object of a polymorphic class is constructed in one
kernel, but used in another kernel.

Author: AlexeySachkov

sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp

@@ -0,0 +1,105 @@
+// UNSUPPORTED: cuda, hip, acc
+// FIXME: replace unsupported with an aspect check once we have it
+//
+// RUN: %{build} -o %t.out %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:
+  BaseIncrement(int Mod, int /* unused */ = 42) : Mod(Mod) {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void increment(int *Data) { *Data += 1 + Mod; }
+
+protected:
+  int Mod = 0;
+};
+
+class IncrementBy2 : public BaseIncrement {
+public:
+  IncrementBy2(int Mod, int /* unused */) : BaseIncrement(Mod) {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 2 + Mod; }
+};
+
+class IncrementBy4 : public BaseIncrement {
+public:
+  IncrementBy4(int Mod, int ExtraMod)
+      : BaseIncrement(Mod), ExtraMod(ExtraMod) {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 4 + Mod + ExtraMod; }
+
+private:
+  int ExtraMod = 0;
+};
+
+class IncrementBy8 : public BaseIncrement {
+public:
+  IncrementBy8(int Mod, int /* unused */) : BaseIncrement(Mod) {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void increment(int *Data) override { *Data += 8 + Mod; }
+};
+
+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);
+
+  // TODO: cover uses case when objects are passed through USM
+  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);
+      CGH.single_task([=]() {
+        StorageAcc[0].construct</* ret type = */ BaseIncrement>(TestCase, 19,
+                                                                23);
+      });
+    });
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::read_write);
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
+      CGH.single_task(props, [=]() {
+        auto *Ptr = StorageAcc[0].getAs<BaseIncrement>();
+        Ptr->increment(
+            DataAcc.get_multi_ptr<sycl::access::decorated::no>().get());
+      });
+    });
+
+    auto *Ptr =
+        HostStorage.construct</* ret type = */ BaseIncrement>(TestCase, 19, 23);
+    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/helpers.hpp

@@ -25,29 +25,36 @@ template <typename... T> struct obj_storage_t {
 
   type storage;
 
-  template <typename RetT> RetT *construct(const unsigned int TypeIndex) {
+  template <typename RetT, typename... Args>
+  RetT *construct(const unsigned int TypeIndex, Args... args) {
     if (TypeIndex >= sizeof...(T)) {
 #ifndef __SYCL_DEVICE_ONLY__
       assert(false && "Type index is invalid");
 #endif
       return nullptr;
     }
 
-    return constructHelper<RetT, T...>(TypeIndex, 0);
+    return constructHelper<RetT, T...>(TypeIndex, 0, args...);
+  }
+
+  template <typename RetT> RetT *getAs() {
+    return reinterpret_cast<RetT *>(&storage);
   }
 
 private:
-  template <typename RetT> RetT *constructHelper(const int, const int) {
+  template <typename RetT, typename... Args>
+  RetT *constructHelper(const int, const int, Args...) {
     // 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) {
+  template <typename RetT, typename Type, typename... Rest, typename... Args>
+  RetT *constructHelper(const int TargetIndex, const int CurIndex,
+                        Args... args) {
     if (TargetIndex != CurIndex)
-      return constructHelper<RetT, Rest...>(TargetIndex, CurIndex + 1);
+      return constructHelper<RetT, Rest...>(TargetIndex, CurIndex + 1, args...);
 
-    RetT *Ptr = new (reinterpret_cast<Type *>(&storage)) Type;
+    RetT *Ptr = new (reinterpret_cast<Type *>(&storage)) Type(args...);
     return Ptr;
   }
 };