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
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
|
#include "string"
namespace gbe {
std::string ocl_stdlib_str =
"/* \n"
" * Copyright © 2012 Intel Corporation\n"
" *\n"
" * This library is free software; you can redistribute it and/or\n"
" * modify it under the terms of the GNU Lesser General Public\n"
" * License as published by the Free Software Foundation; either\n"
" * version 2 of the License, or (at your option) any later version.\n"
" *\n"
" * This library is distributed in the hope that it will be useful,\n"
" * but WITHOUT ANY WARRANTY; without even the implied warranty of\n"
" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU\n"
" * Lesser General Public License for more details.\n"
" *\n"
" * You should have received a copy of the GNU Lesser General Public\n"
" * License along with this library. If not, see <http://www.gnu.org/licenses/>.\n"
" *\n"
" * Author: Benjamin Segovia <benjamin.segovia@intel.com>\n"
" */\n"
"\n"
"#ifndef __GEN_OCL_STDLIB_H__\n"
"#define __GEN_OCL_STDLIB_H__\n"
"\n"
"#define INLINE_OVERLOADABLE __attribute__((overloadable,always_inline))\n"
"#define OVERLOADABLE __attribute__((overloadable))\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// OpenCL basic types\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"typedef unsigned int uint;\n"
"typedef unsigned int size_t;\n"
"typedef float float2 __attribute__((ext_vector_type(2)));\n"
"typedef float float3 __attribute__((ext_vector_type(3)));\n"
"typedef float float4 __attribute__((ext_vector_type(4)));\n"
"typedef float float8 __attribute__((ext_vector_type(8)));\n"
"typedef float float16 __attribute__((ext_vector_type(16)));\n"
"typedef int int2 __attribute__((ext_vector_type(2)));\n"
"typedef int int3 __attribute__((ext_vector_type(3)));\n"
"typedef int int4 __attribute__((ext_vector_type(4)));\n"
"typedef int int8 __attribute__((ext_vector_type(8)));\n"
"typedef int int16 __attribute__((ext_vector_type(16)));\n"
"typedef unsigned int uint2 __attribute__((ext_vector_type(2)));\n"
"typedef unsigned uint3 __attribute__((ext_vector_type(3)));\n"
"typedef unsigned uint4 __attribute__((ext_vector_type(4)));\n"
"typedef unsigned uint8 __attribute__((ext_vector_type(8)));\n"
"typedef unsigned uint16 __attribute__((ext_vector_type(16)));\n"
"typedef bool bool2 __attribute__((ext_vector_type(2)));\n"
"typedef bool bool3 __attribute__((ext_vector_type(3)));\n"
"typedef bool bool4 __attribute__((ext_vector_type(4)));\n"
"typedef bool bool8 __attribute__((ext_vector_type(8)));\n"
"typedef bool bool16 __attribute__((ext_vector_type(16)));\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// OpenCL address space\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"#define __private __attribute__((address_space(0)))\n"
"#define __global __attribute__((address_space(1)))\n"
"#define __constant __attribute__((address_space(2)))\n"
"//#define __local __attribute__((address_space(3)))\n"
"#define global __global\n"
"//#define local __local\n"
"#define constant __constant\n"
"#define private __private\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// Work groups and work items functions\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"#define DECL_INTERNAL_WORK_ITEM_FN(NAME) \\\n"
"__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##0(void); \\\n"
"__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##1(void); \\\n"
"__attribute__((pure,const)) unsigned int __gen_ocl_##NAME##2(void);\n"
"DECL_INTERNAL_WORK_ITEM_FN(get_group_id)\n"
"DECL_INTERNAL_WORK_ITEM_FN(get_local_id)\n"
"DECL_INTERNAL_WORK_ITEM_FN(get_local_size)\n"
"DECL_INTERNAL_WORK_ITEM_FN(get_global_size)\n"
"DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)\n"
"#undef DECL_INTERNAL_WORK_ITEM_FN\n"
"\n"
"#define DECL_PUBLIC_WORK_ITEM_FN(NAME) \\\n"
"inline unsigned NAME(unsigned int dim) { \\\n"
" if (dim == 0) return __gen_ocl_##NAME##0(); \\\n"
" else if (dim == 1) return __gen_ocl_##NAME##1(); \\\n"
" else if (dim == 2) return __gen_ocl_##NAME##2(); \\\n"
" else return 0; \\\n"
"}\n"
"DECL_PUBLIC_WORK_ITEM_FN(get_group_id)\n"
"DECL_PUBLIC_WORK_ITEM_FN(get_local_id)\n"
"DECL_PUBLIC_WORK_ITEM_FN(get_local_size)\n"
"DECL_PUBLIC_WORK_ITEM_FN(get_global_size)\n"
"DECL_PUBLIC_WORK_ITEM_FN(get_num_groups)\n"
"#undef DECL_PUBLIC_WORK_ITEM_FN\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// Vector loads and stores\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"\n"
"// These loads and stores will use untyped reads and writes, so we can just\n"
"// cast to vector loads / stores. Not C99 compliant BTW due to aliasing issue.\n"
"// Well we do not care, we do not activate TBAA in the compiler\n"
"#define DECL_UNTYPED_RW_SPACE_N(TYPE, DIM, SPACE) \\\n"
"__attribute__((always_inline, overloadable)) \\\n"
"inline TYPE##DIM vload##DIM(size_t offset, const SPACE TYPE *p) { \\\n"
" return *(SPACE TYPE##DIM *) (p + DIM * offset); \\\n"
"} \\\n"
"__attribute__((always_inline, overloadable)) \\\n"
"inline void vstore##DIM(TYPE##DIM v, size_t offset, SPACE TYPE *p) { \\\n"
" *(SPACE TYPE##DIM *) (p + DIM * offset) = v; \\\n"
"}\n"
"\n"
"#define DECL_UNTYPED_RW_ALL_SPACE(TYPE, SPACE) \\\n"
" DECL_UNTYPED_RW_SPACE_N(TYPE, 2, SPACE) \\\n"
" DECL_UNTYPED_RW_SPACE_N(TYPE, 3, SPACE) \\\n"
" DECL_UNTYPED_RW_SPACE_N(TYPE, 4, SPACE) \\\n"
" DECL_UNTYPED_RW_SPACE_N(TYPE, 8, SPACE) \\\n"
" DECL_UNTYPED_RW_SPACE_N(TYPE, 16, SPACE)\n"
"\n"
"#define DECL_UNTYPED_RW_ALL(TYPE) \\\n"
" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __global) \\\n"
" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __local) \\\n"
" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __constant) \\\n"
" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __private)\n"
"\n"
"DECL_UNTYPED_RW_ALL(float)\n"
"DECL_UNTYPED_RW_ALL(uint)\n"
"DECL_UNTYPED_RW_ALL(int)\n"
"\n"
"#undef DECL_UNTYPED_RW_ALL\n"
"#undef DECL_UNTYPED_RW_ALL_SPACE\n"
"#undef DECL_UNTYPED_RW_SPACE_N\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// Arithmetic functions\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"__attribute__((always_inline))\n"
"inline uint get_global_id(uint dim) {\n"
" return get_local_id(dim) + get_local_size(dim) * get_group_id(dim);\n"
"}\n"
"\n"
"__attribute__ ((pure, const, overloadable)) float mad(float a, float b, float c);\n"
"__attribute__((overloadable, always_inline))\n"
"inline uint select(uint src0, uint src1, uint cond) {\n"
" return cond ? src0 : src1;\n"
"}\n"
"__attribute__((overloadable, always_inline))\n"
"inline int select(int src0, int src1, int cond) {\n"
" return cond ? src0 : src1;\n"
"}\n"
"\n"
"// This will be optimized out by LLVM and will output LLVM select instructions\n"
"#define DECL_SELECT4(TYPE4, TYPE, COND_TYPE4, MASK) \\\n"
"__attribute__((overloadable)) \\\n"
"inline TYPE4 select(TYPE4 src0, TYPE4 src1, COND_TYPE4 cond) { \\\n"
" TYPE4 dst; \\\n"
" const TYPE x0 = src0.x; /* Fix performance issue with CLANG */ \\\n"
" const TYPE x1 = src1.x; \\\n"
" const TYPE y0 = src0.y; \\\n"
" const TYPE y1 = src1.y; \\\n"
" const TYPE z0 = src0.z; \\\n"
" const TYPE z1 = src1.z; \\\n"
" const TYPE w0 = src0.w; \\\n"
" const TYPE w1 = src1.w; \\\n"
" dst.x = (cond.x & MASK) ? x1 : x0; \\\n"
" dst.y = (cond.y & MASK) ? y1 : y0; \\\n"
" dst.z = (cond.z & MASK) ? z1 : z0; \\\n"
" dst.w = (cond.w & MASK) ? w1 : w0; \\\n"
" return dst; \\\n"
"}\n"
"DECL_SELECT4(int4, int, int4, 0x80000000)\n"
"DECL_SELECT4(float4, float, int4, 0x80000000)\n"
"#undef DECL_SELECT4\n"
"\n"
"INLINE_OVERLOADABLE inline float2 mad(float2 a, float2 b, float2 c) {\n"
" return (float2)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y));\n"
"}\n"
"INLINE_OVERLOADABLE inline float3 mad(float3 a, float3 b, float3 c) {\n"
" return (float3)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y), mad(a.z,b.z,c.z));\n"
"}\n"
"INLINE_OVERLOADABLE inline float4 mad(float4 a, float4 b, float4 c) {\n"
" return (float4)(mad(a.x,b.x,c.x), mad(a.y,b.y,c.y),\n"
" mad(a.z,b.z,c.z), mad(a.w,b.w,c.w));\n"
"}\n"
"\n"
"#define DECL_MIN_MAX(TYPE) \\\n"
"INLINE_OVERLOADABLE inline TYPE max(TYPE a, TYPE b) { \\\n"
" return a > b ? a : b; \\\n"
"} \\\n"
"INLINE_OVERLOADABLE inline TYPE min(TYPE a, TYPE b) { \\\n"
" return a < b ? a : b; \\\n"
"}\n"
"DECL_MIN_MAX(float)\n"
"DECL_MIN_MAX(int)\n"
"DECL_MIN_MAX(short)\n"
"DECL_MIN_MAX(char)\n"
"DECL_MIN_MAX(uint)\n"
"DECL_MIN_MAX(unsigned short)\n"
"DECL_MIN_MAX(unsigned char)\n"
"#undef DECL_MIN_MAX\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// Extensions to manipulate the register file\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"\n"
"// Direct addressing register regions\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int);\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int, int);\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int, int, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int, int, int, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_region(int offset, int vstride, int width, int hstride, int, int, int, int, int, int, int, int);\n"
"\n"
"// Gather from register file\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int);\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int, int);\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int, int, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int, int, int, int, int, int, int);\n"
"OVERLOADABLE int __gen_ocl_rgather(unsigned short index, int, int, int, int, int, int, int, int);\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// Extension to have uniform condition per hardware thread\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"\n"
"OVERLOADABLE unsigned short __gen_ocl_any(unsigned short cond);\n"
"OVERLOADABLE unsigned short __gen_ocl_all(unsigned short cond);\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// Extension to support OBlock reads / writes\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"\n"
"OVERLOADABLE int __gen_ocl_obread(const __global void *address);\n"
"OVERLOADABLE int __gen_ocl_obread(const __constant void *address);\n"
"OVERLOADABLE int __gen_ocl_obread(const __local void *address);\n"
"OVERLOADABLE void __gen_ocl_obwrite(const __global void *address, int);\n"
"OVERLOADABLE void __gen_ocl_obwrite(const __local void *address, int);\n"
"\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"// Force the compilation to SIMD8 or SIMD16\n"
"/////////////////////////////////////////////////////////////////////////////\n"
"\n"
"int __gen_ocl_force_simd8(void);\n"
"int __gen_ocl_force_simd16(void);\n"
"\n"
"#define DECL_VOTE(TYPE) \\\n"
"__attribute__((overloadable,always_inline)) \\\n"
"TYPE __gen_ocl_any(TYPE cond) { \\\n"
" return (TYPE) __gen_ocl_any((unsigned short) cond); \\\n"
"} \\\n"
"__attribute__((overloadable,always_inline)) \\\n"
"TYPE __gen_ocl_all(TYPE cond) { \\\n"
" return (TYPE) __gen_ocl_all((unsigned short) cond); \\\n"
"}\n"
"DECL_VOTE(unsigned int)\n"
"DECL_VOTE(unsigned char)\n"
"DECL_VOTE(int)\n"
"DECL_VOTE(char)\n"
"DECL_VOTE(short)\n"
"DECL_VOTE(bool)\n"
"#undef DECL_VOTE\n"
"\n"
"#define NULL ((void*)0)\n"
"#undef INLINE_OVERLOADABLE\n"
"#endif /* __GEN_OCL_STDLIB_H__ */\n"
"\n"
;
}
|
