[SYCL] Remove unwanted changes.

Signed-off-by: Sergey Kanaev <[email protected]>

Author: Sergey Kanaev

sycl/include/CL/sycl/detail/defines.hpp

@@ -40,22 +40,6 @@
 
 #if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) && __has_builtin(__builtin_assume)
 #define __SYCL_ASSUME_INT(x) __builtin_assume(x <= INT_MAX)
-#define __SYCL_ASSUME_ID_INT(dims, x)                                          \
-  do {                                                                         \
-    for (size_t dim = 0; dim < dims; ++dim) {                                  \
-      size_t id = x.get_id(dim);                                               \
-      __SYCL_ASSUME_INT(id);                                                   \
-    }                                                                          \
-  } while (0)
-#define __SYCL_ASSUME_ARRAY_INT(dims, x)                                       \
-  do {                                                                         \
-    for (size_t dim = 0; dim < dims; ++dim) {                                  \
-      size_t id = x[dim];                                                      \
-      __SYCL_ASSUME_INT(id);                                                   \
-    }                                                                          \
-  } while (0)
 #else
 #define __SYCL_ASSUME_INT(x)
-#define __SYCL_ASSUME_ID_INT(dims, x)
-#define __SYCL_ASSUME_ARRAY_INT(dims, x)
 #endif

sycl/include/CL/sycl/item.hpp

@@ -29,10 +29,7 @@ template <int dimensions = 1, bool with_offset = true> class item {
 public:
   item() = delete;
 
-  id<dimensions> ALWAYS_INLINE get_id() const {
-    __SYCL_ASSUME_ARRAY_INT(dimensions, MImpl.MIndex);
-    return MImpl.MIndex;
-  }
+  id<dimensions> ALWAYS_INLINE get_id() const { return MImpl.MIndex; }
 
   size_t ALWAYS_INLINE get_id(int dimension) const {
     size_t id = MImpl.MIndex[dimension];
@@ -46,10 +43,7 @@ template <int dimensions = 1, bool with_offset = true> class item {
     return id;
   }
 
-  range<dimensions> ALWAYS_INLINE get_range() const {
-    __SYCL_ASSUME_ARRAY_INT(dimensions, MImpl.MExtent);
-    return MImpl.MExtent;
-  }
+  range<dimensions> ALWAYS_INLINE get_range() const { return MImpl.MExtent; }
 
   size_t ALWAYS_INLINE get_range(int dimension) const {
     size_t id = MImpl.MExtent[dimension];
@@ -58,9 +52,7 @@ template <int dimensions = 1, bool with_offset = true> class item {
   }
 
   template <bool has_offset = with_offset>
-  detail::enable_if_t<has_offset, id<dimensions>>
-      ALWAYS_INLINE get_offset() const {
-    __SYCL_ASSUME_ARRAY_INT(dimensions, MImpl.MOffset);
+  detail::enable_if_t<has_offset, id<dimensions>> get_offset() const {
     return MImpl.MOffset;
   }
 

sycl/include/CL/sycl/nd_item.hpp

@@ -32,10 +32,7 @@ template <int dimensions = 1> class nd_item {
 public:
   nd_item() = delete;
 
-  id<dimensions> ALWAYS_INLINE get_global_id() const {
-    __SYCL_ASSUME_ID_INT(dimensions, globalItem);
-    return globalItem.get_id();
-  }
+  id<dimensions> get_global_id() const { return globalItem.get_id(); }
 
   size_t ALWAYS_INLINE get_global_id(int dimension) const {
     size_t id = globalItem.get_id(dimension);
@@ -49,10 +46,7 @@ template <int dimensions = 1> class nd_item {
     return id;
   }
 
-  id<dimensions> ALWAYS_INLINE get_local_id() const {
-    __SYCL_ASSUME_ID_INT(dimensions, localItem);
-    return localItem.get_id();
-  }
+  id<dimensions> get_local_id() const { return localItem.get_id(); }
 
   size_t ALWAYS_INLINE get_local_id(int dimension) const {
     size_t id = localItem.get_id(dimension);
@@ -66,10 +60,7 @@ template <int dimensions = 1> class nd_item {
     return id;
   }
 
-  group<dimensions> ALWAYS_INLINE get_group() const {
-    __SYCL_ASSUME_ARRAY_INT(dimensions, Group);
-    return Group;
-  }
+  group<dimensions> get_group() const { return Group; }
 
   intel::sub_group get_sub_group() const { return intel::sub_group(); }
 

sycl/test/basic_tests/id_queries_fit_int.cpp

@@ -15,13 +15,6 @@ int main() {
   // CHECK: call void @llvm.assume(i1 {{.*}})
   int LinearId = TestItem.get_linear_id();
 
-  id<1> IdD = TestItem.get_id();
-  // CHECK: call void @llvm.assume(i1 {{.*}})
-  range<1> RangeD = TestItem.get_range();
-  // CHECK: call void @llvm.assume(i1 {{.*}})
-  id<1> Offset = TestItem.get_offset();
-  // CHECK: call void @llvm.assume(i1 {{.*}})
-
   cl::sycl::nd_item<1> TestNDItem =
       detail::Builder::createNDItem<1>(detail::Builder::createItem<1, false>({4}, {2}),
                                        detail::Builder::createItem<1, false>({2}, {0}),
@@ -46,12 +39,6 @@ int main() {
   // CHECK: call void @llvm.assume(i1 {{.*}})
   int LocalRange = TestNDItem.get_local_range(0);
 
-  id<1> GlobalIdD = TestNDItem.get_global_id();
-  // CHECK: call void @llvm.assume(i1 {{.*}})
-  id<1> LocalIdD = TestNDItem.get_local_id();
-  // CHECK: call void @llvm.assume(i1 {{.*}})
-  group<1> GroupD = TestNDItem.get_group();
-
   return 0;
 }
 // CHECK: }