From 7247f0047884230155b5cd8c4f3a000e1439889a Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 15:25:20 +0200 Subject: [PATCH 01/14] added virtual mem test for basic access from a kernel --- .../basic_access_from_kernel_virtual_mem.cpp | 79 +++++++++++++++++++ 1 file changed, 79 insertions(+) create mode 100644 sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp 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..7d3f0b7bc696f --- /dev/null +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -0,0 +1,79 @@ +// REQUIRES: aspect-ext_oneapi_virtual_mem + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include +#include +#include + +namespace syclext = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue Q; + sycl::context Context = Q.get_context(); + + size_t UsedGranularityInBytes = syclext::get_mem_granularity(Context, syclext::granularity_mode::recommended); + + syclext::physical_mem NewPhysicalMem{Q.get_device(), Context, UsedGranularityInBytes}; + uintptr_t VirtualMemoryPtr = syclext::reserve_virtual_mem(0, UsedGranularityInBytes, Context); + + void *MappedPtr = NewPhysicalMem.map(VirtualMemoryPtr, UsedGranularityInBytes, syclext::address_access_mode::read_write); + + int* DataPtr = reinterpret_cast(MappedPtr); + + sycl::range NumItems{UsedGranularityInBytes/sizeof(int)}; + + std::vector ResultHostData(NumItems.size()); + + constexpr int ExpectedValueAfterFill = 1; + + Q.fill(DataPtr,ExpectedValueAfterFill , NumItems.size()).wait_and_throw(); + { + sycl::buffer CheckBuffer(ResultHostData); + Q.submit([&](auto &handle) { + sycl::accessor A(CheckBuffer, handle, sycl::write_only); + handle.parallel_for(NumItems, [=](auto i) { A[i] = DataPtr[i]; }); + }); + } + + 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] + << " != " < ResultBuffer(ResultHostData); + + Q.submit([&](auto &handle) { + sycl::accessor A(ResultBuffer, handle, sycl::write_only); + handle.parallel_for(NumItems, [=](auto i) { A[i] = DataPtr[i]; }); + }); + } + + for (size_t i = 0; i < NumItems.size(); i++) { + const int ExpectedValue = static_cast(i); + if (ResultHostData[i] != ExpectedValue) { + std::cout << "Comparison failed at index " << i << ": " << ResultHostData[i] + << " != " << ExpectedValue<< std::endl; + return 1; + } + } + + syclext::unmap(MappedPtr, UsedGranularityInBytes, Context); + syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); + + return 0; +} \ No newline at end of file From 1297895737ec69401244986b8c03e01363edb864 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 15:53:10 +0200 Subject: [PATCH 02/14] fixes --- .../basic_access_from_kernel_virtual_mem.cpp | 26 +++++++++---------- 1 file changed, 13 insertions(+), 13 deletions(-) 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 index 7d3f0b7bc696f..146476c1af726 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -3,6 +3,7 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out +#include "sycl/handler.hpp" #include #include #include @@ -12,6 +13,7 @@ namespace syclext = sycl::ext::oneapi::experimental; int main() { sycl::queue Q; sycl::context Context = Q.get_context(); + int Failed = 0; size_t UsedGranularityInBytes = syclext::get_mem_granularity(Context, syclext::granularity_mode::recommended); @@ -31,9 +33,9 @@ int main() { Q.fill(DataPtr,ExpectedValueAfterFill , NumItems.size()).wait_and_throw(); { sycl::buffer CheckBuffer(ResultHostData); - Q.submit([&](auto &handle) { - sycl::accessor A(CheckBuffer, handle, sycl::write_only); - handle.parallel_for(NumItems, [=](auto i) { A[i] = DataPtr[i]; }); + Q.submit([&](sycl::handler &Handle) { + sycl::accessor A(CheckBuffer, Handle, sycl::write_only); + Handle.parallel_for(NumItems, [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); }); } @@ -41,14 +43,12 @@ int main() { if (ResultHostData[i] != ExpectedValueAfterFill) { std::cout << "Comparison failed after fill operation at index " << i << ": " << ResultHostData[i] << " != " < Idx) { + DataPtr[Idx] = Idx; }).wait_and_throw(); syclext::set_access_mode(DataPtr,UsedGranularityInBytes, syclext::address_access_mode::read, Context); @@ -57,9 +57,9 @@ int main() { { sycl::buffer ResultBuffer(ResultHostData); - Q.submit([&](auto &handle) { - sycl::accessor A(ResultBuffer, handle, sycl::write_only); - handle.parallel_for(NumItems, [=](auto i) { A[i] = DataPtr[i]; }); + Q.submit([&](sycl::handler &Handle) { + sycl::accessor A(ResultBuffer, Handle, sycl::write_only); + Handle.parallel_for(NumItems, [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); }); } @@ -68,12 +68,12 @@ int main() { if (ResultHostData[i] != ExpectedValue) { std::cout << "Comparison failed at index " << i << ": " << ResultHostData[i] << " != " << ExpectedValue<< std::endl; - return 1; + ++Failed; } } syclext::unmap(MappedPtr, UsedGranularityInBytes, Context); syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); - return 0; + return Failed; } \ No newline at end of file From 0c990f937f2b20752f37e52363136639ed8f970e Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 15:57:14 +0200 Subject: [PATCH 03/14] fixed formatting --- .../basic_access_from_kernel_virtual_mem.cpp | 22 ++++++++++--------- 1 file changed, 12 insertions(+), 10 deletions(-) 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 index 146476c1af726..114b3d74a7dfc 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -34,8 +34,9 @@ int main() { { sycl::buffer CheckBuffer(ResultHostData); Q.submit([&](sycl::handler &Handle) { - sycl::accessor A(CheckBuffer, Handle, sycl::write_only); - Handle.parallel_for(NumItems, [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); + sycl::accessor A(CheckBuffer, Handle, sycl::write_only); + Handle.parallel_for(NumItems, + [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); }); } @@ -46,20 +47,21 @@ int main() { ++Failed; } } - + Q.parallel_for(NumItems, [=](sycl::id<1> Idx) { - DataPtr[Idx] = Idx; - }).wait_and_throw(); - + DataPtr[Idx] = Idx; + }).wait_and_throw(); + syclext::set_access_mode(DataPtr,UsedGranularityInBytes, 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(NumItems, [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); + sycl::accessor A(ResultBuffer, Handle, sycl::write_only); + Handle.parallel_for(NumItems, + [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); }); } @@ -76,4 +78,4 @@ int main() { syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); return Failed; -} \ No newline at end of file +} From b1b23cdd22354511e41d868b7dc92f00906f5cb8 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 16:17:45 +0200 Subject: [PATCH 04/14] formatting fixed --- .../basic_access_from_kernel_virtual_mem.cpp | 114 +++++++++--------- 1 file changed, 60 insertions(+), 54 deletions(-) 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 index 114b3d74a7dfc..57b32b0cea584 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -11,51 +11,57 @@ namespace syclext = sycl::ext::oneapi::experimental; int main() { - sycl::queue Q; - sycl::context Context = Q.get_context(); - int Failed = 0; - - size_t UsedGranularityInBytes = syclext::get_mem_granularity(Context, syclext::granularity_mode::recommended); - - syclext::physical_mem NewPhysicalMem{Q.get_device(), Context, UsedGranularityInBytes}; - uintptr_t VirtualMemoryPtr = syclext::reserve_virtual_mem(0, UsedGranularityInBytes, Context); - - void *MappedPtr = NewPhysicalMem.map(VirtualMemoryPtr, UsedGranularityInBytes, syclext::address_access_mode::read_write); - - int* DataPtr = reinterpret_cast(MappedPtr); - - sycl::range NumItems{UsedGranularityInBytes/sizeof(int)}; - - std::vector ResultHostData(NumItems.size()); - - constexpr int ExpectedValueAfterFill = 1; - - Q.fill(DataPtr,ExpectedValueAfterFill , NumItems.size()).wait_and_throw(); - { - sycl::buffer CheckBuffer(ResultHostData); - Q.submit([&](sycl::handler &Handle) { - sycl::accessor A(CheckBuffer, Handle, sycl::write_only); - Handle.parallel_for(NumItems, - [=](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] - << " != " <(MappedPtr); + + sycl::range NumItems{UsedGranularityInBytes / sizeof(int)}; + + std::vector ResultHostData(NumItems.size()); + + constexpr int ExpectedValueAfterFill = 1; + + Q.fill(DataPtr, ExpectedValueAfterFill, NumItems.size()).wait_and_throw(); + { + sycl::buffer CheckBuffer(ResultHostData); + Q.submit([&](sycl::handler &Handle) { + sycl::accessor A(CheckBuffer, Handle, sycl::write_only); + Handle.parallel_for(NumItems, + [=](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(NumItems, [=](sycl::id<1> Idx) { - DataPtr[Idx] = Idx; - }).wait_and_throw(); + Q.parallel_for(NumItems, [=](sycl::id<1> Idx) { + DataPtr[Idx] = Idx; + }).wait_and_throw(); - syclext::set_access_mode(DataPtr,UsedGranularityInBytes, syclext::address_access_mode::read, Context); + syclext::set_access_mode(DataPtr, UsedGranularityInBytes, + syclext::address_access_mode::read, Context); - - { + { sycl::buffer ResultBuffer(ResultHostData); Q.submit([&](sycl::handler &Handle) { @@ -63,19 +69,19 @@ int main() { Handle.parallel_for(NumItems, [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); }); + } + + for (size_t i = 0; i < NumItems.size(); i++) { + const int ExpectedValue = static_cast(i); + if (ResultHostData[i] != ExpectedValue) { + std::cout << "Comparison failed at index " << i << ": " + << ResultHostData[i] << " != " << ExpectedValue << std::endl; + ++Failed; } - - for (size_t i = 0; i < NumItems.size(); 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, UsedGranularityInBytes, Context); - syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); + } + + syclext::unmap(MappedPtr, UsedGranularityInBytes, Context); + syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); - return Failed; + return Failed; } From 34d34e77a338f92d8a3b33b7991342e90d33611c Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Tue, 22 Oct 2024 16:36:15 +0200 Subject: [PATCH 05/14] Update sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp Co-authored-by: Alexey Sachkov --- .../VirtualMem/basic_access_from_kernel_virtual_mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 index 57b32b0cea584..e2eeab119589a 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -3,7 +3,7 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include "sycl/handler.hpp" +#include #include #include #include From a9bf4c3abad63a33907e6bec2891ad95913a85e0 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 17:12:55 +0200 Subject: [PATCH 06/14] removed unnecessary includes and added comment explaining test purpose --- .../VirtualMem/basic_access_from_kernel_virtual_mem.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) 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 index e2eeab119589a..15eaffb98fa6a 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -1,10 +1,11 @@ // REQUIRES: aspect-ext_oneapi_virtual_mem +// 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 -#include #include #include From 292bbe6450ff5decf96b9832a5ce1f6536f025f3 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 17:13:38 +0200 Subject: [PATCH 07/14] formatting fixed --- .../VirtualMem/basic_access_from_kernel_virtual_mem.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) 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 index 15eaffb98fa6a..85005356eb796 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -1,6 +1,7 @@ // REQUIRES: aspect-ext_oneapi_virtual_mem -// This test checks whether data can be correctly written to and read from virtual memory. +// This test checks whether data can be correctly written to and read from +// virtual memory. // RUN: %{build} -o %t.out // RUN: %{run} %t.out From f2a7e071bb65e80b5d57916ca63e459f6bcbe845 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 17:43:04 +0200 Subject: [PATCH 08/14] added bytes required and number of elements --- .../basic_access_from_kernel_virtual_mem.cpp | 34 ++++++++++--------- 1 file changed, 18 insertions(+), 16 deletions(-) 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 index 85005356eb796..6b87d90886e93 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -16,33 +16,35 @@ int main() { sycl::queue Q; sycl::context Context = Q.get_context(); int Failed = 0; - - size_t UsedGranularityInBytes = syclext::get_mem_granularity( + constexpr size_t NumberOfElements = 1000; + size_t BytesRequired = NumberOfElements*sizeof(int); + + size_t CtxGranularity = syclext::get_mem_granularity( Context, syclext::granularity_mode::recommended); + size_t AlignedByteSize = ((BytesRequired + CtxGranularity - 1) / CtxGranularity) * CtxGranularity; + syclext::physical_mem NewPhysicalMem{Q.get_device(), Context, - UsedGranularityInBytes}; + AlignedByteSize}; uintptr_t VirtualMemoryPtr = - syclext::reserve_virtual_mem(0, UsedGranularityInBytes, Context); + syclext::reserve_virtual_mem(0, AlignedByteSize, Context); void *MappedPtr = - NewPhysicalMem.map(VirtualMemoryPtr, UsedGranularityInBytes, + NewPhysicalMem.map(VirtualMemoryPtr, AlignedByteSize, syclext::address_access_mode::read_write); int *DataPtr = reinterpret_cast(MappedPtr); - sycl::range NumItems{UsedGranularityInBytes / sizeof(int)}; - - std::vector ResultHostData(NumItems.size()); + std::vector ResultHostData(NumberOfElements); constexpr int ExpectedValueAfterFill = 1; - Q.fill(DataPtr, ExpectedValueAfterFill, NumItems.size()).wait_and_throw(); + 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(NumItems, + Handle.parallel_for(NumberOfElements, [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); }); } @@ -56,11 +58,11 @@ int main() { } } - Q.parallel_for(NumItems, [=](sycl::id<1> Idx) { + Q.parallel_for(NumberOfElements, [=](sycl::id<1> Idx) { DataPtr[Idx] = Idx; }).wait_and_throw(); - syclext::set_access_mode(DataPtr, UsedGranularityInBytes, + syclext::set_access_mode(DataPtr, AlignedByteSize, syclext::address_access_mode::read, Context); { @@ -68,12 +70,12 @@ int main() { Q.submit([&](sycl::handler &Handle) { sycl::accessor A(ResultBuffer, Handle, sycl::write_only); - Handle.parallel_for(NumItems, + Handle.parallel_for(NumberOfElements, [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); }); } - for (size_t i = 0; i < NumItems.size(); i++) { + 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 << ": " @@ -82,8 +84,8 @@ int main() { } } - syclext::unmap(MappedPtr, UsedGranularityInBytes, Context); - syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); + syclext::unmap(MappedPtr, AlignedByteSize, Context); + syclext::free_virtual_mem(VirtualMemoryPtr, AlignedByteSize, Context); return Failed; } From ff092240ea26e1bd56e8522a3aeb99eee6a5bfe3 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 22 Oct 2024 17:43:27 +0200 Subject: [PATCH 09/14] fix --- .../VirtualMem/basic_access_from_kernel_virtual_mem.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) 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 index 6b87d90886e93..7230b11365819 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -17,12 +17,13 @@ int main() { sycl::context Context = Q.get_context(); int Failed = 0; constexpr size_t NumberOfElements = 1000; - size_t BytesRequired = NumberOfElements*sizeof(int); - + size_t BytesRequired = NumberOfElements * sizeof(int); + size_t CtxGranularity = syclext::get_mem_granularity( Context, syclext::granularity_mode::recommended); - size_t AlignedByteSize = ((BytesRequired + CtxGranularity - 1) / CtxGranularity) * CtxGranularity; + size_t AlignedByteSize = + ((BytesRequired + CtxGranularity - 1) / CtxGranularity) * CtxGranularity; syclext::physical_mem NewPhysicalMem{Q.get_device(), Context, AlignedByteSize}; From 1fd26cef6ed495def84a4539bb9864edebaa4000 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 23 Oct 2024 12:15:31 +0200 Subject: [PATCH 10/14] fix and added some improvements for future tests --- .../basic_access_from_kernel_virtual_mem.cpp | 14 ++++------ sycl/test-e2e/VirtualMem/helpers.hpp | 23 ++++++++++++++++ .../VirtualMem/vector_with_virtual_mem.cpp | 27 +++---------------- 3 files changed, 32 insertions(+), 32 deletions(-) create mode 100644 sycl/test-e2e/VirtualMem/helpers.hpp 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 index 7230b11365819..a4e4a176ae7b8 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -6,26 +6,22 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include -#include -#include - -namespace syclext = sycl::ext::oneapi::experimental; +#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 CtxGranularity = syclext::get_mem_granularity( - Context, syclext::granularity_mode::recommended); + size_t UsedGranularity = GetLCMGranularity(Device,Context); size_t AlignedByteSize = - ((BytesRequired + CtxGranularity - 1) / CtxGranularity) * CtxGranularity; + ((BytesRequired + UsedGranularity - 1) / UsedGranularity) * UsedGranularity; - syclext::physical_mem NewPhysicalMem{Q.get_device(), Context, + syclext::physical_mem NewPhysicalMem{Device, Context, AlignedByteSize}; uintptr_t VirtualMemoryPtr = syclext::reserve_virtual_mem(0, AlignedByteSize, Context); diff --git a/sycl/test-e2e/VirtualMem/helpers.hpp b/sycl/test-e2e/VirtualMem/helpers.hpp new file mode 100644 index 0000000000000..534bbffb7e317 --- /dev/null +++ b/sycl/test-e2e/VirtualMem/helpers.hpp @@ -0,0 +1,23 @@ +#pragma once + +#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, 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; +} diff --git a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp index cbbcf52e3ab25..36bedd2064a75 100644 --- a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp @@ -1,32 +1,13 @@ -// REQUIRES: aspect-ext_oneapi_virtual_mem, usm_shared_allocations +// REQUIRES: aspect-ext_oneapi_virtual_mem, aspect-usm_shared_allocations // 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(MContext); - size_t DevGranularity = syclext::get_mem_granularity(MDevice, MContext); - - size_t GCD = CtxGranularity; - size_t Rem = DevGranularity % GCD; - while (Rem != 0) { - std::swap(GCD, Rem); - Rem %= GCD; - } - return (DevGranularity / GCD) * LCMGranularity; -} + +#include "helpers.hpp" template class VirtualVector { public: @@ -233,4 +214,4 @@ int main() { } return 0; -} +} \ No newline at end of file From a09c57fc215992d29097b07834e06ca6b63e0160 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Wed, 23 Oct 2024 12:17:54 +0200 Subject: [PATCH 11/14] formatting fixed --- .../VirtualMem/basic_access_from_kernel_virtual_mem.cpp | 8 ++++---- sycl/test-e2e/VirtualMem/helpers.hpp | 4 +++- 2 files changed, 7 insertions(+), 5 deletions(-) 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 index a4e4a176ae7b8..b56ee90b69776 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -16,13 +16,13 @@ int main() { constexpr size_t NumberOfElements = 1000; size_t BytesRequired = NumberOfElements * sizeof(int); - size_t UsedGranularity = GetLCMGranularity(Device,Context); + size_t UsedGranularity = GetLCMGranularity(Device, Context); size_t AlignedByteSize = - ((BytesRequired + UsedGranularity - 1) / UsedGranularity) * UsedGranularity; + ((BytesRequired + UsedGranularity - 1) / UsedGranularity) * + UsedGranularity; - syclext::physical_mem NewPhysicalMem{Device, Context, - AlignedByteSize}; + syclext::physical_mem NewPhysicalMem{Device, Context, AlignedByteSize}; uintptr_t VirtualMemoryPtr = syclext::reserve_virtual_mem(0, AlignedByteSize, Context); diff --git a/sycl/test-e2e/VirtualMem/helpers.hpp b/sycl/test-e2e/VirtualMem/helpers.hpp index 534bbffb7e317..d655723a180fd 100644 --- a/sycl/test-e2e/VirtualMem/helpers.hpp +++ b/sycl/test-e2e/VirtualMem/helpers.hpp @@ -9,7 +9,9 @@ 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 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); From b143c437ea3a963c72d1237a5477c5d9e322f40c Mon Sep 17 00:00:00 2001 From: Daniel Skrobot Date: Mon, 28 Oct 2024 11:24:30 +0100 Subject: [PATCH 12/14] Update sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp Co-authored-by: Steffen Larsen --- sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp index dc7c189440420..f1697e98cd617 100644 --- a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp @@ -218,4 +218,4 @@ int main() { } return 0; -} \ No newline at end of file +} From 7ecefbc6bc9385d019f0258f2acd482d4cb5bc3c Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Mon, 28 Oct 2024 11:57:39 +0100 Subject: [PATCH 13/14] formatting fixed --- sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp index f1697e98cd617..946a2f7ba0092 100644 --- a/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/vector_with_virtual_mem.cpp @@ -12,7 +12,6 @@ #include "helpers.hpp" - template class VirtualVector { public: VirtualVector(sycl::queue &Q) From 8b47a4b5b93ffc8975639f54bdc74db11980780c Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Thu, 31 Oct 2024 01:39:08 -0700 Subject: [PATCH 14/14] common fix because helpers header --- .../VirtualMem/basic_access_from_kernel_virtual_mem.cpp | 2 -- sycl/test-e2e/VirtualMem/remapping_virtual_memory_range.cpp | 2 -- 2 files changed, 4 deletions(-) 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 index b56ee90b69776..4ad528b40d965 100644 --- a/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp +++ b/sycl/test-e2e/VirtualMem/basic_access_from_kernel_virtual_mem.cpp @@ -1,5 +1,3 @@ -// REQUIRES: aspect-ext_oneapi_virtual_mem - // This test checks whether data can be correctly written to and read from // virtual memory. 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"