From 5ee3eb6490e0927ce059ee899877455f3448d26b Mon Sep 17 00:00:00 2001
From: amochalo <anastasiya.mochalova@intel.com>
Date: Wed, 19 Feb 2020 14:56:10 +0300
Subject: [PATCH 1/6] [SYCL] Fix restriction and some function

-Replacing "static_assetr" with "enable_if" for correct work with pointers
-fixing some functions for correct "make sycl-check"

Signed-off-by: amochalo <anastasiya.mochalova@intel.com>
---
 sycl/include/CL/sycl/accessor.hpp           | 26 +++++++++++----------
 sycl/include/CL/sycl/buffer.hpp             |  5 ++--
 sycl/include/CL/sycl/detail/buffer_impl.hpp |  2 +-
 sycl/include/CL/sycl/handler.hpp            |  2 +-
 4 files changed, 19 insertions(+), 16 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 74bb361bb96f2..cda8d9a5ef99c 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -773,11 +773,14 @@ class accessor :
 #endif
   }
 
-  template <int Dims = Dimensions, typename AllocatorT>
-  accessor(buffer<DataT, 1, AllocatorT> &BufferRef,
-           detail::enable_if_t<Dims == 0 &&
-                               (!IsPlaceH && (IsGlobalBuf || IsConstantBuf)),
-                               handler> &CommandGroupHandler)
+  template <int Dims = Dimensions, typename AllocatorT,
+	   typename = typename std::enable_if<
+		   (Dims == 0) && 
+                    (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>
+		    			::type>
+  accessor(buffer<DataT,1,AllocatorT> &BufferRef,
+		  handler &CommandGroupHandler)
+   
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
   }
