|
11 | 11 | namespace syclext = sycl::ext::oneapi::experimental; |
12 | 12 |
|
13 | 13 | int main() { |
14 | | - sycl::queue Q; |
15 | | - sycl::context Context = Q.get_context(); |
16 | | - int Failed = 0; |
17 | | - |
18 | | - size_t UsedGranularityInBytes = syclext::get_mem_granularity(Context, syclext::granularity_mode::recommended); |
19 | | - |
20 | | - syclext::physical_mem NewPhysicalMem{Q.get_device(), Context, UsedGranularityInBytes}; |
21 | | - uintptr_t VirtualMemoryPtr = syclext::reserve_virtual_mem(0, UsedGranularityInBytes, Context); |
22 | | - |
23 | | - void *MappedPtr = NewPhysicalMem.map(VirtualMemoryPtr, UsedGranularityInBytes, syclext::address_access_mode::read_write); |
24 | | - |
25 | | - int* DataPtr = reinterpret_cast<int*>(MappedPtr); |
26 | | - |
27 | | - sycl::range NumItems{UsedGranularityInBytes/sizeof(int)}; |
28 | | - |
29 | | - std::vector<int> ResultHostData(NumItems.size()); |
30 | | - |
31 | | - constexpr int ExpectedValueAfterFill = 1; |
32 | | - |
33 | | - Q.fill(DataPtr,ExpectedValueAfterFill , NumItems.size()).wait_and_throw(); |
34 | | - { |
35 | | - sycl::buffer<int> CheckBuffer(ResultHostData); |
36 | | - Q.submit([&](sycl::handler &Handle) { |
37 | | - sycl::accessor A(CheckBuffer, Handle, sycl::write_only); |
38 | | - Handle.parallel_for(NumItems, |
39 | | - [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); |
40 | | - }); |
41 | | - } |
42 | | - |
43 | | - for (size_t i = 0; i < ResultHostData.size(); i++) { |
44 | | - if (ResultHostData[i] != ExpectedValueAfterFill) { |
45 | | - std::cout << "Comparison failed after fill operation at index " << i << ": " << ResultHostData[i] |
46 | | - << " != " <<ExpectedValueAfterFill << std::endl; |
47 | | - ++Failed; |
48 | | - } |
| 14 | + sycl::queue Q; |
| 15 | + sycl::context Context = Q.get_context(); |
| 16 | + int Failed = 0; |
| 17 | + |
| 18 | + size_t UsedGranularityInBytes = syclext::get_mem_granularity( |
| 19 | + Context, syclext::granularity_mode::recommended); |
| 20 | + |
| 21 | + syclext::physical_mem NewPhysicalMem{Q.get_device(), Context, |
| 22 | + UsedGranularityInBytes}; |
| 23 | + uintptr_t VirtualMemoryPtr = |
| 24 | + syclext::reserve_virtual_mem(0, UsedGranularityInBytes, Context); |
| 25 | + |
| 26 | + void *MappedPtr = |
| 27 | + NewPhysicalMem.map(VirtualMemoryPtr, UsedGranularityInBytes, |
| 28 | + syclext::address_access_mode::read_write); |
| 29 | + |
| 30 | + int *DataPtr = reinterpret_cast<int *>(MappedPtr); |
| 31 | + |
| 32 | + sycl::range NumItems{UsedGranularityInBytes / sizeof(int)}; |
| 33 | + |
| 34 | + std::vector<int> ResultHostData(NumItems.size()); |
| 35 | + |
| 36 | + constexpr int ExpectedValueAfterFill = 1; |
| 37 | + |
| 38 | + Q.fill(DataPtr, ExpectedValueAfterFill, NumItems.size()).wait_and_throw(); |
| 39 | + { |
| 40 | + sycl::buffer<int> CheckBuffer(ResultHostData); |
| 41 | + Q.submit([&](sycl::handler &Handle) { |
| 42 | + sycl::accessor A(CheckBuffer, Handle, sycl::write_only); |
| 43 | + Handle.parallel_for(NumItems, |
| 44 | + [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); |
| 45 | + }); |
| 46 | + } |
| 47 | + |
| 48 | + for (size_t i = 0; i < ResultHostData.size(); i++) { |
| 49 | + if (ResultHostData[i] != ExpectedValueAfterFill) { |
| 50 | + std::cout << "Comparison failed after fill operation at index " << i |
| 51 | + << ": " << ResultHostData[i] << " != " << ExpectedValueAfterFill |
| 52 | + << std::endl; |
| 53 | + ++Failed; |
49 | 54 | } |
| 55 | + } |
50 | 56 |
|
51 | | - Q.parallel_for(NumItems, [=](sycl::id<1> Idx) { |
52 | | - DataPtr[Idx] = Idx; |
53 | | - }).wait_and_throw(); |
| 57 | + Q.parallel_for(NumItems, [=](sycl::id<1> Idx) { |
| 58 | + DataPtr[Idx] = Idx; |
| 59 | + }).wait_and_throw(); |
54 | 60 |
|
55 | | - syclext::set_access_mode(DataPtr,UsedGranularityInBytes, syclext::address_access_mode::read, Context); |
| 61 | + syclext::set_access_mode(DataPtr, UsedGranularityInBytes, |
| 62 | + syclext::address_access_mode::read, Context); |
56 | 63 |
|
57 | | - |
58 | | - { |
| 64 | + { |
59 | 65 | sycl::buffer<int> ResultBuffer(ResultHostData); |
60 | 66 |
|
61 | 67 | Q.submit([&](sycl::handler &Handle) { |
62 | 68 | sycl::accessor A(ResultBuffer, Handle, sycl::write_only); |
63 | 69 | Handle.parallel_for(NumItems, |
64 | 70 | [=](sycl::id<1> Idx) { A[Idx] = DataPtr[Idx]; }); |
65 | 71 | }); |
| 72 | + } |
| 73 | + |
| 74 | + for (size_t i = 0; i < NumItems.size(); i++) { |
| 75 | + const int ExpectedValue = static_cast<int>(i); |
| 76 | + if (ResultHostData[i] != ExpectedValue) { |
| 77 | + std::cout << "Comparison failed at index " << i << ": " |
| 78 | + << ResultHostData[i] << " != " << ExpectedValue << std::endl; |
| 79 | + ++Failed; |
66 | 80 | } |
67 | | - |
68 | | - for (size_t i = 0; i < NumItems.size(); i++) { |
69 | | - const int ExpectedValue = static_cast<int>(i); |
70 | | - if (ResultHostData[i] != ExpectedValue) { |
71 | | - std::cout << "Comparison failed at index " << i << ": " << ResultHostData[i] |
72 | | - << " != " << ExpectedValue<< std::endl; |
73 | | - ++Failed; |
74 | | - } |
75 | | - } |
76 | | - |
77 | | - syclext::unmap(MappedPtr, UsedGranularityInBytes, Context); |
78 | | - syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); |
| 81 | + } |
| 82 | + |
| 83 | + syclext::unmap(MappedPtr, UsedGranularityInBytes, Context); |
| 84 | + syclext::free_virtual_mem(VirtualMemoryPtr, UsedGranularityInBytes, Context); |
79 | 85 |
|
80 | | - return Failed; |
| 86 | + return Failed; |
81 | 87 | } |
0 commit comments