[SYCL] Add runtime support for fsycl-id-queries-fit-in-int

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

@@ -8,6 +8,8 @@
 
 #pragma once
 
+#include <climits>
+
 #ifndef __SYCL_DISABLE_NAMESPACE_INLINE__
 #define __SYCL_INLINE_NAMESPACE(X) inline namespace X
 #else
@@ -18,6 +20,10 @@
 #define __has_attribute(x) 0
 #endif
 
+#ifndef __has_builtin
+#define __has_builtin(x) 0
+#endif
+
 #if __has_attribute(always_inline)
 #define ALWAYS_INLINE __attribute__((always_inline))
 #else
@@ -31,3 +37,12 @@
 #ifndef SYCL_EXTERNAL
 #define SYCL_EXTERNAL
 #endif
+
+#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) && __has_builtin(__builtin_assume)
+#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= INT_MAX)
+#else
+#define __SYCL_ASSUME_INT(x)
+#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) && !__has_builtin(__builtin_assume)
+#warning "No assumptions will be emitted due to no __builtin_assume available"
+#endif
+#endif

sycl/include/CL/sycl/id.hpp

@@ -96,8 +96,10 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
    * conversion:
    * int a = id<1>(value); */
 
-  operator EnableIfT<(dimensions == 1), size_t>() const {
-    return this->common_array[0];
+  ALWAYS_INLINE operator EnableIfT<(dimensions == 1), size_t>() const {
+    size_t Result = this->common_array[0];
+    __SYCL_ASSUME_INT(Result);
+    return Result;
   }
 #endif // __SYCL_DISABLE_ID_TO_INT_CONV__
 

sycl/include/CL/sycl/item.hpp

@@ -8,12 +8,15 @@
 
 #pragma once
 
+#include <CL/sycl/detail/defines.hpp>
 #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 <cstddef>
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
@@ -28,22 +31,37 @@ template <int dimensions = 1, bool with_offset = true> class item {
 
   id<dimensions> get_id() const { return MImpl.MIndex; }
 
-  size_t get_id(int dimension) const { return MImpl.MIndex[dimension]; }
+  size_t ALWAYS_INLINE get_id(int dimension) const {
+    size_t Id = MImpl.MIndex[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
-  size_t operator[](int dimension) const { return MImpl.MIndex[dimension]; }
+  size_t ALWAYS_INLINE operator[](int dimension) const {
+    size_t Id = MImpl.MIndex[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   range<dimensions> get_range() const { return MImpl.MExtent; }
 
-  size_t get_range(int dimension) const { return MImpl.MExtent[dimension]; }
+  size_t ALWAYS_INLINE get_range(int dimension) const {
+    size_t Id = MImpl.MExtent[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   template <bool has_offset = with_offset>
   detail::enable_if_t<has_offset, id<dimensions>> get_offset() const {
     return MImpl.MOffset;
   }
 
   template <bool has_offset = with_offset>
-  detail::enable_if_t<has_offset, size_t> get_offset(int dimension) const {
-    return MImpl.MOffset[dimension];
+  detail::enable_if_t<has_offset, size_t>
+      ALWAYS_INLINE get_offset(int dimension) const {
+    size_t Id = MImpl.MOffset[dimension];
+    __SYCL_ASSUME_INT(Id);
+    return Id;
   }
 
   template <bool has_offset = with_offset>
@@ -52,7 +70,11 @@ template <int dimensions = 1, bool with_offset = true> class item {
         MImpl.MExtent, MImpl.MIndex, /*Offset*/ {});
   }
 
-  size_t get_linear_id() const { return MImpl.get_linear_id(); }
+  size_t ALWAYS_INLINE get_linear_id() const {
+    size_t Id = MImpl.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   item(const item &rhs) = default;
 

sycl/include/CL/sycl/nd_item.hpp

@@ -19,6 +19,7 @@
 #include <CL/sycl/nd_range.hpp>
 #include <CL/sycl/range.hpp>
 
+#include <cstddef>
 #include <stdexcept>
 #include <type_traits>
 
@@ -33,34 +34,57 @@ template <int dimensions = 1> class nd_item {
 
   id<dimensions> get_global_id() const { return globalItem.get_id(); }
 
-  size_t get_global_id(int dimension) const {
-    return globalItem.get_id(dimension);
+  size_t ALWAYS_INLINE get_global_id(int dimension) const {
+    size_t Id = globalItem.get_id(dimension);
+    __SYCL_ASSUME_INT(Id);
+    return Id;
   }
 
-  size_t get_global_linear_id() const { return globalItem.get_linear_id(); }
+  size_t ALWAYS_INLINE get_global_linear_id() const {
+    size_t Id = globalItem.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   id<dimensions> get_local_id() const { return localItem.get_id(); }
 
-  size_t get_local_id(int dimension) const {
-    return localItem.get_id(dimension);
+  size_t ALWAYS_INLINE get_local_id(int dimension) const {
+    size_t Id = localItem.get_id(dimension);
+    __SYCL_ASSUME_INT(Id);
+    return Id;
   }
 
-  size_t get_local_linear_id() const { return localItem.get_linear_id(); }
+  size_t get_local_linear_id() const {
+    size_t Id = localItem.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   group<dimensions> get_group() const { return Group; }
 
   intel::sub_group get_sub_group() const { return intel::sub_group(); }
 
-  size_t get_group(int dimension) const { return Group[dimension]; }
+  size_t ALWAYS_INLINE get_group(int dimension) const {
+    size_t Size = Group[dimension];
+    __SYCL_ASSUME_INT(Size);
+    return Size;
+  }
 
-  size_t get_group_linear_id() const { return Group.get_linear_id(); }
+  size_t ALWAYS_INLINE get_group_linear_id() const {
+    size_t Id = Group.get_linear_id();
+    __SYCL_ASSUME_INT(Id);
+    return Id;
+  }
 
   range<dimensions> get_group_range() const {
     return Group.get_global_range() / Group.get_local_range();
   }
 
-  size_t get_group_range(int dimension) const {
-    return Group.get_global_range(dimension) / Group.get_local_range(dimension);
+  size_t ALWAYS_INLINE get_group_range(int dimension) const {
+    size_t Range =
+        Group.get_global_range(dimension) / Group.get_local_range(dimension);
+    __SYCL_ASSUME_INT(Range);
+    return Range;
   }
 
   range<dimensions> get_global_range() const { return globalItem.get_range(); }
@@ -101,39 +125,36 @@ template <int dimensions = 1> class nd_item {
     Group.mem_fence();
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(local_ptr<dataT> dest,
                                      global_ptr<dataT> src,
                                      size_t numElements) const {
     return Group.async_work_group_copy(dest, src, numElements);
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(global_ptr<dataT> dest,
                                      local_ptr<dataT> src,
                                      size_t numElements) const {
     return Group.async_work_group_copy(dest, src, numElements);
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(local_ptr<dataT> dest,
-                                     global_ptr<dataT> src,
-                                     size_t numElements,
+                                     global_ptr<dataT> src, size_t numElements,
                                      size_t srcStride) const {
 
     return Group.async_work_group_copy(dest, src, numElements, srcStride);
   }
 
-  template<typename dataT>
+  template <typename dataT>
   device_event async_work_group_copy(global_ptr<dataT> dest,
-                                     local_ptr<dataT> src,
-                                     size_t numElements,
+                                     local_ptr<dataT> src, size_t numElements,
                                      size_t destStride) const {
     return Group.async_work_group_copy(dest, src, numElements, destStride);
   }
 
-  template<typename... eventTN>
-  void wait_for(eventTN... events) const {
+  template <typename... eventTN> void wait_for(eventTN... events) const {
     Group.wait_for(events...);
   }
 

sycl/test/check_device_code/id_queries_fit_int.cpp

@@ -0,0 +1,51 @@
+// RUN: %clangxx -fsycl -Xclang -fsycl-is-host -O1 -c -S -emit-llvm -o %t.ll -D__SYCL_ID_QUERIES_FIT_IN_INT__=1 %s
+// RUN: FileCheck %s --input-file %t.ll
+
+#include <CL/sycl.hpp>
+
+using namespace sycl;
+
+// CHECK: define dso_local i32 @main() {{.*}} {
+int main() {
+  item<1, true> TestItem = detail::Builder::createItem<1, true>({3}, {2}, {1});
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int Id = TestItem.get_id(0);
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int Range = TestItem.get_range(0);
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int LinearId = TestItem.get_linear_id();
+
+  cl::sycl::nd_item<1> TestNDItem =
+      detail::Builder::createNDItem<1>(detail::Builder::createItem<1, false>({4}, {2}),
+                                       detail::Builder::createItem<1, false>({2}, {0}),
+                                       detail::Builder::createGroup<1>({4}, {2}, {1}));
+
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int GlobalId = TestNDItem.get_global_id(0);
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int GlobalLinearId = TestNDItem.get_global_linear_id();
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int LocalId = TestNDItem.get_local_id(0);
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int LocalLinearId = TestNDItem.get_local_linear_id();
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int GroupRange = TestNDItem.get_group_range(0);
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int GroupId = TestNDItem.get_group(0);
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int GroupLinearId = TestNDItem.get_group_linear_id();
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int GlobalRange = TestNDItem.get_global_range(0);
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int LocalRange = TestNDItem.get_local_range(0);
+
+  int GlobalIdConverted = TestNDItem.get_global_id();
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int LocalIdConverted = TestNDItem.get_local_id();
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+  int OffsetConferted = TestNDItem.get_offset();
+  // CHECK: call void @llvm.assume(i1 {{.*}})
+
+  return 0;
+}
+// CHECK: }