[SYCL] Remove the offset field from item<dimensions, false> class.

sycl/include/CL/sycl/detail/helpers.hpp

@@ -75,6 +75,12 @@ class Builder {
     return cl::sycl::item<dimensions, with_offset>(R, I);
   }
 
+  template <int dimensions, bool with_offset>
+  static void updateItemIndex(cl::sycl::item<dimensions, with_offset> &Item,
+                              const id<dimensions> &NextIndex) {
+    Item.MImpl.MIndex = NextIndex;
+  }
+
   template <int dimensions>
   static nd_item<dimensions>
   createNDItem(const cl::sycl::item<dimensions, true> &GL,

sycl/include/CL/sycl/detail/item_base.hpp

@@ -0,0 +1,76 @@
+//==---------- item_base.hpp --- SYCL iteration ItemBase ------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/id.hpp>
+#include <CL/sycl/range.hpp>
+
+namespace cl {
+namespace sycl {
+template <int dimensions> class id;
+template <int dimensions> class range;
+
+namespace detail {
+template <int Dims, bool WithOffset> struct ItemBase;
+
+template <int Dims> struct ItemBase<Dims, true> {
+
+  bool operator==(const ItemBase &Rhs) const {
+    return (Rhs.MIndex == MIndex) && (Rhs.MExtent == MExtent) &&
+           (Rhs.MOffset == MOffset);
+  }
+
+  bool operator!=(const ItemBase &Rhs) const { return !((*this) == Rhs); }
+
+  size_t get_linear_id() const {
+    if (1 == Dims) {
+      return MIndex[0] - MOffset[0];
+    }
+    if (2 == Dims) {
+      return (MIndex[0] - MOffset[0]) * MExtent[1] + (MIndex[1] - MOffset[1]);
+    }
+    return ((MIndex[0] - MOffset[0]) * MExtent[1] * MExtent[2]) +
+           ((MIndex[1] - MOffset[1]) * MExtent[2]) + (MIndex[2] - MOffset[2]);
+  }
+
+  range<Dims> MExtent;
+  id<Dims> MIndex;
+  id<Dims> MOffset;
+};
+
+template <int Dims> struct ItemBase<Dims, false> {
+
+  bool operator==(const ItemBase &Rhs) const {
+    return (Rhs.MIndex == MIndex) && (Rhs.MExtent == MExtent);
+  }
+
+  bool operator!=(const ItemBase &Rhs) const { return !((*this) == Rhs); }
+
+  operator ItemBase<Dims, true>() const {
+    return ItemBase<Dims, true>(MExtent, MIndex, id<Dims>{});
+  }
+
+  size_t get_linear_id() const {
+    if (1 == Dims) {
+      return MIndex[0];
+    }
+    if (2 == Dims) {
+      return MIndex[0] * MExtent[1] + MIndex[1];
+    }
+    return (MIndex[0] * MExtent[1] * MExtent[2]) + (MIndex[1] * MExtent[2]) +
+           MIndex[2];
+  }
+
+  range<Dims> MExtent;
+  id<Dims> MIndex;
+};
+
+} // namespace detail
+} // namespace sycl
+} // namespace cl

sycl/include/CL/sycl/h_item.hpp

@@ -8,8 +8,10 @@
 
 #pragma once
 
+#include <CL/sycl/detail/helpers.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/item.hpp>
+#include <CL/sycl/range.hpp>
 
 namespace cl {
 namespace sycl {
@@ -22,9 +24,9 @@ template <int dimensions> class h_item {
 public:
   h_item() = delete;
 
-  h_item(const h_item<dimensions> &hi) = default;
+  h_item(const h_item &hi) = default;
 
-  h_item<dimensions> &operator=(const h_item<dimensions> &hi) = default;
+  h_item &operator=(const h_item &hi) = default;
 
   /* -- public interface members -- */
   item<dimensions, false> get_global() const { return globalItem; }
@@ -93,29 +95,29 @@ template <int dimensions> class h_item {
     return get_physical_local().get_id(dimension);
   }
 
-  bool operator==(const h_item<dimensions> &rhs) const {
+  bool operator==(const h_item &rhs) const {
     return (rhs.localItem == localItem) && (rhs.globalItem == globalItem) &&
            (rhs.logicalLocalItem == logicalLocalItem);
   }
 
-  bool operator!=(const h_item<dimensions> &rhs) const {
-    return !((*this) == rhs);
-  }
+  bool operator!=(const h_item &rhs) const { return !((*this) == rhs); }
 
 protected:
   friend class detail::Builder;
   friend class group<dimensions>;
   h_item(const item<dimensions, false> &GL, const item<dimensions, false> &L,
          const range<dimensions> &flexLocalRange)
       : globalItem(GL), localItem(L),
-        logicalLocalItem(flexLocalRange, L.get_id()) {}
+        logicalLocalItem(detail::Builder::createItem<dimensions, false>(
+            flexLocalRange, L.get_id())) {}
 
   h_item(const item<dimensions, false> &GL, const item<dimensions, false> &L)
       : globalItem(GL), localItem(L),
-        logicalLocalItem(localItem.get_range(), localItem.get_id()) {}
+        logicalLocalItem(detail::Builder::createItem<dimensions, false>(
+            localItem.get_range(), localItem.get_id())) {}
 
   void setLogicalLocalID(const id<dimensions> &ID) {
-    logicalLocalItem.setID(ID);
+    detail::Builder::updateItemIndex(logicalLocalItem, ID);
   }
 
 private:

sycl/include/CL/sycl/item.hpp

@@ -8,10 +8,11 @@
 
 #pragma once
 
+#include <CL/sycl/detail/helpers.hpp>
+#include <CL/sycl/detail/item_base.hpp>
+#include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/range.hpp>
-#include <stdexcept>
-#include <type_traits>
 
 namespace cl {
 namespace sycl {
@@ -20,93 +21,66 @@ class Builder;
 }
 template <int dimensions> class id;
 template <int dimensions> class range;
-template <int dimensions> class h_item;
 
 template <int dimensions = 1, bool with_offset = true> class item {
 public:
   item() = delete;
 
-  id<dimensions> get_id() const { return index; }
+  id<dimensions> get_id() const { return MImpl.MIndex; }
 
-  size_t get_id(int dimension) const { return index[dimension]; }
+  size_t get_id(int dimension) const { return MImpl.MIndex[dimension]; }
 
-  size_t operator[](int dimension) const { return index[dimension]; }
+  size_t operator[](int dimension) const { return MImpl.MIndex[dimension]; }
 
-  range<dimensions> get_range() const { return extent; }
+  range<dimensions> get_range() const { return MImpl.MExtent; }
 
-  size_t get_range(int dimension) const { return extent.get(dimension); }
+  size_t get_range(int dimension) const { return MImpl.MExtent[dimension]; }
 
-  // only available if with_offset is true;
-  template <bool W = with_offset,
-            typename = typename std::enable_if<(W == true)>::type>
-  id<dimensions> get_offset() const {
-    return offset;
+  template <bool has_offset = with_offset>
+  detail::enable_if_t<has_offset, id<dimensions>> get_offset() const {
+    return MImpl.MOffset;
   }
 
-  template <bool W = with_offset>
-  operator typename std::enable_if<W == false, item<dimensions, true>>::type()
-      const {
-    return item<dimensions, true>(extent, index, offset);
+  template <bool has_offset = with_offset>
+  detail::enable_if_t<has_offset, size_t> get_offset(int dimension) const {
+    return MImpl.MOffset[dimension];
   }
 
-  /* The following member function is only available in the id class
-   * specialization where: dimensions>0 and dimensions<4 */
-  template <int N = dimensions,
-            typename = typename std::enable_if<((N > 0) && (N < 4))>::type>
-  size_t get_linear_id() const {
-    if (1 == dimensions) {
-      return index[0] - offset[0];
-    }
-    if (2 == dimensions) {
-      return (index[0] - offset[0]) * extent[1] + (index[1] - offset[1]);
-    }
-    return ((index[0] - offset[0]) * extent[1] * extent[2]) +
-           ((index[1] - offset[1]) * extent[2]) + (index[2] - offset[2]);
+  template <bool has_offset = with_offset>
+  operator detail::enable_if_t<!has_offset, item<dimensions, true>>() const {
+    return detail::Builder::createItem<dimensions, true>(
+        MImpl.MExtent, MImpl.MIndex, /*Offset*/ {});
   }
 
-  item<dimensions, with_offset>(const item<dimensions, with_offset> &rhs) =
-      default;
+  size_t get_linear_id() const { return MImpl.get_linear_id(); }
 
-  item<dimensions, with_offset>(item<dimensions, with_offset> &&rhs) = default;
+  item(const item &rhs) = default;
 
-  item<dimensions, with_offset> &
-  operator=(const item<dimensions, with_offset> &rhs) = default;
+  item(item<dimensions, with_offset> &&rhs) = default;
 
-  item<dimensions, with_offset> &
-  operator=(item<dimensions, with_offset> &&rhs) = default;
+  item &operator=(const item &rhs) = default;
 
-  bool operator==(const item<dimensions, with_offset> &rhs) const {
-    return (rhs.index == this->index) && (rhs.extent == this->extent) &&
-           (rhs.offset == this->offset);
-  }
+  item &operator=(item &&rhs) = default;
 
-  bool operator!=(const item<dimensions, with_offset> &rhs) const {
-    return !((*this) == rhs);
-  }
+  bool operator==(const item &rhs) const { return rhs.MImpl == MImpl; }
 
-protected:
-  // For call constructor inside conversion operator
-  friend class item<dimensions, false>;
-  friend class item<dimensions, true>;
-  friend class h_item<dimensions>;
-  friend class detail::Builder;
+  bool operator!=(const item &rhs) const { return rhs.MImpl != MImpl; }
 
-  template <size_t W = with_offset>
-  item(typename std::enable_if<(W == true), const range<dimensions>>::type &R,
-       const id<dimensions> &I, const id<dimensions> &O)
-      : extent(R), index(I), offset(O) {}
+protected:
+  template <bool has_offset = with_offset>
+  item(detail::enable_if_t<has_offset, const range<dimensions>> &extent,
+       const id<dimensions> &index, const id<dimensions> &offset)
+      : MImpl{extent, index, offset} {}
 
-  template <size_t W = with_offset>
-  item(typename std::enable_if<(W == false), const range<dimensions>>::type &R,
-       const id<dimensions> &I)
-      : extent(R), index(I), offset() {}
+  template <bool has_offset = with_offset>
+  item(detail::enable_if_t<!has_offset, const range<dimensions>> &extent,
+       const id<dimensions> &index)
+      : MImpl{extent, index} {}
 
-  void setID(const id<dimensions> &ID) { index = ID; }
+  friend struct detail::Builder;
 
 private:
-  range<dimensions> extent;
-  id<dimensions> index;
-  id<dimensions> offset;
+  detail::ItemBase<dimensions, with_offset> MImpl;
 };
 
 } // namespace sycl