forked from pocl/pocl
-
Notifications
You must be signed in to change notification settings - Fork 0
/
CHANGES
254 lines (215 loc) · 9.16 KB
/
CHANGES
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
0.9 unreleased
===============
[changes checked from bzr history until 2013-05-23]
The OpenCL Runtime (OpenCL 1.2 Chapter 5)
-----------------------------------------
- clCreateImage2D() and clCreateImage3D() implementation moved to
clCreateImage()
- Image creation now uses clCreateBuffer()
- clBuildProgram: Propagate the supported -cl* compiler options to Clang's
OpenCL frontend.
- clFinish: works with commands with event wait lists.
- Refactored clEnqueue* functions to use pocl_create_event/command() functions.
Builtin Function Implementations (OpenCL 1.2 Section 6.12)
----------------------------------------------------------
- Refactored read/write_image()-functions to support refactored device image
object. (Only functions used by SimpleImage test)
- Introduced new macro based implementation for read/write_image()-functions
- Added sampler implementation for CLK_ADDRESS_CLAMP and
CLK_ADDRESS_CLAMP_TO_EDGE (Only integer coords supported)
Device Drivers
--------------
- removed redundancy in pthread and basic:
* pthread now uses read/write_rect implementation from basic
- Added fill_rect function to device driver.
Performance
-----------
- kernel compiler now tries to avoid replicating uniform variables,
this leads to less context data to be saved per work-item and cleaner
kernel bitcode for later optimizations
Misc.
-----
- poclu: helpers for converting between the C float and OpenCL cl_half
types
- clEnqueueNativeKernel implemented
0.8 August 2013
================
This lists only the most interesting changes. Please
refer to the version control log for a full listing.
Overall
-------
- Added support for LLVM/Clang 3.3.
- Dropped support for LLVM/Clang v3.1.
- Removed the depedency on llvm-ld (which was copied to
pocl-llvm-ld to pocl tree). Now uses llvm-link instead.
- Project renamed to Portable Computing Language (pocl).
- Luxmark v2.0 now works.
- x86_64 can now use efficient math built-in function
implementations from the vecmathlib project to avoid libm
calls and to exploit the SIMD instructions more efficiently
in case of vector datatypes in the kernel.
- Parallelize kernel inner loops "horizontally", if possible.
This converts possibly sequential inner kernel loops to parallel
loops by effectively performing "loop interchange" of the
work-item loop and the kernel's inner loop.
- Added VexCL tests to the test suite. All but one of them
work with pocl.
Major bugfixes
--------------
- Fixed passing NULL as a buffer argument to clSetKernelArg
(this time with a regression test added).
- Constant BitCast expressions broken to variables to avoid
crashing when copying a kernel with casts on automatic
local pointers.
- Fixes for i386/i686. Tested on Pentium4/Ubuntu 10.04 LTS.
- Lots of API error checking added (found by the Piglit testing suite).
- Fixed bug in select producing incorrect results when the third
conditional argument is an unsigned scalar or vector.
- Replaced deprecated SSE 4.1 assembly mneunomics in x86-64 min/max
kernel functions that have since been removed in more recent
versions of gas and llvm-as.
- SPIR/LLVM IR 'byval' attributes are now handled correctly on
kernel function arguments, allowing for structs and oversized
vectors to be passed in with value semantics.
- Fixed to work with the latest Khronos OpenCL headers for 1.2.
Some issues fixed with the new cl.hpp.
- The ICD dispatch table was too small which might have caused
"interesting" behavior when calling the later functions in
the table and not using ocl-icd as the dispatcher.
- Several kernel compiler bugs fixed.
- A multithreaded host application could free the same object
multiple times due to a race issue.
Platform Layer implementations (OpenCL 1.2 Chapter 4)
-----------------------------------------------------
- Return correctly formatted CL_DEVICE_VERSION and
CL_DEVICE_OPENCL_C_VERSION.
- clGetDeviceInfo: Use the 'cpufreq' sys interface of Linux for
querying the CPU clock frequency, if available.
The OpenCL Runtime (OpenCL 1.2 Chapter 5)
-----------------------------------------
- clGetEventInfo: Querying the command type, command queue,
and the reference count of the event.
Builtin Function Implementations (OpenCL 1.2 Section 6.12)
----------------------------------------------------------
- convert_type* builtins now generated with a Python script by
Victor Oliveira.
- length() fingerprint was assuming two arguments instead of one.
- The kernel bitcode library is now optimized when built in pocl. Speeds
up kernel optimization for cases which use the kernel functions
a lot.
- Fix mul_hi() implementation
ICD
---
- Fixed pocl tests to work when executed through the Khronos
supplied icd loader (needs a patch applied to the loader be able to
override the .icd search path).
Misc.
-----
- Fix to the helper script search logic:
Search from the BUILDDIR only if env POCL_BUILDING is defined.
Otherwise search from PKGDATADIR first, then from the PATH.
- Fixed memory leaks in clCreateContext* and clCreateKernel
- Ensured that stored arguments are adequately aligned in
clSetKernelArg and clEnqueueNDRangeKernel.
0.7 January 2013
=================
This lists only the most interesting changes. Please
refer to the version control log for a full listing.
Overall
-------
- Support for LLVM 3.2.
- Multi-WI work group functions can be now generated
using loops which are only partially unrolled. Reduces
code size explosion with large WGs in comparison to
the full replication method.
- PowerPC 64 support (tested on Cell/Debian Sid/PS3).
- PowerPC 32 support (tested on Cell/Debian Sid/PS3).
- ARM v7 support (on Linux)
- Beginning of Cell SPU support (very experimental!).
- Most of the AMD APP SDK OpenCL examples now work and have been
added to the pocl test suite.
- Most of the Parboil benchmark cases added to the test
suite.
Kernel Compiler Passes
----------------------
- Several miscompilations and compiler crashes fixed.
- Multiple bugs fixed from the work group vectorizer.
- Updated metadata format pocl uses to pass information
to vectorization and TCE backend to simplify debuging.
- Kernel pointer arguments are not always marked 'noalias' (restricted).
Doing this previously was a specs misunderstanding.
- ConstantGEPs to static variables generated from automated
locals caused problems. Now converting them to normal GEPs
using a pass from the SAFECode project.
OpenCL Platform Layer implementations (OpenCL 1.2 Chapter 4)
-------------------------------------------------------
- clGetDeviceInfo now uses the hwloc lib for device property
queries. Many new queries implemented.
- clGetKernelInfo (initial implementation)
- clGetMemObjectInfo (initial implementation)
- clGetCommandQueueInfo (initial implementation)
- clReleaseDevice
- clRetainDevice
- Proper freeing of devices in clReleaseContext
The OpenCL Runtime Implementations (OpenCL 1.2 Chapter 5)
---------------------------------------------------------
- clBuildProgram: support for passing options to the compiler.
- clEnqueueMarker
OpenCL C Builtin Function Implementations (OpenCL 1.2 Section 6.12)
-------------------------------------------------------------------
- Atomic Functions (6.12.11)
- get_global_offset() was not linked correctly
Framework
---------
- Made it possible to override the .cl -> .bc build command
called by clBuildProgram per device.
Device Drivers
--------------
- pthread/basic:
* extract CPU clock frequency from /proc/cpuinfo, if available
* return cl_khr_fp64 if doubles supported by the CPU
- ttasim: support for explicitly calling custom/special operations
through the vendor extensions API
Misc.
-----
- Fixes for MacOSX builds.
- Fixed passing NULL as a buffer argument to clSetKernelArguments
- Fixed a major bug when launching the same kernel multiple times:
the arguments very not copied to the command object.
- Fixed several issues with ICD, it is now considered stable to be
used by default.
0.6 August 2012
=================
Kernel library
--------------
- Added initial optimized kernel library for X86_64/SSE.
- Preliminary support for ARM architectures on Linux
(briefly tested on MeeGo/Nokia N9).
Pthread device driver
---------------------
- Multithreading at the work group granularity using pthreads.
- Tries to figure out the optimal maximum number of
threads for the system based on the available hardware
threads. Currently works only in Linux using the
/proc/cpuinfo interface.
- Region-based customized memory allocator for speeding up buffer
allocations.
Kernel compiler
---------------
- Most of the tricky work group barrier cases (barriers inside
for-loops etc) now supported.
- Support for local variables, also automatic locals.
- Reuse previous compilation results, if available.
- Automatic vectorization of work groups (multiple work items
in parallel).
Miscellaneous
-------------
- Installable Client Driver (icd) support.
- Event profiling support (incomplete, works only for kernel and
buffer read/write/map/unmap events).
Known issues
------------
- Non-pointer struct kernel arguments fail due to varying ABIs
* https://bugs.launchpad.net/pocl/+bug/987905
- Produces always "fully unrolled" chains of work items for
work groups causing code size explosion for large WGs.