[SYCL][E2E] Add more tests for virtual functions (#15067)

This commit still doesn't bring an exhaustive coverage for the feature,
but still improves the situation by checking the following scenarios:
- using math built-ins from virtual functions
- using group barriers from virtual functions
- using virtual functions in nd-range kernels where every work-item
calls a different virtual function
- using virtual functions when the code is scattered across several
translation units

Some tests are disabled, because we do not support those scenarios yet
and more changes are required to make them work.

sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp

@@ -0,0 +1,162 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// On CPU it segfaults within the kernel that performs virtual function call.
+// XFAIL: cpu
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/15080
+// UNSUPPORTED: gpu
+// On GPU this test (its older version which used nd_item instead of group)
+// used to fail with UR_RESULT_ERROR_PROGRAM_LINK_FAILURE.
+// SPIR-V files produced by SYCL_DUMP_IMAGES could be linked just fine (using
+// both llvm-spirv -r + llvm-link and ocloc).
+// Current version hangs and therefore it is marked as unsupported to avoid
+// wasting time in CI and potentially blocking a machine.
+// Reported in https://github.com/intel/llvm/issues/15068
+//
+// This test checks that group operations (barrier in this case) work correctly
+// inside virtual functions.
+//
+// RUN: %{build} -o %t.out %helper-includes
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/group_algorithm.hpp>
+#include <sycl/group_barrier.hpp>
+#include <sycl/usm.hpp>
+
+#include "helpers.hpp"
+
+#include <iostream>
+#include <numeric>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual int apply(int *, sycl::group<1>) = 0;
+
+  virtual int computeReference(sycl::range<1> LocalRange, int Init) = 0;
+};
+
+class SumOp : public BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  int apply(int *LocalData, sycl::group<1> WG) override {
+    LocalData[WG.get_local_id()] = WG.get_local_id() + WG.get_group_id();
+    sycl::group_barrier(WG);
+    if (WG.leader()) {
+      int Res = 0;
+      for (size_t I = 0; I < WG.get_local_range().size(); ++I) {
+        Res += LocalData[I];
+      }
+      LocalData[0] = Res;
+    }
+    sycl::group_barrier(WG);
+
+    return LocalData[0];
+  }
+
+  int computeReference(sycl::range<1> LocalRange, int WGID) override {
+    std::vector<int> LocalData(LocalRange.size());
+    for (size_t LID = 0; LID < LocalRange.size(); ++LID)
+      LocalData[LID] = LID + WGID;
+
+    int Res = 0;
+    for (size_t LID = 0; LID < LocalRange.size(); ++LID)
+      Res += LocalData[LID];
+
+    return Res;
+  }
+};
+
+class MultiplyOp : public BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  int apply(int *LocalData, sycl::group<1> WG) override {
+    // +1 to avoid multiplying by 0 below
+    LocalData[WG.get_local_id()] = WG.get_local_id() + WG.get_group_id() + 1;
+    sycl::group_barrier(WG);
+    if (WG.leader()) {
+      int Res = 1;
+      for (size_t I = 0; I < WG.get_local_range().size(); ++I) {
+        Res *= LocalData[I];
+      }
+      LocalData[0] = Res;
+    }
+    sycl::group_barrier(WG);
+
+    return LocalData[0];
+  }
+
+  int computeReference(sycl::range<1> LocalRange, int WGID) override {
+    std::vector<int> LocalData(LocalRange.size());
+    for (size_t LID = 0; LID < LocalRange.size(); ++LID)
+      LocalData[LID] = LID + WGID + 1;
+
+    int Res = 1;
+    for (size_t LID = 0; LID < LocalRange.size(); ++LID)
+      Res *= LocalData[LID];
+
+    return Res;
+  }
+};
+
+int main() try {
+  using storage_t = obj_storage_t<SumOp, MultiplyOp>;
+
+  sycl::queue q;
+
+  storage_t HostStorage;
+  auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
+  // Let's keep ranges small, or otherwise we will encounter integer overflow
+  // (which is a UB) in MultiplyOp::apply.
+  sycl::range G{16};
+  sycl::range L{4};
+
+  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+  for (unsigned TestCase = 0; TestCase < 2; ++TestCase) {
+    sycl::buffer<int> DataStorage(G);
+
+    q.submit([&](sycl::handler &CGH) {
+       CGH.single_task([=]() {
+         DeviceStorage->construct</* ret type = */ BaseOp>(TestCase);
+       });
+     }).wait_and_throw();
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
+      sycl::local_accessor<int> LocalAcc(L, CGH);
+      CGH.parallel_for(sycl::nd_range{G, L}, props, [=](auto It) {
+        auto *Ptr = DeviceStorage->getAs<BaseOp>();
+        DataAcc[It.get_global_id()] = Ptr->apply(
+            LocalAcc.get_multi_ptr<sycl::access::decorated::no>().get(),
+            It.get_group());
+      });
+    }).wait_and_throw();
+
+    auto *Ptr = HostStorage.construct</* ret type = */ BaseOp>(TestCase);
+    sycl::host_accessor HostAcc(DataStorage);
+
+    // All work-items in a group produce the same result, so we do verification
+    // per work-group.
+    for (size_t WorkGroupID = 0; WorkGroupID < G.size() / L.size();
+         ++WorkGroupID) {
+      int Reference = Ptr->computeReference(L, WorkGroupID);
+      for (size_t I = 0; I < L.size(); ++I) {
+        size_t GID = WorkGroupID * L.size() + I;
+        if (HostAcc[GID] != Reference) {
+          std::cout << "Mismatch at index " << I << ": " << HostAcc[I]
+                    << " != " << Reference << std::endl;
+          assert(HostAcc[I] == Reference);
+        }
+      }
+    }
+  }
+
+  sycl::free(DeviceStorage, q);
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}

