diff --git a/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp b/sycl/test-e2e/VirtualFunctions/misc/group-barrier.cpp
new file mode 100644
index 0000000000000..a69aada295c25
--- /dev/null
+++ b/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;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/misc/math.cpp b/sycl/test-e2e/VirtualFunctions/misc/math.cpp
new file mode 100644
index 0000000000000..71b34c23cef1f
--- /dev/null
+++ b/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;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp
new file mode 100644
index 0000000000000..45b56916a5c1d
--- /dev/null
+++ b/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;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp
new file mode 100644
index 0000000000000..453a3aee81fa6
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp
@@ -0,0 +1,98 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// This test checks that virtual functions work correctly within simple range
+// kernels when different work-items perform a virtual function calls using
+// different objects.
+//
+// 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>;
+
+  std::array<storage_t, 3> HostStorage;
+
+  sycl::queue q;
+
+  auto *DeviceStorage = sycl::malloc_shared<storage_t>(3, q);
+  sycl::range R{1024};
+
+  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+  {
+    std::vector<float> HostData(R.size());
+    for (size_t I = 1; I < HostData.size(); ++I)
+      HostData[I] = HostData[I - 1] + 0.7;
+    std::vector<float> DeviceData = HostData;
+    sycl::buffer<float> DataStorage(DeviceData.data(), R);
+
+    q.submit([&](sycl::handler &CGH) {
+       CGH.single_task([=]() {
+         DeviceStorage[0].construct</* ret type = */ BaseOp>(0);
+         DeviceStorage[1].construct</* ret type = */ BaseOp>(1);
+         DeviceStorage[2].construct</* ret type = */ BaseOp>(2);
+       });
+     }).wait_and_throw();
+
+    q.submit([&](sycl::handler &CGH) {
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
+      CGH.parallel_for(R, props, [=](auto it) {
+        // Select an object that corresponds to this work-item
+        auto Ind = it % 3;
+        auto *Ptr = DeviceStorage[Ind].template getAs<BaseOp>();
+        DataAcc[it] = Ptr->apply(DataAcc[it]);
+      });
+    });
+
+    BaseOp *Ptr[] = {HostStorage[0].construct</* ret type = */ BaseOp>(0),
+                     HostStorage[1].construct</* ret type = */ BaseOp>(1),
+                     HostStorage[2].construct</* ret type = */ BaseOp>(2)};
+
+    for (size_t I = 0; I < HostData.size(); ++I)
+      HostData[I] = Ptr[I % 3]->apply(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;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp b/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp
new file mode 100644
index 0000000000000..66db6a0c5af7a
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp
@@ -0,0 +1,91 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// This test checks that virtual functions work correctly within simple range
+// kernels when every work-item calls the same virtual function on the same
+// object.
+//
+// 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);
+  sycl::range R{1024};
+
+  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+  for (unsigned TestCase = 0; TestCase < 3; ++TestCase) {
+    std::vector<float> HostData(R.size());
+    for (size_t I = 1; I < HostData.size(); ++I)
+      HostData[I] = HostData[I - 1] + 0.7;
+    std::vector<float> DeviceData = HostData;
+    sycl::buffer<float> 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) {
+        auto *Ptr = DeviceStorage->getAs<BaseOp>();
+        DataAcc[it] = Ptr->apply(DataAcc[it]);
+      });
+    });
+
+    auto *Ptr = HostStorage.construct</* ret type = */ BaseOp>(TestCase);
+    for (size_t I = 0; I < HostData.size(); ++I)
+      HostData[I] = Ptr->apply(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;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp
new file mode 100644
index 0000000000000..9ce59931405d6
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp
@@ -0,0 +1,19 @@
+#include "declarations.hpp"
+
+int call(sycl::queue Q, storage_t *DeviceStorage, int Init) {
+  int Data = Init;
+  {
+    sycl::buffer<int> DataStorage(&Data, sycl::range{1});
+    constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+    Q.submit([&](sycl::handler &CGH) {
+      sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
+      CGH.single_task(props, [=]() {
+        auto *Ptr = DeviceStorage->getAs<BaseIncrement>();
+        Ptr->increment(
+            DataAcc.get_multi_ptr<sycl::access::decorated::no>().get());
+      });
+    });
+  }
+
+  return Data;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/construct.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/construct.cpp
new file mode 100644
index 0000000000000..ae980b307a1ae
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/construct.cpp
@@ -0,0 +1,10 @@
+#include "declarations.hpp"
+
+void construct(sycl::queue Q, storage_t *DeviceStorage, unsigned TestCase) {
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.single_task([=]() {
+       DeviceStorage->construct</* ret type = */ BaseIncrement>(TestCase, 19,
+                                                                23);
+     });
+   }).wait_and_throw();
+}
diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/declarations.hpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/declarations.hpp
new file mode 100644
index 0000000000000..a03eaeb66d9ed
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/declarations.hpp
@@ -0,0 +1,53 @@
+#pragma once
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include "helpers.hpp"
+
+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);
+
+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;
+};
+
+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;
+
+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;
+};
+
+using storage_t =
+    obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4, IncrementBy8>;
+
+void construct(sycl::queue Q, storage_t *DeviceStorage, unsigned TestCase);
+int call(sycl::queue Q, storage_t *DeviceStorage, int Init);
diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/vf.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/vf.cpp
new file mode 100644
index 0000000000000..73eb7fb91744a
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/vf.cpp
@@ -0,0 +1,13 @@
+#include "declarations.hpp"
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+void BaseIncrement::increment(int *Data) { *Data += 1 + Mod; }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+void IncrementBy2::increment(int *Data) { *Data += 2 + Mod; }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+void IncrementBy4::increment(int *Data) { *Data += 4 + Mod; }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable)
+void IncrementBy8::increment(int *Data) { *Data += 8 + Mod; }
diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-call.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-call.cpp
new file mode 100644
index 0000000000000..ac6cfc234567d
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-call.cpp
@@ -0,0 +1,55 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// VTables are global variables with possibly external linkage and that causes
+// them to be copied into every module we produce during device code split
+// which in turn leads to multiple definitions error at runtime.
+// https://github.com/intel/llvm/issues/15069
+// XFAIL: *
+//
+// This test covers a scenario where virtual functions defintion and their uses
+// are split into different translation units. In particular:
+// - both virtual functions and construct kernel are in the same translation
+//   unit
+// - but use kernel is outlined into a separate translation unit
+//
+// RUN: %{build} %S/Inputs/call.cpp -o %t.out %helper-includes
+// RUN: %{run} %t.out
+
+#include "Inputs/declarations.hpp"
+
+#include <iostream>
+
+#include "Inputs/construct.cpp"
+#include "Inputs/vf.cpp"
+
+int main() try {
+  storage_t HostStorage;
+
+  auto asyncHandler = [](sycl::exception_list list) {
+    for (auto &e : list)
+      std::rethrow_exception(e);
+  };
+
+  sycl::queue q(asyncHandler);
+  auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
+
+  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+  for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
+    int HostData = 42;
+    construct(q, DeviceStorage, TestCase);
+    int Result = call(q, DeviceStorage, HostData);
+
+    auto *Ptr =
+        HostStorage.construct</* ret type = */ BaseIncrement>(TestCase, 19, 23);
+    Ptr->increment(&HostData);
+
+    assert(Result == HostData);
+  }
+
+  sycl::free(DeviceStorage, q);
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-vf-defs-and-call.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-vf-defs-and-call.cpp
new file mode 100644
index 0000000000000..ec24d395c92fd
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-vf-defs-and-call.cpp
@@ -0,0 +1,52 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// We attach calls-indirectly attribute (and therefore device image property)
+// to construct kernels at compile step. At that stage we may not see virtual
+// function definitions and therefore we won't mark construct kernel as using
+// virtual functions and link operation at runtime will fail due to undefined
+// references to virtual functions from vtable.
+// https://github.com/intel/llvm/issues/15071
+// XFAIL: *
+//
+// This test covers a scenario where virtual functions defintion and their uses
+// are all split into different translation units.
+//
+// RUN: %{build} %S/Inputs/call.cpp %S/Inputs/vf.cpp -o %t.out %helper-includes
+// RUN: %{run} %t.out
+
+#include "Inputs/declarations.hpp"
+
+#include <iostream>
+
+#include "Inputs/construct.cpp"
+
+int main() try {
+  storage_t HostStorage;
+
+  auto asyncHandler = [](sycl::exception_list list) {
+    for (auto &e : list)
+      std::rethrow_exception(e);
+  };
+
+  sycl::queue q(asyncHandler);
+  auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
+
+  for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
+    int HostData = 42;
+    construct(q, DeviceStorage, TestCase);
+    int Result = call(q, DeviceStorage, HostData);
+
+    auto *Ptr =
+        HostStorage.construct</* ret type = */ BaseIncrement>(TestCase, 19, 23);
+    Ptr->increment(&HostData);
+
+    assert(Result == HostData);
+  }
+
+  sycl::free(DeviceStorage, q);
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}
diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-vf-defs.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-vf-defs.cpp
new file mode 100644
index 0000000000000..99645973593e8
--- /dev/null
+++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/separate-vf-defs.cpp
@@ -0,0 +1,56 @@
+// REQUIRES: aspect-usm_shared_allocations
+//
+// We attach calls-indirectly attribute (and therefore device image property)
+// to construct kernels at compile step. At that stage we may not see virtual
+// function definitions and therefore we won't mark construct kernel as using
+// virtual functions and link operation at runtime will fail due to undefined
+// references to virtual functions from vtable.
+// https://github.com/intel/llvm/issues/15071
+// XFAIL: *
+//
+// This test covers a scenario where virtual functions defintion and their uses
+// are split into different translation units. In particular:
+// - use and construct kernesl are in the same translation unit
+// - but virtual functions are defined in a separate translation unit
+//
+// RUN: %{build} %S/Inputs/vf.cpp -o %t.out %helper-includes
+// RUN: %{run} %t.out
+
+#include "Inputs/declarations.hpp"
+
+#include <iostream>
+
+#include "Inputs/call.cpp"
+#include "Inputs/construct.cpp"
+
+int main() try {
+  storage_t HostStorage;
+
+  auto asyncHandler = [](sycl::exception_list list) {
+    for (auto &e : list)
+      std::rethrow_exception(e);
+  };
+
+  sycl::queue q(asyncHandler);
+  auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
+
+  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
+  for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
+    int HostData = 42;
+    construct(q, DeviceStorage, TestCase);
+    int Result = call(q, DeviceStorage, HostData);
+
+    auto *Ptr =
+        HostStorage.construct</* ret type = */ BaseIncrement>(TestCase, 19, 23);
+    Ptr->increment(&HostData);
+
+    assert(Result == HostData);
+  }
+
+  sycl::free(DeviceStorage, q);
+
+  return 0;
+} catch (sycl::exception &e) {
+  std::cout << "Unexpected exception was thrown: " << e.what() << std::endl;
+  return 1;
+}