diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp index 5ca49f0067..b9231519cb 100644 --- a/SYCL/Basic/accessor/accessor.cpp +++ b/SYCL/Basic/accessor/accessor.cpp @@ -61,6 +61,62 @@ template struct Wrapper2 { template struct Wrapper3 { Wrapper2 w2; }; +template +void testLocalAccItersImpl(sycl::handler &cgh, GlobAcc &globAcc, LocAcc &locAcc, + bool testConstIter) { + if (testConstIter) { + cgh.single_task([=]() { + size_t Idx = 0; + for (auto &It : locAcc) { + It = globAcc[Idx++]; + } + Idx = 0; + for (auto It = locAcc.cbegin(); It != locAcc.cend(); It++) + globAcc[Idx++] = *It * 2 + 1; + Idx = locAcc.size() - 1; + for (auto It = locAcc.crbegin(); It != locAcc.crend(); It++) + globAcc[Idx--] += *It; + }); + } else { + cgh.single_task([=]() { + size_t Idx = 0; + for (auto It = locAcc.begin(); It != locAcc.end(); It++) + *It = globAcc[Idx++] * 2; + for (auto &It : locAcc) { + It++; + } + for (auto It = locAcc.rbegin(); It != locAcc.rend(); It++) { + *It *= 2; + *It += 1; + } + Idx = 0; + for (auto &It : locAcc) { + globAcc[Idx++] = It; + } + }); + } +} + +void testLocalAccIters(std::vector &vec, bool testConstIter = false, + bool test2D = false) { + try { + sycl::queue queue; + sycl::buffer buf(vec.data(), vec.size()); + queue.submit([&](sycl::handler &cgh) { + auto globAcc = buf.get_access(cgh); + if (test2D) { + sycl::local_accessor locAcc(sycl::range<2>{2, 16}, cgh); + testLocalAccItersImpl(cgh, globAcc, locAcc, testConstIter); + } else { + sycl::local_accessor locAcc(32, cgh); + testLocalAccItersImpl(cgh, globAcc, locAcc, testConstIter); + } + }); + } catch (sycl::exception &e) { + std::cout << e.what() << std::endl; + } +} + int main() { // Host accessor. { @@ -787,5 +843,40 @@ int main() { } } + // Test iterator methods with 1D local_accessor + { + std::vector v(32); + for (int i = 0; i < v.size(); ++i) { + v[i] = i; + } + testLocalAccIters(v); + for (int i = 0; i < v.size(); ++i) + assert(v[i] == ((i * 2 + 1) * 2 + 1)); + + for (int i = 0; i < v.size(); ++i) { + v[i] = i; + } + testLocalAccIters(v, true); + for (int i = 0; i < v.size(); ++i) + assert(v[i] == ((i * 2 + 1) + i)); + } + // Test iterator methods with 2D local_accessor + { + std::vector v(32); + for (int i = 0; i < v.size(); ++i) { + v[i] = i; + } + testLocalAccIters(v, false, true); + for (int i = 0; i < v.size(); ++i) + assert(v[i] == ((i * 2 + 1) * 2 + 1)); + + for (int i = 0; i < v.size(); ++i) { + v[i] = i; + } + testLocalAccIters(v, true, true); + for (int i = 0; i < v.size(); ++i) + assert(v[i] == ((i * 2 + 1) + i)); + } + std::cout << "Test passed" << std::endl; }