1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
|
/*========================== begin_copyright_notice ============================
Copyright (C) 2017-2021 Intel Corporation
SPDX-License-Identifier: MIT
============================= end_copyright_notice ===========================*/
//*****************************************************************************/
// Work-Item functions
//*****************************************************************************/
#define MAX_DIM 2
INLINE size_t OVERLOADABLE get_enqueued_local_size(uint dim) {
if (dim > MAX_DIM) {
return 1;
}
return __builtin_IB_get_enqueued_local_size(dim);
}
INLINE size_t OVERLOADABLE get_global_id(uint dim) {
if (dim > MAX_DIM) {
return 0;
}
return get_group_id(dim) * get_enqueued_local_size(dim) + get_local_id(dim) + get_global_offset(dim);
}
INLINE size_t OVERLOADABLE get_group_id(uint dim) {
if (dim > MAX_DIM) {
return 0;
}
return __builtin_IB_get_group_id(dim);
}
INLINE size_t OVERLOADABLE get_local_id(uint dim) {
if (dim > MAX_DIM) {
return 0;
} else if (dim == 0) {
return __builtin_IB_get_local_id_x();
} else if (dim == 1) {
return __builtin_IB_get_local_id_y();
} else if (dim == 2) {
return __builtin_IB_get_local_id_z();
}
}
INLINE size_t OVERLOADABLE get_num_groups(uint dim) {
if (dim > MAX_DIM) {
return 1;
}
return __builtin_IB_get_num_groups(dim);
}
INLINE size_t OVERLOADABLE get_global_size(uint dim) {
if (dim > MAX_DIM) {
return 1;
}
return __builtin_IB_get_global_size(dim);
}
INLINE size_t OVERLOADABLE get_local_size(uint dim) {
if (dim > MAX_DIM) {
return 1;
}
return __builtin_IB_get_local_size(dim);
}
INLINE size_t OVERLOADABLE get_global_offset(uint dim) {
if (dim > MAX_DIM) {
return 0;
}
return __builtin_IB_get_global_offset(dim);
}
INLINE size_t OVERLOADABLE get_global_linear_id( void ) {
uint dim = get_work_dim();
size_t result = 0;
switch (dim) {
default:
case 1:
result = get_global_id(0) - get_global_offset(0);
break;
case 2:
result = (get_global_id(1) - get_global_offset(1))*
get_global_size (0) + (get_global_id(0) - get_global_offset(0));
break;
case 3:
result = ((get_global_id(2) - get_global_offset(2)) * get_global_size(1) * get_global_size(0)) +
((get_global_id(1) - get_global_offset(1)) * get_global_size (0)) +
(get_global_id(0) - get_global_offset(0));
break;
}
return result;
}
INLINE size_t OVERLOADABLE get_local_linear_id( void ) {
#if 0
// This doesn't work right now due to a bug in the runtime.
// If/when they fix their bug we can experiment if spending the
// register(s) for get_local_linear_id() is better than spending
// the math to compute the linear local ID.
return __builtin_IB_get_local_linear_id();
#else
uint llid;
llid = (uint)get_local_id(2);
llid *= (uint)get_local_size(1);
llid += (uint)get_local_id(1);
llid *= (uint)get_local_size(0);
llid += (uint)get_local_id(0);
return llid;
#endif
}
uint __intel_get_local_size( void )
{
uint totalWorkGroupSize =
(uint)get_local_size(0) *
(uint)get_local_size(1) *
(uint)get_local_size(2);
return totalWorkGroupSize;
}
uint __intel_get_enqueued_local_size( void )
{
uint totalWorkGroupSize =
(uint)get_enqueued_local_size(0) *
(uint)get_enqueued_local_size(1) *
(uint)get_enqueued_local_size(2);
return totalWorkGroupSize;
}
uint __intel_get_local_linear_id( void )
{
return get_local_linear_id();
}
bool __intel_is_first_work_group_item( void )
{
return get_local_id(0) == 0 &
get_local_id(1) == 0 &
get_local_id(2) == 0;
}
|
