diff --git a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp new file mode 100644 index 0000000000000..4ad528b40d965 --- /dev/null +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -0,0 +1,86 @@ +// This test checks whether data can be correctly written to and read from +// virtual memory. + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include "helpers.hpp" + +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; + size_t BytesRequired = NumberOfElements * sizeof(int); + + size_t UsedGranularity = GetLCMGranularity(Device, Context); + + size_t AlignedByteSize = + ((BytesRequired + UsedGranularity - 1) / UsedGranularity) * + UsedGranularity; + + syclext::physical_mem NewPhysicalMem{Device, Context, AlignedByteSize}; + uintptr_t VirtualMemoryPtr = + syclext::reserve_virtual_mem(0, AlignedByteSize, Context); + + void *MappedPtr = + NewPhysicalMem.map(VirtualMemoryPtr, AlignedByteSize, + syclext::address_access_mode::read_write); + + int *DataPtr = reinterpret_cast(MappedPtr); + + std::vector ResultHostData(NumberOfElements); + + constexpr int ExpectedValueAfterFill = 1; + + Q.fill(DataPtr, ExpectedValueAfterFill, NumberOfElements).wait_and_throw(); + { + sycl::buffer CheckBuffer(ResultHostData); + Q.submit([&](sycl::handler &Handle) { + sycl::accessor A(CheckBuffer, Handle, sycl::write_only); + Handle.parallel_for(NumberOfElements, + [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); + }); + } + + for (size_t i = 0; i < ResultHostData.size(); i++) { + if (ResultHostData[i] != ExpectedValueAfterFill) { + std::cout << "Comparison failed after fill operation at index " << i + << ": " << ResultHostData[i] << " != " << ExpectedValueAfterFill + << std::endl; + ++Failed; + } + } + + Q.parallel_for(NumberOfElements, [=](sycl::id<1> Idx) { + DataPtr[Idx] = Idx; + }).wait_and_throw(); + + syclext::set_access_mode(DataPtr, AlignedByteSize, + syclext::address_access_mode::read, Context); + + { + sycl::buffer 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++) { + const int ExpectedValue = static_cast(i); + if (ResultHostData[i] != ExpectedValue) { + std::cout << "Comparison failed at index " << i << ": " + << ResultHostData[i] << " != " << ExpectedValue << std::endl; + ++Failed; + } + } + + syclext::unmap(MappedPtr, AlignedByteSize, Context); + syclext::free_virtual_mem(VirtualMemoryPtr, AlignedByteSize, Context); + + return Failed; +} diff --git a/sycl/test-e2e/VirtualMem/remapping_virtual_memory_range.cpp b/sycl/test-e2e/VirtualMem/remapping_virtual_memory_range.cpp index 07678a180a78b..31817cb00314e 100644 --- a/sycl/test-e2e/VirtualMem/remapping_virtual_memory_range.cpp +++ b/sycl/test-e2e/VirtualMem/remapping_virtual_memory_range.cpp @@ -4,8 +4,6 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include - #include #include "helpers.hpp" diff --git a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp index 85fbd3fbdabbc..4d80ade89df04 100644 --- a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp @@ -6,30 +6,11 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include #include #include -#include -#include - -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) { - size_t CtxGranularity = syclext::get_mem_granularity(Ctx); - size_t DevGranularity = syclext::get_mem_granularity(Dev, Ctx); - - size_t GCD = CtxGranularity; - size_t Rem = DevGranularity % GCD; - while (Rem != 0) { - std::swap(GCD, Rem); - Rem %= GCD; - } - return (DevGranularity / GCD) * CtxGranularity; -} + +#include "helpers.hpp" template class VirtualVector { public: