Zhigang Gong [Thu, 8 Aug 2013 07:15:43 +0000 (15:15 +0800)]
GBE: fix insntruction scheduling related bugs in read64/write64.
In read64 and write64, we allocate some temporary registers and
we should put all of those temporary registers may be modified
to the instruction's dst array. Otherwise, the latter post instruction
scheduling may rearrange the instruction incorrectly.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: "Song, Ruiling" <ruiling.song@intel.com>
Zhigang Gong [Tue, 6 Aug 2013 04:01:46 +0000 (12:01 +0800)]
GBE: enable double vector load/store support.
We have some accurate problem for double calculation
on GPU side. I have to change the test case for double
type to add a tolerate error when check the double
data result.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: "Xing, Homer" <homer.xing@intel.com>
Homer Hsing [Wed, 7 Aug 2013 08:28:34 +0000 (16:28 +0800)]
test 64bit-integer selection operator
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 7 Aug 2013 08:28:33 +0000 (16:28 +0800)]
support 64bit-integer selection operator "?:"
v2: reuse MOV to move 64bit integer. not add MOV_INT64 instruction.
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Ruiling Song [Wed, 7 Aug 2013 07:15:50 +0000 (15:15 +0800)]
Implement spill/unspill
The current implementation works like below:
I reserve a pool of registers for spill/reload. Currently 6 registers
are reserved to handle SelectionVector with at most 5 elements.
The other one is used as scratch message header register. The register
after header register was used as the payload for scratch write.
To do spill, just iterate the instructions. If the virtual register
was used as src, insert reload instruction before it. If the virtual
register was used as dst, insert spill instruction to write the register
content to scratch memory.
Limitations yet:
64bit not support.
SelectionVector > 5 not handled.
Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Ruiling Song [Wed, 7 Aug 2013 07:07:40 +0000 (15:07 +0800)]
enable scratch memory allocation and read/write
v2: refine function naming.
Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 7 Aug 2013 05:19:25 +0000 (13:19 +0800)]
keep address space qualifier of pointers
for built-in function with address-space-qualified pointers,
keep address space when accessing the pointers.
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Homer Hsing [Wed, 7 Aug 2013 04:19:54 +0000 (12:19 +0800)]
test 64bit-integer immediate value, and "and", "or", "xor" arithmetic
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 7 Aug 2013 04:19:53 +0000 (12:19 +0800)]
support 64bit-integer AND(&), OR(|), XOR(^) arithmetic
v3: folded similar code into a "for-loop"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 7 Aug 2013 04:19:52 +0000 (12:19 +0800)]
support 64bit-integer immediate value
v3: folded similar code into a "for-loop"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Yi Sun [Tue, 6 Aug 2013 14:02:49 +0000 (22:02 +0800)]
Improve the accuracy of built-in function asin.
Method: asin(x) = x + (1 * x^3)/(2 * 3) + (1 * 3 * x^5)/(2*4 * 5) + \
(1 * 3 * 5 * x^7)/(2*4*6 * 7) + ...
Iterate this for 30 times.
Signed-off-by: Yi Sun <yi.sun@intel.com>
Reviewed-by: Homer Hsing <homer.xing@intel.com>
Homer Hsing [Tue, 6 Aug 2013 07:41:40 +0000 (15:41 +0800)]
support 64bit-integer addition, subtraction
also enable GPU command "subb" (subtract with borrow)
also add test cases
v2: renamed GEN_TYPE_UQ/GEN_TYPE_Q to GEN_TYPE_UL/GEN_TYPE_L
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Ruiling Song [Mon, 5 Aug 2013 07:14:39 +0000 (15:14 +0800)]
Fix a bug in stack calculation.
1. the thread_id is located in r0.5[0-8], so we need to get the correct bits.
2. also, we don't need so much stack size, max_compute_unit have already
been treated as: #EU * max_thread_per_eu.
Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Xing, Homer <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Tue, 6 Aug 2013 06:24:34 +0000 (14:24 +0800)]
support 64bit-integer reading(writing)
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 31 Jul 2013 01:36:57 +0000 (09:36 +0800)]
test if register allocation and 64-bit reading are fixed
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Zhigang Gong [Fri, 2 Aug 2013 17:53:44 +0000 (01:53 +0800)]
GBE: refactor double support.
There are two major issues in double support:
1. Doesn't work at SIMD16 mode.
2. The incorrect usage of vectors. We only need to allocate
those temporary register to contiguous registers.
If you look at the previous implementation of
READ_FLOAT64/WRITE_FLOAT64 in gen_encoder.cpp. You can easily
find it contains many duplicate code and considering the SIMD16
code path never work correctly, it's so difficult to work based
on that code. So I choose to refactor those two major functions.
And refine other parts in the instruction selection stage to fix
the above two major problem with a cleaner code.
Now, it works well on both SIMD16/SIMD8 mode.
Another minor improvement is for the READ_FLOAT64 on SIMD8 mode,
this patch saves one time of send instruction to read all the
8 double data into registers.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Homer Hsing <homer.xing@intel.com>
Homer Hsing [Wed, 31 Jul 2013 02:39:12 +0000 (10:39 +0800)]
test built-in function "shuffle2"
v3: tested two implement of "shuffle2"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Homer Hsing [Wed, 31 Jul 2013 02:39:11 +0000 (10:39 +0800)]
add built-in function "shuffle2"
v3: convert address of "x" to a pointer, then select element by mask
v3: add two-component return-value overloaded version
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Zhigang Gong [Tue, 30 Jul 2013 08:44:47 +0000 (16:44 +0800)]
Need to define local to __local.
It seems that the clang 3.3 already support local/global/private
memory space qualifiers directly. But the previous versions don't
support, we still need to define local as __local here.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Sun, Yi <yi.sun@intel.com>
Zhigang Gong [Thu, 25 Jul 2013 06:52:26 +0000 (14:52 +0800)]
Added memory space parameters support at the autogeneration script.
Enhance the python script to support pointer with memory space
type, such as :
gentype fract (gentype x, __global gentype *iptr)
gentype fract (gentype x, __local gentype *iptr)
gentype fract (gentype x, __private gentype *iptr)
So enable the following function at the builtin function spec file
fract/frexp/modf/nextafter/remquo/sincos.
Remove the duplicate at the ocl_stdlib.tmp.h.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Yi Sun [Tue, 30 Jul 2013 05:11:04 +0000 (13:11 +0800)]
utests: Add a test case for built-in functions get_num_groups.
Signed-off-by: Yi Sun <yi.sun@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Yi Sun [Mon, 29 Jul 2013 15:19:53 +0000 (23:19 +0800)]
utest: Add test for built-in function get_local_id.
Signed-off-by: Yi Sun <yi.sun@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Yi Sun [Mon, 29 Jul 2013 06:52:34 +0000 (14:52 +0800)]
utest: Add test for built-in function get_local_size.
Signed-off-by: Yi Sun <yi.sun@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Zhigang Gong [Mon, 29 Jul 2013 05:11:04 +0000 (13:11 +0800)]
Enable islessgreater/isordered/isunordered builtin vector functions.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Mon, 29 Jul 2013 02:25:18 +0000 (10:25 +0800)]
add built-in function "isordered", "isunordered"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Mon, 29 Jul 2013 02:20:33 +0000 (10:20 +0800)]
add built-in function "islessgreater"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Simon Richter [Thu, 25 Jul 2013 08:31:43 +0000 (10:31 +0200)]
Add generated header and PCH to gitignore
Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Simon Richter [Thu, 25 Jul 2013 08:31:42 +0000 (10:31 +0200)]
Use access() instead of fopen() to search for PCH
Since all we want to do is find out whether the file exists and can be
read, use the appropriate function.
Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Zhigang Gong [Wed, 24 Jul 2013 10:06:06 +0000 (18:06 +0800)]
Fix the indention handling in vector builtin function generator.
As OpenCL use .sa rather than .s10 to represent the tenth element
of a vector. We don't need to adjust one space when we cross the
tenth element (.sa) boundary.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 24 Jul 2013 07:27:10 +0000 (15:27 +0800)]
add address space qualifier to "modf"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 24 Jul 2013 07:20:56 +0000 (15:20 +0800)]
add address space qualifier to "remquo"
renamed origin "remquo" to "__gen_ocl_remquo",
added new "remquo" with address space qualifier
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 24 Jul 2013 06:58:28 +0000 (14:58 +0800)]
revise built-in function "shuffle"
v2 from Zhigang:
Delete the 3-component vectors from the shuffle functions according
to the OpenCL spec.
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Zhigang Gong [Fri, 19 Jul 2013 10:44:08 +0000 (18:44 +0800)]
Add misc builtin vector functions.
Although we don't support them yet.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Simon Richter <Simon.Richter@hogyros.de>
Zhigang Gong [Fri, 19 Jul 2013 09:59:42 +0000 (17:59 +0800)]
check whether python is installed.
We need python to do some source code generation things.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Simon Richter <Simon.Richter@hogyros.de>
Junyan He [Wed, 17 Jul 2013 08:38:41 +0000 (16:38 +0800)]
Add the PCH support when building the source.
Because the utest grows bigger, the runtime compiling time seems a big
problem to cause the utest_run very slow. Most of the time wastes on
re-parsing the big header file such as ocl_stdlib.h.
Using the PCH feature of Clang can shorten the compiling time a lot.
Because we build the pch file at building time, but some source
has runtime building option for LLVM, and these options may cause
the compitable problem between the pch header and the CL source code.
We fallback to build the whole header and source when we find the
options are not compitable.
v2 from Zhigang:
Based on junyan's patch, I made the following changes and fixing:
1. Change the way to find correct pch file, now we just define
two possible locations of the header file at build time. Don't
need to query the linker to get correct pch file directory.
2. Use python to generate a big blob of all ocl standard headers
into one file.
3. Fix those broken dependencies in the CMakeLists.txt, now it
works fine.
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Simon Richter <Simon.Richter@hogyros.de>
Zhigang Gong [Fri, 19 Jul 2013 06:04:48 +0000 (14:04 +0800)]
Split the thounsands autogenerated code out from ocl_stdlib header file.
This patch split the three auto generated code segments from the
huge ocl_stdlib.h. And we rename ocl_stdlib.h to ocl_stdlib.tmpl.h.
The final ocl_stdlib.h will be generated at runtime. It will
insert the ocl_as.h/ocl_convert.h/ocl_vector (which is also
autogenerated at runtime) into ocl_stdlib.tmpl.h's proper positon.
After this patch, we will get a maintainable header file size.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Simon Richter <Simon.Richter@hogyros.de>
Zhigang Gong [Wed, 17 Jul 2013 12:36:36 +0000 (20:36 +0800)]
Implement a pyton script to auto generate those builtin vector functions.
As Beignet is using the SOA model, and we need to lower down all
the vector builtin functions to the scalar version. This type of
thing is ideal to use a script to generate all the code according
to the OpenCL's spec. I just copy/paste most of the prototypes from
the OpenCL spec and put them into builtin_vector_proto.def. Then
I wrote a python script to parse the spec and generate all the
vector inline functions and I removed all existing duplicate
functions in ocl_stdlib.h.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Simon Richter <Simon.Richter@hogyros.de>
Homer Hsing [Wed, 24 Jul 2013 03:05:11 +0000 (11:05 +0800)]
test builtin function "shuffle"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Homer Hsing [Wed, 24 Jul 2013 03:04:57 +0000 (11:04 +0800)]
add builtin function "shuffle"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Zhigang Gong [Wed, 24 Jul 2013 03:09:06 +0000 (11:09 +0800)]
Frexp support global memory directly
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Mon, 22 Jul 2013 08:06:51 +0000 (16:06 +0800)]
add address_space modifier to builtin functions' pointer
I forgot that builtin functions' pointer should be modified
by address_space such like "global", "local", "private"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Homer Hsing [Thu, 18 Jul 2013 06:46:33 +0000 (14:46 +0800)]
test built-in function "remquo"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Homer Hsing [Thu, 18 Jul 2013 06:46:32 +0000 (14:46 +0800)]
add built-in function "remquo"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
Homer Hsing [Wed, 17 Jul 2013 03:00:25 +0000 (11:00 +0800)]
test builtin function "modf"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 17 Jul 2013 03:00:24 +0000 (11:00 +0800)]
add builtin function "modf"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Tue, 16 Jul 2013 08:55:17 +0000 (16:55 +0800)]
test built-in function "nextafter"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Tue, 16 Jul 2013 08:55:16 +0000 (16:55 +0800)]
add built-in function "nextafter"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Yi Sun [Mon, 22 Jul 2013 07:59:42 +0000 (15:59 +0800)]
utest: add built-in test case for get_global_id.
v2. Remove the useless argument in kernel.
Signed-off-by: Yi Sun <yi.sun@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Yang Rong [Mon, 22 Jul 2013 05:24:18 +0000 (13:24 +0800)]
Add build clang option fno-builtin to disable intrinsics.
Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Tue, 16 Jul 2013 06:34:20 +0000 (14:34 +0800)]
test built-in function "frexp"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Tue, 16 Jul 2013 06:34:19 +0000 (14:34 +0800)]
add built-in function "frexp"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Junyan He [Fri, 12 Jul 2013 09:52:59 +0000 (17:52 +0800)]
Improve the clGetMemObjectInfo API, add more info option
Improve the clGetMemObjectInfo API, add more info option.
CL_MEM_ASSOCIATED_MEMOBJECT and CL_MEM_OFFSET need create
subbuffer implememted firstly.
Attach the test case in get_cl_info.cpp
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: "Xing, Homer" <homer.xing@intel.com>
Junyan He [Fri, 12 Jul 2013 08:02:32 +0000 (16:02 +0800)]
Add the support for clSetMemObjectDestructorCallback API
Reviewed-by: "Xing, Homer" <homer.xing@intel.com>
Junyan He [Fri, 12 Jul 2013 06:31:14 +0000 (14:31 +0800)]
Improve the clEnqueueMapBuffer and clCreateBuffer API
In clCreateBuffer API, add the CL_MEM_ALLOC_HOST_PTR and
CL_MEM_USE_HOST_PTR flag support.
CL_MEM_ALLOC_HOST_PTR flag seem nothings special to do.
CL_MEM_USE_HOST_PTR flag will request clEnqueueMapBuffer API:
1> The host_ptr specified in clCreateBuffer is guaranteed to
contain the latest bits in the region being mapped when the
clEnqueueMapBuffer command has completed.
2> The pointer value returned by clEnqueueMapBuffer will be
derived from the host_ptr specified when the buffer object is created.
We improve the clEnqueueMapBuffer to setup a map for the mapped
address and do the data sync problem based on the address when
mapped and unmapped.
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 10 Jul 2013 04:38:58 +0000 (12:38 +0800)]
fix vectorial built-in functions "min, max, clamp"
vectorial versions of "min,max,clamp" was missing.
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Fri, 12 Jul 2013 04:10:26 +0000 (12:10 +0800)]
test built-in function "sign"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Fri, 12 Jul 2013 04:10:25 +0000 (12:10 +0800)]
built-in function "sign"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 10 Jul 2013 02:09:40 +0000 (10:09 +0800)]
test function "mad_sat"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Wed, 10 Jul 2013 02:09:39 +0000 (10:09 +0800)]
add built-in function "mad_sat"
currently only "char, uchar, short, ushort" types of arguments are supported
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Zhigang Gong [Mon, 8 Jul 2013 09:36:46 +0000 (17:36 +0800)]
utest: Query the device driver version and the open cl c version.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Zhigang Gong [Mon, 8 Jul 2013 09:36:45 +0000 (17:36 +0800)]
CL: Refine the version string handling.
Now concentrate the version assignment at the root cmake files.
All the other place will refer the specified macros other than
hard coded a number.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Homer Hsing [Tue, 9 Jul 2013 01:41:59 +0000 (09:41 +0800)]
test function "smoothstep"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Tue, 9 Jul 2013 01:41:58 +0000 (09:41 +0800)]
support built-in function "smoothstep"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Thu, 11 Jul 2013 02:29:39 +0000 (10:29 +0800)]
test API function "clGetImageInfo"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Homer Hsing [Wed, 10 Jul 2013 06:20:57 +0000 (14:20 +0800)]
support clGetImageInfo
clGetImageInfo() is an OpenCL API. It returns information of an image.
This patch makes Piglit test case "clGetImageInfo" pass.
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Homer Hsing [Wed, 10 Jul 2013 05:22:01 +0000 (13:22 +0800)]
improve clCreateContext conformance
OpenCL specification says "clCreateContext" function returns
CL_INVALID_PROPERTY if the same property name appears more than once.
But "clCreateContext" did not follow that.
This patch uses a local temp integer to track whether a property name has
appeared. If so, returns CL_INVALID_PROPERTY.
This patch makes Piglit test case "clCreateContext" pass.
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Homer Hsing [Wed, 10 Jul 2013 01:29:42 +0000 (09:29 +0800)]
test built-in function "bitselect"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang, Rong <rong.r.yang@intel.com>
Homer Hsing [Wed, 10 Jul 2013 01:29:41 +0000 (09:29 +0800)]
support built-in function "bitselect"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Homer Hsing [Mon, 8 Jul 2013 07:50:40 +0000 (15:50 +0800)]
test const-indexed global constant array
A test case is updated, to test three feature:
variable-indexed global constant array
const-indexed global constant array
global constant scalar
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Homer Hsing [Mon, 8 Jul 2013 07:50:39 +0000 (15:50 +0800)]
support const indexed global constant array
support reading global constant arrays by CONST index.
example:
constant int o[3] = {71, 72, 73};
kernel void f(global int *dst) {
dst[get_global_id(0)] = o[2]; // const index: 2
}
in llvm converting phase, calculate offset from const index,
then add the offset to array head
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Homer Hsing [Mon, 8 Jul 2013 02:35:19 +0000 (10:35 +0800)]
test built-in functions "degrees" and "radians"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Yang, Rong R <rong.r.yang@intel.com>
Homer Hsing [Mon, 8 Jul 2013 02:17:47 +0000 (10:17 +0800)]
support built-in functions "degrees" and "radians"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Yang, Rong R <rong.r.yang@intel.com>
Homer Hsing [Thu, 4 Jul 2013 03:19:02 +0000 (11:19 +0800)]
test cases for "mul24", "mad24"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Homer Hsing [Thu, 4 Jul 2013 03:19:01 +0000 (11:19 +0800)]
support built-in functions "mul24", "mad24"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Simon Richter [Fri, 5 Jul 2013 14:50:32 +0000 (16:50 +0200)]
Fix OpenCL C version format
Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Zhigang Gong [Fri, 5 Jul 2013 10:09:33 +0000 (18:09 +0800)]
Bump the version number.
The library version is now 0.2, and the Opencl spec number is 1.1.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Fri, 5 Jul 2013 07:38:05 +0000 (15:38 +0800)]
test built-in function "upsample"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Homer Hsing [Fri, 5 Jul 2013 07:38:04 +0000 (15:38 +0800)]
support built-in function "upsample"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Zhigang Gong [Tue, 25 Jun 2013 10:34:34 +0000 (18:34 +0800)]
Update beignet docs for release v0.2.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Junyan He [Thu, 4 Jul 2013 07:35:34 +0000 (15:35 +0800)]
Add the test case for builtin step() function
The step function has two kind of prototype:
gentypen step(gentypen edge, gentypen x)
and
gentypen step(float edge, gentypen x)
The first's test name is compiler_step_floatX
The second's test name is compiler_stepf_floatX
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Junyan He [Thu, 4 Jul 2013 07:35:28 +0000 (15:35 +0800)]
Add the step builtin function support
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Zhigang Gong [Thu, 4 Jul 2013 12:15:08 +0000 (20:15 +0800)]
CLGL: Refine the hack of gbm extension initialization.
Previous implementation need to refer a EGL internal symbol.
This refinement is also a hack. It just avoid use the EGL internal
symbol, but it does use the internal EGL data structure.
Anyway, before we made some changes to gbm to support our use
model, this ugly method seems the only way to achive sharing
gl 2d/3d textures.
This patch also fix the bug when it failed to get a valid egl
context it may crash.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Zhigang Gong [Thu, 4 Jul 2013 11:12:52 +0000 (19:12 +0800)]
GBE: Clear the value map when start a new scalarize pass.
The scalarize pass is a function pass, and the valueMap should
be a per-function data rather than a per-unit data. The reason
we put it in the unit data structure is that the scalarize pass
is before the GenWriter pass thus there is no ir::Function exists.
As there may be multiple kernel functions in one unit, if we don't
clear the valueMap each time running a new scalarize pass, the previous
data may cause some unexpected behaviour. For example, the previous
instructions have been already erased, then latter a new instruction
in this function may be created in the same position of the erased
instruction, then it breaks this valueMap. That's the root cause why
we run the unit test several times and may encounter an assertion
sometime.
This commit also modify the ir::unit layer implementation to remove
the dependency of llvm from that layer. In general, we should not add
llvm related code to the ir layer.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: He Junyan <junyan.he@linux.intel.com>
Homer Hsing [Thu, 4 Jul 2013 02:45:42 +0000 (10:45 +0800)]
test scalar global constants
a test case was updated, to test not only constant arrays,
but also constant scalar.
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Homer Hsing [Thu, 4 Jul 2013 02:45:41 +0000 (10:45 +0800)]
support global scalar constants
support kind of "constant int n = 1;"
collect scalar value in constant collector
also do virtual register translation for scalar constants
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Junyan He [Wed, 3 Jul 2013 07:17:10 +0000 (15:17 +0800)]
Add the test case for builtin abs_diff() function
All the integer value types check are supported.
Please use the case named compiler_abs_diff_xxxx,
where xxxx means the data type such as int2, char4
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Junyan He [Wed, 3 Jul 2013 07:17:05 +0000 (15:17 +0800)]
Add the abs_diff builtin function support
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Junyan He [Wed, 3 Jul 2013 07:16:59 +0000 (15:16 +0800)]
Add the vector3 support for builtin abs function
Add the forgetten abs vector3 for all the types.
Because the kernel input alignment, improve the test
case to match the alignment request.
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Zhigang Gong [Wed, 3 Jul 2013 09:36:16 +0000 (17:36 +0800)]
CL: destroy the EGL image which is created for gl sharing when delete the mem object.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Junyan He [Wed, 3 Jul 2013 04:41:17 +0000 (12:41 +0800)]
Modify all the builtin function vect return to (vect_name)(e1, e2, e3)
Some builtin functions has the prototype like:
int3 function_name (int3 x) { return (x.s0, x.s1, x.s2);}
which not comply with CL spec and will cause the clang IR
be translated error.
The vector declare should be (vect)(e1, e2, e3)
Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Ruiling Song [Tue, 2 Jul 2013 08:44:43 +0000 (16:44 +0800)]
Disable error message output in release version.
As piglit will got the error message we output to stderr and mark the case 'WARN'.
so, we disable the message to stderr, and use release version to run piglit.
also fix a minor compile fail under release version.
Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Yang Rong [Tue, 2 Jul 2013 07:22:24 +0000 (15:22 +0800)]
Fix atomic test failed in GT1.
Barrier only ensure one work group finish, can't guarantee all work item's atomic ops
have finished before the last atomic_add.
So use atomic_xchg to update first work group's local buffer to other global buffer position.
Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Xing, Homer <homer.xing@intel.com>
Zhigang Gong [Tue, 2 Jul 2013 10:44:19 +0000 (18:44 +0800)]
utests: increase local size in the two barrier test cases.
Increasing the local size to 256 to bring more pressure
to barrier testing.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Zou, Nanhai <nanhai.zou@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Tested-by: Sun, Yi <yi.sun@intel.com>
Zhigang Gong [Tue, 2 Jul 2013 10:40:08 +0000 (18:40 +0800)]
GBE: fixed a barrier related bug.
Actually, this commit fixed two bugs related to barrier.
1. We should set useSLM to true if we use barrier.
2. We need to set barrier id to the barrierMsg payload according to
r0.2. And we don't need to reprogram the barrierCount.
And after this fix, we don't need the work around for the local
memory barrier, thus we don't need the memory fence for local memory
barrier.
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Zou, Nanhai <nanhai.zou@intel.com>
Reviewed-by: Yang, Rong R <rong.r.yang@intel.com>
Tested-by: Sun, Yi <yi.sun@intel.com>
Homer Hsing [Tue, 2 Jul 2013 06:45:16 +0000 (14:45 +0800)]
test cases for built-in functions "mul_hi", "mad_hi"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Homer Hsing [Tue, 2 Jul 2013 06:45:10 +0000 (14:45 +0800)]
support built-in functions "mul_hi", "mad_hi"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Homer Hsing [Tue, 2 Jul 2013 06:44:36 +0000 (14:44 +0800)]
test cases for "hadd", "rhadd"
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Homer Hsing [Tue, 2 Jul 2013 06:44:30 +0000 (14:44 +0800)]
support built-in functions "hadd", "rhadd"
backend now support GPU opcode "addc".
add built-in functions "hadd", "rhadd".
Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
Ruiling Song [Mon, 1 Jul 2013 05:46:21 +0000 (13:46 +0800)]
add test case for clGetContextInfo/clGetKernelInfo
Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Ruiling Song [Mon, 1 Jul 2013 05:46:20 +0000 (13:46 +0800)]
Implement API clGetKernelInfo
Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>