@@ -793,11 +796,10 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions, typename AllocatorT,
-            typename detail::enable_if_t<
+            typename = typename std::enable_if<
                 (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
-                               (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>
-                * = nullptr>
-  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef)
+                               (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>::type>
+	accessor(buffer<DataT, Dims, AllocatorT> &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
   }
@@ -816,7 +818,7 @@ class accessor :
   template <int Dims = Dimensions, typename AllocatorT,
             typename = detail::enable_if_t<
                 (Dims > 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>>
-  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
            handler &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.get_range()) {
@@ -836,7 +838,7 @@ class accessor :
             typename = detail::enable_if_t<
                 (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
                                (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
            range<Dimensions> AccessRange, id<Dimensions> AccessOffset = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AccessOffset, AccessRange, BufferRef.get_range()) {
@@ -856,7 +858,7 @@ class accessor :
   template <int Dims = Dimensions, typename AllocatorT,
             typename = detail::enable_if_t<
                 (Dims > 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>>
-  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
            handler &CommandGroupHandler, range<Dimensions> AccessRange,
            id<Dimensions> AccessOffset = {})
 #ifdef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index b894bfeda9961..c995906193a9e 100644
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -21,7 +21,8 @@ class queue;
 template <int dimensions> class range;
 
 template <typename T, int dimensions = 1,
-          typename AllocatorT = cl::sycl::buffer_allocator>
+          typename AllocatorT = cl::sycl::buffer_allocator,
+	  typename = typename std::enable_if<(dimensions >0)&&(dimensions <= 3)>::type											>
 class buffer {
 public:
   using value_type = T;
@@ -300,7 +301,7 @@ class buffer {
   shared_ptr_class<detail::buffer_impl> impl;
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
-  template <typename A, int dims, typename C> friend class buffer;
+  template <typename A, int dims, typename C,typename Enable> friend class buffer;
   template <typename DataT, int dims, access::mode mode,
             access::target target, access::placeholder isPlaceholder>
   friend class accessor;
diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp
index f9c06b96709eb..e0e2609dd7bd4 100644
--- a/sycl/include/CL/sycl/detail/buffer_impl.hpp
+++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp
@@ -30,7 +30,7 @@ namespace sycl {
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::target AccessTarget, access::placeholder IsPlaceholder>
 class accessor;
-template <typename T, int Dimensions, typename AllocatorT> class buffer;
+template <typename T, int Dimensions, typename AllocatorT, typename Enable> class buffer;
 class handler;
 
 using buffer_allocator = detail::sycl_memory_object_allocator;
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index dd464de3d851c..43fe082fefb3d 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -60,7 +60,7 @@ namespace sycl {
 // Forward declaration
 
 class handler;
-template <typename T, int Dimensions, typename AllocatorT> class buffer;
+template <typename T, int Dimensions, typename AllocatorT, typename Enable> class buffer;
 namespace detail {
 
 /// This class is the default KernelName template parameter type for kernel

From 693c49d0f786c2ef952d5ee44a13d7e19abe9603 Mon Sep 17 00:00:00 2001
From: amochalo <anastasiya.mochalova@intel.com>
Date: Thu, 20 Feb 2020 21:54:47 +0300
Subject: [PATCH 2/6] Apply comments

Signed-off-by: amochalo <anastasiya.mochalova@intel.com>
---
 sycl/include/CL/sycl/accessor.hpp | 18 +++++++++---------
 1 file changed, 9 insertions(+), 9 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index cda8d9a5ef99c..25cfa250be20c 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -774,10 +774,10 @@ class accessor :
   }
 
   template <int Dims = Dimensions, typename AllocatorT,
-	   typename = typename std::enable_if<
+	   typename = typename detail::enable_if_t<
 		   (Dims == 0) && 
                     (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>
-		    			::type>
+		    			>
   accessor(buffer<DataT,1,AllocatorT> &BufferRef,
 		  handler &CommandGroupHandler)
    
@@ -796,9 +796,9 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions, typename AllocatorT,
-            typename = typename std::enable_if<
+            typename = typename detail::enable_if_t<
                 (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
-                               (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>::type>
+                               (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
 	accessor(buffer<DataT, Dims, AllocatorT> &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
@@ -934,17 +934,17 @@ class accessor :
   }
 
   template <int Dims = Dimensions>
-  operator typename std::enable_if<Dims == 0 &&
+  operator typename detail::enable_if_t<Dims == 0 &&
                                        AccessMode == access::mode::atomic,
-                                   atomic<DataT, AS>>::type() const {
+                                   atomic<DataT, AS>>() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
   }
 
   template <int Dims = Dimensions>
-  typename std::enable_if<(Dims > 0) && AccessMode == access::mode::atomic,
-                          atomic<DataT, AS>>::type
+  typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
+                          atomic<DataT, AS>>
   operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return atomic<DataT, AS>(
@@ -953,7 +953,7 @@ class accessor :
 
   template <int Dims = Dimensions>
   typename detail::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
-                               atomic<DataT, AS>>::type
+                               atomic<DataT, AS>>
   operator[](size_t Index) const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
     return atomic<DataT, AS>(

From 80222d87791388845439f08792301ab351951230 Mon Sep 17 00:00:00 2001
From: amochalo <anastasiya.mochalova@intel.com>
Date: Wed, 26 Feb 2020 19:36:43 +0300
Subject: [PATCH 3/6] Apply comment

Signed-off-by: amochalo <anastasiya.mochalova@intel.com>
---
 sycl/include/CL/sycl/accessor.hpp | 9 +++------
 1 file changed, 3 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 25cfa250be20c..4c68c2c6be4a4 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -763,8 +763,7 @@ class accessor :
       : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
 #else
       : AccessorBaseHost(
-            /*Offset=*/{0, 0, 0},
-            detail::convertToArrayOfN<3, 1>(range<1>{1}),
+            /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode,
             detail::getSyclObjImpl(BufferRef).get(), AdjustedDim, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
@@ -780,14 +779,12 @@ class accessor :
 		    			>
   accessor(buffer<DataT,1,AllocatorT> &BufferRef,
 		  handler &CommandGroupHandler)
-   
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), range<1>{1}, BufferRef.get_range()) {
   }
 #else
       : AccessorBaseHost(
-            /*Offset=*/{0, 0, 0},
-            detail::convertToArrayOfN<3, 1>(range<1>{1}),
+            /*Offset=*/{0, 0, 0}, detail::convertToArrayOfN<3, 1>(range<1>{1}),
             detail::convertToArrayOfN<3, 1>(BufferRef.get_range()), AccessMode,
             detail::getSyclObjImpl(BufferRef).get(), Dimensions, sizeof(DataT),
             BufferRef.OffsetInBytes, BufferRef.IsSubBuffer) {
@@ -1276,4 +1273,4 @@ struct hash<cl::sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
   }
 };
 
-} // namespace std
+} // __SYCL_INLINE_NAMESPACE(cl)

From 9197f4a3fe3f9b4dc62c49f1f7d7ab55a25ebf81 Mon Sep 17 00:00:00 2001
From: amochalo <anastasiya.mochalova@intel.com>
Date: Thu, 27 Feb 2020 16:46:42 +0300
Subject: [PATCH 4/6] Apply comment

Signed-off-by: amochalo <anastasiya.mochalova@intel.com>
---
 sycl/include/CL/sycl/accessor.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 4c68c2c6be4a4..797ca76d04007 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -1273,4 +1273,4 @@ struct hash<cl::sycl::accessor<DataT, Dimensions, AccessMode, AccessTarget,
   }
 };
 
-} // __SYCL_INLINE_NAMESPACE(cl)
+} // namespace std

From d1d14afa7ddb93e6b6ddb23ea23d66543c61d588 Mon Sep 17 00:00:00 2001
From: amochalo <anastasiya.mochalova@intel.com>
Date: Mon, 2 Mar 2020 10:01:05 +0300
Subject: [PATCH 5/6] Clang-format changes

Signed-off-by: amochalo <anastasiya.mochalova@intel.com>
---
 sycl/include/CL/sycl/accessor.hpp           | 10 +++++-----
 sycl/include/CL/sycl/buffer.hpp             | 10 ++++++----
 sycl/include/CL/sycl/detail/buffer_impl.hpp |  3 ++-
 sycl/include/CL/sycl/handler.hpp            |  3 ++-
 4 files changed, 15 insertions(+), 11 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 797ca76d04007..487363a21e773 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -796,7 +796,7 @@ class accessor :
             typename = typename detail::enable_if_t<
                 (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
                                (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
-	accessor(buffer<DataT, Dims, AllocatorT> &BufferRef)
+  accessor(buffer<DataT, Dims, AllocatorT> &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
   }
@@ -931,9 +931,9 @@ class accessor :
   }
 
   template <int Dims = Dimensions>
-  operator typename detail::enable_if_t<Dims == 0 &&
-                                       AccessMode == access::mode::atomic,
-                                   atomic<DataT, AS>>() const {
+  operator typename detail::enable_if_t<
+      Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
+      const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
@@ -941,7 +941,7 @@ class accessor :
 
   template <int Dims = Dimensions>
   typename detail::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
-                          atomic<DataT, AS>>
+                               atomic<DataT, AS>>
   operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return atomic<DataT, AS>(
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index c995906193a9e..6fe59191e059b 100644
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -22,7 +22,8 @@ template <int dimensions> class range;
 
 template <typename T, int dimensions = 1,
           typename AllocatorT = cl::sycl::buffer_allocator,
-	  typename = typename std::enable_if<(dimensions >0)&&(dimensions <= 3)>::type											>
+          typename = typename std::enable_if<(dimensions > 0) &&
+                                             (dimensions <= 3)>::type>
 class buffer {
 public:
   using value_type = T;
@@ -301,9 +302,10 @@ class buffer {
   shared_ptr_class<detail::buffer_impl> impl;
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
-  template <typename A, int dims, typename C,typename Enable> friend class buffer;
-  template <typename DataT, int dims, access::mode mode,
-            access::target target, access::placeholder isPlaceholder>
+  template <typename A, int dims, typename C, typename Enable>
+  friend class buffer;
+  template <typename DataT, int dims, access::mode mode, access::target target,
+            access::placeholder isPlaceholder>
   friend class accessor;
   range<dimensions> Range;
   // Offset field specifies the origin of the sub buffer inside the parent
diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp
index e0e2609dd7bd4..34e44369e00e0 100644
--- a/sycl/include/CL/sycl/detail/buffer_impl.hpp
+++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp
@@ -30,7 +30,8 @@ namespace sycl {
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::target AccessTarget, access::placeholder IsPlaceholder>
 class accessor;
-template <typename T, int Dimensions, typename AllocatorT, typename Enable> class buffer;
+template <typename T, int Dimensions, typename AllocatorT, typename Enable>
+class buffer;
 class handler;
 
 using buffer_allocator = detail::sycl_memory_object_allocator;
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 43fe082fefb3d..fccc7cff68348 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -60,7 +60,8 @@ namespace sycl {
 // Forward declaration
 
 class handler;
-template <typename T, int Dimensions, typename AllocatorT, typename Enable> class buffer;
+template <typename T, int Dimensions, typename AllocatorT, typename Enable>
+class buffer;
 namespace detail {
 
 /// This class is the default KernelName template parameter type for kernel

From fc54ac77dbb4d830a67e4535851f03a276e3aed5 Mon Sep 17 00:00:00 2001
From: amochalo <anastasiya.mochalova@intel.com>
Date: Wed, 4 Mar 2020 12:18:38 +0300
Subject: [PATCH 6/6] Add some checks for safety

Signed-off-by: amochalo <anastasiya.mochalova@intel.com>
---
 sycl/include/CL/sycl/accessor.hpp | 24 ++++++++++++++----------
 1 file changed, 14 insertions(+), 10 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 487363a21e773..700c074340261 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -793,9 +793,10 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions, typename AllocatorT,
-            typename = typename detail::enable_if_t<
-                (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
-                               (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
+            typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) &&
+                                           ((!IsPlaceH && IsHostBuf) ||
+                                            (IsPlaceH &&
+                                             (IsGlobalBuf || IsConstantBuf)))>>
   accessor(buffer<DataT, Dims, AllocatorT> &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.get_range()) {
@@ -813,8 +814,9 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions, typename AllocatorT,
-            typename = detail::enable_if_t<
-                (Dims > 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>>
+            typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) &&
+                                           (!IsPlaceH &&
+                                            (IsGlobalBuf || IsConstantBuf))>>
   accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
            handler &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
@@ -832,9 +834,10 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions, typename AllocatorT,
-            typename = detail::enable_if_t<
-                (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
-                               (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
+            typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) &&
+                                           ((!IsPlaceH && IsHostBuf) ||
+                                            (IsPlaceH &&
+                                             (IsGlobalBuf || IsConstantBuf)))>>
   accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
            range<Dimensions> AccessRange, id<Dimensions> AccessOffset = {})
 #ifdef __SYCL_DEVICE_ONLY__
@@ -853,8 +856,9 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions, typename AllocatorT,
-            typename = detail::enable_if_t<
-                (Dims > 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>>
+            typename = detail::enable_if_t<(Dims > 0) && (Dims == Dimensions) &&
+                                           (!IsPlaceH &&
+                                            (IsGlobalBuf || IsConstantBuf))>>
   accessor(buffer<DataT, Dims, AllocatorT> &BufferRef,
            handler &CommandGroupHandler, range<Dimensions> AccessRange,
            id<Dimensions> AccessOffset = {})