Skip to content

Commit fdcaeae

Browse files
s-kanaevs-kanaev
authored
[SYCL] Add runtime support for fsycl-id-queries-fit-in-int (#1685)
Signed-off-by: Alexander Batashev <alexander.batashev@intel.com> Signed-off-by: Sergey Kanaev <sergey.kanaev@intel.com>
1 parent 9d4c284 commit fdcaeae

File tree

5 files changed

+139
-28
lines changed

5 files changed

+139
-28
lines changed

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

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,8 @@
88

99
#pragma once
1010

11+
#include <climits>
12+
1113
#ifndef __SYCL_DISABLE_NAMESPACE_INLINE__
1214
#define __SYCL_INLINE_NAMESPACE(X) inline namespace X
1315
#else
@@ -18,6 +20,10 @@
1820
#define __has_attribute(x) 0
1921
#endif
2022

23+
#ifndef __has_builtin
24+
#define __has_builtin(x) 0
25+
#endif
26+
2127
#if __has_attribute(always_inline)
2228
#define ALWAYS_INLINE __attribute__((always_inline))
2329
#else
@@ -31,3 +37,12 @@
3137
#ifndef SYCL_EXTERNAL
3238
#define SYCL_EXTERNAL
3339
#endif
40+
41+
#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) && __has_builtin(__builtin_assume)
42+
#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= INT_MAX)
43+
#else
44+
#define __SYCL_ASSUME_INT(x)
45+
#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__) && !__has_builtin(__builtin_assume)
46+
#warning "No assumptions will be emitted due to no __builtin_assume available"
47+
#endif
48+
#endif

‎sycl/include/CL/sycl/id.hpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -96,8 +96,10 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
9696
* conversion:
9797
* int a = id<1>(value); */
9898

99-
operator EnableIfT<(dimensions == 1), size_t>() const {
100-
return this->common_array[0];
99+
ALWAYS_INLINE operator EnableIfT<(dimensions == 1), size_t>() const {
100+
size_t Result = this->common_array[0];
101+
__SYCL_ASSUME_INT(Result);
102+
return Result;
101103
}
102104
#endif // __SYCL_DISABLE_ID_TO_INT_CONV__
103105

‎sycl/include/CL/sycl/item.hpp

Lines changed: 28 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -8,12 +8,15 @@
88

99
#pragma once
1010

11+
#include <CL/sycl/detail/defines.hpp>
1112
#include <CL/sycl/detail/helpers.hpp>
1213
#include <CL/sycl/detail/item_base.hpp>
1314
#include <CL/sycl/detail/type_traits.hpp>
1415
#include <CL/sycl/id.hpp>
1516
#include <CL/sycl/range.hpp>
1617

18+
#include <cstddef>
19+
1720
__SYCL_INLINE_NAMESPACE(cl) {
1821
namespace sycl {
1922
namespace detail {
@@ -28,22 +31,37 @@ template <int dimensions = 1, bool with_offset = true> class item {
2831

2932
id<dimensions> get_id() const { return MImpl.MIndex; }
3033

31-
size_t get_id(int dimension) const { return MImpl.MIndex[dimension]; }
34+
size_t ALWAYS_INLINE get_id(int dimension) const {
35+
size_t Id = MImpl.MIndex[dimension];
36+
__SYCL_ASSUME_INT(Id);
37+
return Id;
38+
}
3239

33-
size_t operator[](int dimension) const { return MImpl.MIndex[dimension]; }
40+
size_t ALWAYS_INLINE operator[](int dimension) const {
41+
size_t Id = MImpl.MIndex[dimension];
42+
__SYCL_ASSUME_INT(Id);
43+
return Id;
44+
}
3445

3546
range<dimensions> get_range() const { return MImpl.MExtent; }
3647

37-
size_t get_range(int dimension) const { return MImpl.MExtent[dimension]; }
48+
size_t ALWAYS_INLINE get_range(int dimension) const {
49+
size_t Id = MImpl.MExtent[dimension];
50+
__SYCL_ASSUME_INT(Id);
51+
return Id;
52+
}
3853

3954
template <bool has_offset = with_offset>
4055
detail::enable_if_t<has_offset, id<dimensions>> get_offset() const {
4156
return MImpl.MOffset;
4257
}
4358

4459
template <bool has_offset = with_offset>
45-
detail::enable_if_t<has_offset, size_t> get_offset(int dimension) const {
46-
return MImpl.MOffset[dimension];
60+
detail::enable_if_t<has_offset, size_t>
61+
ALWAYS_INLINE get_offset(int dimension) const {
62+
size_t Id = MImpl.MOffset[dimension];
63+
__SYCL_ASSUME_INT(Id);
64+
return Id;
4765
}
4866

4967
template <bool has_offset = with_offset>
@@ -52,7 +70,11 @@ template <int dimensions = 1, bool with_offset = true> class item {
5270
MImpl.MExtent, MImpl.MIndex, /*Offset*/ {});
5371
}
5472

55-
size_t get_linear_id() const { return MImpl.get_linear_id(); }
73+
size_t ALWAYS_INLINE get_linear_id() const {
74+
size_t Id = MImpl.get_linear_id();
75+
__SYCL_ASSUME_INT(Id);
76+
return Id;
77+
}
5678

5779
item(const item &rhs) = default;
5880

‎sycl/include/CL/sycl/nd_item.hpp

Lines changed: 41 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@
1919
#include <CL/sycl/nd_range.hpp>
2020
#include <CL/sycl/range.hpp>
2121

22+
#include <cstddef>
2223
#include <stdexcept>
2324
#include <type_traits>
2425

@@ -33,34 +34,57 @@ template <int dimensions = 1> class nd_item {
3334

3435
id<dimensions> get_global_id() const { return globalItem.get_id(); }
3536

36-
size_t get_global_id(int dimension) const {
37-
return globalItem.get_id(dimension);
37+
size_t ALWAYS_INLINE get_global_id(int dimension) const {
38+
size_t Id = globalItem.get_id(dimension);
39+
__SYCL_ASSUME_INT(Id);
40+
return Id;
3841
}
3942

40-
size_t get_global_linear_id() const { return globalItem.get_linear_id(); }
43+
size_t ALWAYS_INLINE get_global_linear_id() const {
44+
size_t Id = globalItem.get_linear_id();
45+
__SYCL_ASSUME_INT(Id);
46+
return Id;
47+
}
4148

4249
id<dimensions> get_local_id() const { return localItem.get_id(); }
4350

44-
size_t get_local_id(int dimension) const {
45-
return localItem.get_id(dimension);
51+
size_t ALWAYS_INLINE get_local_id(int dimension) const {
52+
size_t Id = localItem.get_id(dimension);
53+
__SYCL_ASSUME_INT(Id);
54+
return Id;
4655
}
4756

48-
size_t get_local_linear_id() const { return localItem.get_linear_id(); }
57+
size_t get_local_linear_id() const {
58+
size_t Id = localItem.get_linear_id();
59+
__SYCL_ASSUME_INT(Id);
60+
return Id;
61+
}
4962

5063
group<dimensions> get_group() const { return Group; }
5164

5265
intel::sub_group get_sub_group() const { return intel::sub_group(); }
5366

54-
size_t get_group(int dimension) const { return Group[dimension]; }
67+
size_t ALWAYS_INLINE get_group(int dimension) const {
68+
size_t Size = Group[dimension];
69+
__SYCL_ASSUME_INT(Size);
70+
return Size;
71+
}
5572

56-
size_t get_group_linear_id() const { return Group.get_linear_id(); }
73+
size_t ALWAYS_INLINE get_group_linear_id() const {
74+
size_t Id = Group.get_linear_id();
75+
__SYCL_ASSUME_INT(Id);
76+
return Id;
77+
}
5778

5879
range<dimensions> get_group_range() const {
5980
return Group.get_global_range() / Group.get_local_range();
6081
}
6182

62-
size_t get_group_range(int dimension) const {
63-
return Group.get_global_range(dimension) / Group.get_local_range(dimension);
83+
size_t ALWAYS_INLINE get_group_range(int dimension) const {
84+
size_t Range =
85+
Group.get_global_range(dimension) / Group.get_local_range(dimension);
86+
__SYCL_ASSUME_INT(Range);
87+
return Range;
6488
}
6589

6690
range<dimensions> get_global_range() const { return globalItem.get_range(); }
@@ -101,39 +125,36 @@ template <int dimensions = 1> class nd_item {
101125
Group.mem_fence();
102126
}
103127

104-
template<typename dataT>
128+
template <typename dataT>
105129
device_event async_work_group_copy(local_ptr<dataT> dest,
106130
global_ptr<dataT> src,
107131
size_t numElements) const {
108132
return Group.async_work_group_copy(dest, src, numElements);
109133
}
110134

111-
template<typename dataT>
135+
template <typename dataT>
112136
device_event async_work_group_copy(global_ptr<dataT> dest,
113137
local_ptr<dataT> src,
114138
size_t numElements) const {
115139
return Group.async_work_group_copy(dest, src, numElements);
116140
}
117141

118-
template<typename dataT>
142+
template <typename dataT>
119143
device_event async_work_group_copy(local_ptr<dataT> dest,
120-
global_ptr<dataT> src,
121-
size_t numElements,
144+
global_ptr<dataT> src, size_t numElements,
122145
size_t srcStride) const {
123146

124147
return Group.async_work_group_copy(dest, src, numElements, srcStride);
125148
}
126149

127-
template<typename dataT>
150+
template <typename dataT>
128151
device_event async_work_group_copy(global_ptr<dataT> dest,
129-
local_ptr<dataT> src,
130-
size_t numElements,
152+
local_ptr<dataT> src, size_t numElements,
131153
size_t destStride) const {
132154
return Group.async_work_group_copy(dest, src, numElements, destStride);
133155
}
134156

135-
template<typename... eventTN>
136-
void wait_for(eventTN... events) const {
157+
template <typename... eventTN> void wait_for(eventTN... events) const {
137158
Group.wait_for(events...);
138159
}
139160

‎sycl/test/check_device_code/id_queries_fit_int.cpp

Lines changed: 51 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,51 @@
1+
// RUN: %clangxx -fsycl -Xclang -fsycl-is-host -O1 -c -S -emit-llvm -o %t.ll -D__SYCL_ID_QUERIES_FIT_IN_INT__=1 %s
2+
// RUN: FileCheck %s --input-file %t.ll
3+
4+
#include <CL/sycl.hpp>
5+
6+
using namespace sycl;
7+
8+
// CHECK: define dso_local i32 @main() {{.*}} {
9+
int main() {
10+
item<1, true> TestItem = detail::Builder::createItem<1, true>({3}, {2}, {1});
11+
// CHECK: call void @llvm.assume(i1 {{.*}})
12+
int Id = TestItem.get_id(0);
13+
// CHECK: call void @llvm.assume(i1 {{.*}})
14+
int Range = TestItem.get_range(0);
15+
// CHECK: call void @llvm.assume(i1 {{.*}})
16+
int LinearId = TestItem.get_linear_id();
17+
18+
cl::sycl::nd_item<1> TestNDItem =
19+
detail::Builder::createNDItem<1>(detail::Builder::createItem<1, false>({4}, {2}),
20+
detail::Builder::createItem<1, false>({2}, {0}),
21+
detail::Builder::createGroup<1>({4}, {2}, {1}));
22+
23+
// CHECK: call void @llvm.assume(i1 {{.*}})
24+
int GlobalId = TestNDItem.get_global_id(0);
25+
// CHECK: call void @llvm.assume(i1 {{.*}})
26+
int GlobalLinearId = TestNDItem.get_global_linear_id();
27+
// CHECK: call void @llvm.assume(i1 {{.*}})
28+
int LocalId = TestNDItem.get_local_id(0);
29+
// CHECK: call void @llvm.assume(i1 {{.*}})
30+
int LocalLinearId = TestNDItem.get_local_linear_id();
31+
// CHECK: call void @llvm.assume(i1 {{.*}})
32+
int GroupRange = TestNDItem.get_group_range(0);
33+
// CHECK: call void @llvm.assume(i1 {{.*}})
34+
int GroupId = TestNDItem.get_group(0);
35+
// CHECK: call void @llvm.assume(i1 {{.*}})
36+
int GroupLinearId = TestNDItem.get_group_linear_id();
37+
// CHECK: call void @llvm.assume(i1 {{.*}})
38+
int GlobalRange = TestNDItem.get_global_range(0);
39+
// CHECK: call void @llvm.assume(i1 {{.*}})
40+
int LocalRange = TestNDItem.get_local_range(0);
41+
42+
int GlobalIdConverted = TestNDItem.get_global_id();
43+
// CHECK: call void @llvm.assume(i1 {{.*}})
44+
int LocalIdConverted = TestNDItem.get_local_id();
45+
// CHECK: call void @llvm.assume(i1 {{.*}})
46+
int OffsetConferted = TestNDItem.get_offset();
47+
// CHECK: call void @llvm.assume(i1 {{.*}})
48+
49+
return 0;
50+
}
51+
// CHECK: }

0 commit comments

Comments
 (0)