[SYCL][E2E] Add re-mapping virtual memory range test for sycl_ext_oneapi_virtual_mem extension

sycl/test-e2e/VirtualMem/helpers.hpp

@@ -0,0 +1,31 @@
+#pragma once
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/virtual_mem/physical_mem.hpp>
+#include <sycl/ext/oneapi/virtual_mem/virtual_mem.hpp>
+
+namespace syclext = sycl::ext::oneapi::experimental;
+
+// Find the least common multiple of the context and device granularities. This
+// value can be used for aligning both physical memory allocations and for
+// reserving virtual memory ranges.
+size_t GetLCMGranularity(
+    const sycl::device &Dev, const sycl::context &Ctx,
+    syclext::granularity_mode Gm = syclext::granularity_mode::recommended) {
+  size_t CtxGranularity = syclext::get_mem_granularity(Ctx, Gm);
+  size_t DevGranularity = syclext::get_mem_granularity(Dev, Ctx, Gm);
+
+  size_t GCD = CtxGranularity;
+  size_t Rem = DevGranularity % GCD;
+  while (Rem != 0) {
+    std::swap(GCD, Rem);
+    Rem %= GCD;
+  }
+  return (DevGranularity / GCD) * CtxGranularity;
+}
+
+size_t GetAlignedByteSize(const size_t UnalignedBytes,
+                          const size_t AligmentGranularity) {
+  return ((UnalignedBytes + AligmentGranularity - 1) / AligmentGranularity) *
+         AligmentGranularity;
+}

sycl/test-e2e/VirtualMem/remapping_virtual_memory_range.cpp

@@ -0,0 +1,82 @@
+// This test checks whether virtual memory range can correctly be accessed
+// even if it was re-mapped to a different physical range.
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+
+#include <cassert>
+
+#include "helpers.hpp"
+
+namespace syclext = sycl::ext::oneapi::experimental;
+
+int main() {
+
+  sycl::queue Q;
+  sycl::context Context = Q.get_context();
+  sycl::device Device = Q.get_device();
+
+  int Failed = 0;
+
+  constexpr size_t NumberOfElements = 1000;
+  constexpr int ValueSetInFirstKernel = 555;
+  constexpr int ValueSetInSecondKernel = 999;
+
+  size_t BytesRequired = NumberOfElements * sizeof(int);
+
+  size_t UsedGranularity = GetLCMGranularity(Device, Context);
+  size_t AlignedByteSize = GetAlignedByteSize(BytesRequired, UsedGranularity);
+
+  syclext::physical_mem FirstPhysicalMemory{Device, Context, AlignedByteSize};
+  uintptr_t VirtualMemoryPtr =
+      syclext::reserve_virtual_mem(0, AlignedByteSize, Context);
+
+  void *MappedPtr =
+      FirstPhysicalMemory.map(VirtualMemoryPtr, AlignedByteSize,
+                              syclext::address_access_mode::read_write);
+
+  int *DataPtr = reinterpret_cast<int *>(MappedPtr);
+
+  std::vector<int> ResultHostData(NumberOfElements);
+
+  Q.parallel_for(NumberOfElements, [=](sycl::id<1> Idx) {
+     DataPtr[Idx] = ValueSetInFirstKernel;
+   }).wait_and_throw();
+
+  syclext::unmap(MappedPtr, AlignedByteSize, Context);
+
+  syclext::physical_mem SecondPhysicalMemory{Device, Context, AlignedByteSize};
+  MappedPtr =
+      SecondPhysicalMemory.map(VirtualMemoryPtr, AlignedByteSize,
+                               syclext::address_access_mode::read_write);
+
+  Q.parallel_for(NumberOfElements, [=](sycl::id<1> Idx) {
+     DataPtr[Idx] = ValueSetInSecondKernel;
+   }).wait_and_throw();
+
+  {
+    sycl::buffer<int> ResultBuffer(ResultHostData);
+
+    Q.submit([&](sycl::handler &Handle) {
+      sycl::accessor A(ResultBuffer, Handle, sycl::write_only);
+      Handle.parallel_for(NumberOfElements,
+                          [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; });
+    });
+  }
+
+  for (size_t i = 0; i < NumberOfElements; i++) {
+    if (ResultHostData[i] != ValueSetInSecondKernel) {
+      std::cout << "Comparison failed at index " << i << ": "
+                << ResultHostData[i] << " != " << ValueSetInSecondKernel
+                << std::endl;
+      ++Failed;
+    }
+  }
+
+  syclext::unmap(MappedPtr, AlignedByteSize, Context);
+  syclext::free_virtual_mem(VirtualMemoryPtr, AlignedByteSize, Context);
+
+  return Failed;
+}