[SYCL][E2E] Add more tests for virtual functions

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;
+}