sycl/test-e2e/VirtualFunctions/misc/math.cpp

@@ -0,0 +1,85 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// This test checks that SYCL math built-in functions work correctly
+// inside virtual functions.
+//
+// RUN: %{build} -o %t.out %helper-includes
+// RUN: %{run} %t.out
+
+#include <sycl/builtins.hpp>
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include "helpers.hpp"
+
+#include <iostream>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual float apply(float) = 0;
+};
+
+class FloorOp : public BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual float apply(float V) { return sycl::floor(V); }
+};
+
+class CeilOp : public BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual float apply(float V) { return sycl::ceil(V); }
+};
+
+class RoundOp : public BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual float apply(float V) { return sycl::round(V); }
+};
+
+int main() try {
+  using storage_t = obj_storage_t<FloorOp, CeilOp, RoundOp>;
+
+  storage_t HostStorage;
+
+  sycl::queue q;
+
+  auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
+
+  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+  for (unsigned TestCase = 0; TestCase < 3; ++TestCase) {
+    float HostData = 3.56;
+    float Data = HostData;
+    sycl::buffer<float> DataStorage(&Data, sycl::range{1});
+
+    q.submit([&](sycl::handler &CGH) {
+       CGH.single_task([=]() {
+         DeviceStorage->construct</* ret type = */ BaseOp>(TestCase);
+       });
+     }).wait_and_throw();
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
+      CGH.single_task(props, [=]() {
+        auto *Ptr = DeviceStorage->getAs<BaseOp>();
+        DataAcc[0] = Ptr->apply(DataAcc[0]);
+      });
+    });
+
+    auto *Ptr = HostStorage.construct</* ret type = */ BaseOp>(TestCase);
+    HostData = Ptr->apply(HostData);
+
+    sycl::host_accessor HostAcc(DataStorage);
+    assert(HostAcc[0] == HostData);
+  }
+
+  sycl::free(DeviceStorage, q);
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}

sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp

@@ -0,0 +1,102 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// This test checks that virtual functions work correctly in simple range
+// kernels when different work-items perform calls to different virtual
+// functions using the same object.
+//
+// RUN: %{build} -o %t.out %helper-includes
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include "helpers.hpp"
+
+#include <iostream>
+#include <numeric>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual int foo(int) = 0;
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual int bar(int) = 0;
+};
+
+class OpA : public BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual int foo(int V) { return V + 2; }
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual int bar(int V) { return V - 2; }
+};
+
+class OpB : public BaseOp {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual int foo(int V) { return V * 2; }
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+  virtual int bar(int V) { return V / 2; }
+};
+
+int main() try {
+  using storage_t = obj_storage_t<OpA, OpB>;
+
+  storage_t HostStorage;
+
+  sycl::queue q;
+
+  auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
+  sycl::range R{1024};
+
+  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+  for (size_t TestCase = 0; TestCase < 2; ++TestCase) {
+    std::vector<int> HostData(R.size());
+    std::iota(HostData.begin(), HostData.end(), 0);
+    std::vector<int> DeviceData = HostData;
+    sycl::buffer<int> DataStorage(DeviceData.data(), R);
+
+    q.submit([&](sycl::handler &CGH) {
+       CGH.single_task([=]() {
+         DeviceStorage->construct</* ret type = */ BaseOp>(TestCase);
+       });
+     }).wait_and_throw();
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
+      CGH.parallel_for(R, props, [=](auto It) {
+        // Select method that corresponds to this work-item
+        auto *Ptr = DeviceStorage->template getAs<BaseOp>();
+        if (It % 2)
+          DataAcc[It] = Ptr->foo(DataAcc[It]);
+        else
+          DataAcc[It] = Ptr->bar(DataAcc[It]);
+      });
+    });
+
+    BaseOp *Ptr = HostStorage.construct</* ret type = */ BaseOp>(TestCase);
+
+    for (size_t I = 0; I < HostData.size(); ++I) {
+      if (I % 2)
+        HostData[I] = Ptr->foo(HostData[I]);
+      else
+        HostData[I] = Ptr->bar(HostData[I]);
+    }
+
+    sycl::host_accessor HostAcc(DataStorage);
+    for (size_t I = 0; I < HostData.size(); ++I)
+      assert(HostAcc[I] == HostData[I]);
+  }
+
+  sycl::free(DeviceStorage, q);
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}