From 03db009a2c3a916734dddee6679bcaae0b454343 Mon Sep 17 00:00:00 2001
From: Romanov Vlad <17316488+romanovvlad@users.noreply.github.com>
Date: Tue, 21 Jul 2020 19:25:16 +0300
Subject: [PATCH] [SYCL] Change read only accessor return const reference
 (#2142)

This should improve performance in some cases.
Also fixed memcpy-in-vec-as.cpp as the patch reveals problems
in the test.
---
 sycl/include/CL/sycl/accessor.hpp                    | 10 ++++++----
 .../test/basic_tests/accessor/addrspace_exposure.cpp |  4 ++--
 sycl/test/regression/memcpy-in-vec-as.cpp            | 12 ++++++++----
 3 files changed, 16 insertions(+), 10 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 7b26a13f475e7..d32a67ee2589b 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -260,6 +260,7 @@ class accessor_common {
       AccessMode == access::mode::read_write;
 
   using RefType = detail::const_if_const_AS<AS, DataT> &;
+  using ConstRefType = const DataT &;
   using PtrType = detail::const_if_const_AS<AS, DataT> *;
 
   using AccType =
@@ -299,7 +300,7 @@ class accessor_common {
 
     template <int CurDims = SubDims,
               typename = detail::enable_if_t<CurDims == 1 && IsAccessReadOnly>>
-    DataT operator[](size_t Index) const {
+    ConstRefType operator[](size_t Index) const {
       MIDs[Dims - SubDims] = Index;
       return MAccessor[MIDs];
     }
@@ -767,6 +768,7 @@ class accessor :
   using ConcreteASPtrType = typename detail::PtrValueType<DataT, AS>::type *;
 
   using RefType = detail::const_if_const_AS<AS, DataT> &;
+  using ConstRefType = const DataT &;
   using PtrType = detail::const_if_const_AS<AS, DataT> *;
 
   template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
@@ -1199,9 +1201,9 @@ class accessor :
     return *(getQualifiedPtr() + LinearIndex);
   }
 
-  template <int Dims = Dimensions,
-            typename = detail::enable_if_t<(Dims > 0) && IsAccessReadOnly>>
-  DataT operator[](id<Dimensions> Index) const {
+  template <int Dims = Dimensions>
+  typename detail::enable_if_t<(Dims > 0) && IsAccessReadOnly, ConstRefType>
+  operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return getQualifiedPtr()[LinearIndex];
   }
diff --git a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp
index ce73bf0296c10..ebc2540ee44d9 100644
--- a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp
+++ b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp
@@ -33,9 +33,9 @@ int main() {
     Cgh.single_task<class test>([=]() {
       static_assert(std::is_same<decltype(GlobalRWAcc[0]), int &>::value,
                     "Incorrect type from global read-write accessor");
-      static_assert(std::is_same<decltype(GlobalRAcc[0]), int>::value,
+      static_assert(std::is_same<decltype(GlobalRAcc[0]), const int &>::value,
                     "Incorrect type from global read accessor");
-      static_assert(std::is_same<decltype(ConstantAcc[0]), int>::value,
+      static_assert(std::is_same<decltype(ConstantAcc[0]), const int &>::value,
                     "Incorrect type from constant accessor");
       static_assert(std::is_same<decltype(LocalAcc[0]), int &>::value,
                     "Incorrect type from local accessor");
diff --git a/sycl/test/regression/memcpy-in-vec-as.cpp b/sycl/test/regression/memcpy-in-vec-as.cpp
index d53921234ffc1..b63b8643f3822 100644
--- a/sycl/test/regression/memcpy-in-vec-as.cpp
+++ b/sycl/test/regression/memcpy-in-vec-as.cpp
@@ -13,13 +13,17 @@ int main() {
                                                  cl::sycl::range<1>(1));
     cl::sycl::queue Queue;
     Queue.submit([&](cl::sycl::handler &cgh) {
-      auto In = InBuf.get_access<cl::sycl::access::mode::write>(cgh);
-      auto Out = OutBuf.get_access<cl::sycl::access::mode::read>(cgh);
+      auto In = InBuf.get_access<cl::sycl::access::mode::read>(cgh);
+      auto Out = OutBuf.get_access<cl::sycl::access::mode::write>(cgh);
       cgh.single_task<class as_op>(
           [=]() { Out[0] = In[0].as<res_vec_type>(); });
     });
   }
-  std::cout << res.s0() << " " << res.s1() << " " << res.s2() << " " << res.s3()
-            << std::endl;
+
+  if (res.s0() != 513 || res.s1() != 1027 || res.s2() != 1541 || res.s3() != 2055) {
+    std::cerr << "Incorrect result" << std::endl;
+    return 1;
+  }
+
   return 0;
 }