diff --git a/sycl/include/CL/sycl/detail/defines.hpp b/sycl/include/CL/sycl/detail/defines.hpp index 352b027890154..7c4381fae189e 100644 --- a/sycl/include/CL/sycl/detail/defines.hpp +++ b/sycl/include/CL/sycl/detail/defines.hpp @@ -8,6 +8,8 @@ #pragma once +#include + #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 diff --git a/sycl/include/CL/sycl/id.hpp b/sycl/include/CL/sycl/id.hpp index 71985b57ad7a1..b6882ca294bca 100644 --- a/sycl/include/CL/sycl/id.hpp +++ b/sycl/include/CL/sycl/id.hpp @@ -96,8 +96,10 @@ template class id : public detail::array { * 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__ diff --git a/sycl/include/CL/sycl/item.hpp b/sycl/include/CL/sycl/item.hpp index ed7a0d89a47ad..9e6d34acced18 100644 --- a/sycl/include/CL/sycl/item.hpp +++ b/sycl/include/CL/sycl/item.hpp @@ -8,12 +8,15 @@ #pragma once +#include #include #include #include #include #include +#include + __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace detail { @@ -28,13 +31,25 @@ template class item { id 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 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 detail::enable_if_t> get_offset() const { @@ -42,8 +57,11 @@ template class item { } template - detail::enable_if_t get_offset(int dimension) const { - return MImpl.MOffset[dimension]; + detail::enable_if_t + ALWAYS_INLINE get_offset(int dimension) const { + size_t Id = MImpl.MOffset[dimension]; + __SYCL_ASSUME_INT(Id); + return Id; } template @@ -52,7 +70,11 @@ template 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; diff --git a/sycl/include/CL/sycl/nd_item.hpp b/sycl/include/CL/sycl/nd_item.hpp index 101e394f5e943..4ac682ea6e2b9 100644 --- a/sycl/include/CL/sycl/nd_item.hpp +++ b/sycl/include/CL/sycl/nd_item.hpp @@ -19,6 +19,7 @@ #include #include +#include #include #include @@ -33,34 +34,57 @@ template class nd_item { id 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 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 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 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 get_global_range() const { return globalItem.get_range(); } @@ -101,39 +125,36 @@ template class nd_item { Group.mem_fence(); } - template + template device_event async_work_group_copy(local_ptr dest, global_ptr src, size_t numElements) const { return Group.async_work_group_copy(dest, src, numElements); } - template + template device_event async_work_group_copy(global_ptr dest, local_ptr src, size_t numElements) const { return Group.async_work_group_copy(dest, src, numElements); } - template + template device_event async_work_group_copy(local_ptr dest, - global_ptr src, - size_t numElements, + global_ptr src, size_t numElements, size_t srcStride) const { return Group.async_work_group_copy(dest, src, numElements, srcStride); } - template + template device_event async_work_group_copy(global_ptr dest, - local_ptr src, - size_t numElements, + local_ptr src, size_t numElements, size_t destStride) const { return Group.async_work_group_copy(dest, src, numElements, destStride); } - template - void wait_for(eventTN... events) const { + template void wait_for(eventTN... events) const { Group.wait_for(events...); } diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp new file mode 100644 index 0000000000000..8389b7e1de355 --- /dev/null +++ b/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 + +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: }