contrib/beignet.git
11 years agoadd built-in function "frexp"
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>
11 years agoImprove the clGetMemObjectInfo API, add more info option
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>
11 years agoAdd the support for clSetMemObjectDestructorCallback API
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>
11 years agoImprove the clEnqueueMapBuffer and clCreateBuffer API
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>
11 years agofix vectorial built-in functions "min, max, clamp"
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>
11 years agotest built-in function "sign"
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>
11 years agobuilt-in function "sign"
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>
11 years agotest function "mad_sat"
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>
11 years agoadd built-in function "mad_sat"
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>
11 years agoutest: Query the device driver version and the open cl c version.
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>
11 years agoCL: Refine the version string handling.
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>
11 years agotest function "smoothstep"
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>
11 years agosupport built-in function "smoothstep"
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>
11 years agotest API function "clGetImageInfo"
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>
11 years agosupport clGetImageInfo
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>
11 years agoimprove clCreateContext conformance
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>
11 years agotest built-in function "bitselect"
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>
11 years agosupport built-in function "bitselect"
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>
11 years agotest const-indexed global constant array
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>
11 years agosupport const indexed global constant array
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>
11 years agotest built-in functions "degrees" and "radians"
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>
11 years agosupport built-in functions "degrees" and "radians"
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>
11 years agotest cases for "mul24", "mad24"
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>
11 years agosupport built-in functions "mul24", "mad24"
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>
11 years agoFix OpenCL C version format
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>
11 years agoBump the version number.
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>
11 years agotest built-in function "upsample"
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>
11 years agosupport built-in function "upsample"
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>
11 years agoUpdate beignet docs for release v0.2.
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>
11 years agoAdd the test case for builtin step() function
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>
11 years agoAdd the step builtin function support
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>
11 years agoCLGL: Refine the hack of gbm extension initialization.
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>
11 years agoGBE: Clear the value map when start a new scalarize pass.
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>
11 years agotest scalar global constants
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>
11 years agosupport global scalar constants
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>
11 years agoAdd the test case for builtin abs_diff() function
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>
11 years agoAdd the abs_diff builtin function support
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>
11 years agoAdd the vector3 support for builtin abs function
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>
11 years agoCL: destroy the EGL image which is created for gl sharing when delete the mem object.
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>
11 years agoModify all the builtin function vect return to (vect_name)(e1, e2, e3)
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>
11 years agoDisable error message output in release version.
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>
11 years agoFix atomic test failed in GT1.
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>
11 years agoutests: increase local size in the two barrier test cases.
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>
11 years agoGBE: fixed a barrier related bug.
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>
11 years agotest cases for built-in functions "mul_hi", "mad_hi"
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>
11 years agosupport built-in functions "mul_hi", "mad_hi"
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>
11 years agotest cases for "hadd", "rhadd"
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>
11 years agosupport built-in functions "hadd", "rhadd"
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>
11 years agoadd test case for clGetContextInfo/clGetKernelInfo
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>
11 years agoImplement API clGetKernelInfo
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>
11 years agoImplement clGetContextInfo
Ruiling Song [Mon, 1 Jul 2013 05:46:19 +0000 (13:46 +0800)]
Implement clGetContextInfo

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoput 64-bit float test cases at tail
Homer Hsing [Fri, 28 Jun 2013 04:30:44 +0000 (12:30 +0800)]
put 64-bit float test cases at tail

64-bit float test cases set SIMD width to 8,
causing failure in GT1 type of Intel GPU card.

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd the test case for builtin abs() function
Junyan He [Mon, 1 Jul 2013 07:00:35 +0000 (15:00 +0800)]
Add the test case for builtin abs() function

All the integer value types check are supported.
Please use the case named compiler_abs_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>
Reviewed-by: Xing, Homer <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
11 years agoAdd the builtin function abs() support
Junyan He [Mon, 1 Jul 2013 07:00:27 +0000 (15:00 +0800)]
Add the builtin function abs() support

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Xing, Homer <homer.xing@intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
11 years agoClear atomic dst buffer to fix atomic random fail.
Yang Rong [Mon, 1 Jul 2013 05:38:48 +0000 (13:38 +0800)]
Clear atomic dst buffer to fix atomic random fail.

Because atomic's address used as src and dst, so need to clear this address.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoFix options parse infinite loop bug.
Yang Rong [Mon, 1 Jul 2013 02:06:17 +0000 (10:06 +0800)]
Fix options parse infinite loop bug.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoTest case for vector type comparison results.
Edward Ching [Thu, 27 Jun 2013 05:09:21 +0000 (22:09 -0700)]
Test case for vector type comparison results.

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoGBE: fixed the bug when sext a i1 to i8/i16/i32.
Zhigang Gong [Wed, 26 Jun 2013 05:00:58 +0000 (13:00 +0800)]
GBE: fixed the bug when sext a i1 to i8/i16/i32.

We need to extent it to -1 rather than 1.
Reported by Edward Ching <edward.k.ching@gmail.com>.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Edward Ching <edward.k.ching@gmail.com>
11 years agoEnable int32 atomic and fp64 extensions.
Yang Rong [Fri, 28 Jun 2013 02:50:18 +0000 (10:50 +0800)]
Enable int32 atomic and fp64 extensions.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoCL: remove the deprecated function clSetCommandQueueProperty.
Zhigang Gong [Thu, 27 Jun 2013 11:38:25 +0000 (19:38 +0800)]
CL: remove the deprecated function clSetCommandQueueProperty.

According to the Open CL spec:
The clSetCommandQueueProperty API is no longer supported in OpenCL 1.1.
so we remove the corresponding function here.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoFix some math function error in simd16.
Yang Rong [Thu, 27 Jun 2013 08:47:58 +0000 (16:47 +0800)]
Fix some math function error in simd16.

INT DIV splite to simd8 but forget to set quarter_control.
Will fail when predication enable.
Change the atomic test case to trigger this bug.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd atomic test case.
Yang Rong [Thu, 27 Jun 2013 08:47:57 +0000 (16:47 +0800)]
Add atomic test case.

The test case include local memory and global memory, atomic operations from
different threads and different work groups.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd all atomic built-in functions.
Yang Rong [Thu, 27 Jun 2013 08:47:56 +0000 (16:47 +0800)]
Add all atomic built-in functions.

Treat all atomic function's operands as unsigned int, except imin/imax.
So use the different function __gen_ocl_atomic_umin and __gen_ocl_atomic_imin.
Overload different address space, local and global.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd atomic help functions.
Yang Rong [Thu, 27 Jun 2013 08:47:55 +0000 (16:47 +0800)]
Add atomic help functions.

Use the data port message "Untyped Atomic Operation" to implement openCL atomic functions.
This message can handle both global memory and SLM.
For all atomic functions, the operands will be treated as unsigned int, except imax/imin.
Only add one opcode Atomic for all atomic functions in ir, and identify the atomic function
by ir::AtomicOps.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agotest case for function "rotate"
Homer Hsing [Thu, 27 Jun 2013 02:58:16 +0000 (10:58 +0800)]
test case for function "rotate"

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
11 years agosupport build-in function "rotate"
Homer Hsing [Thu, 27 Jun 2013 02:58:15 +0000 (10:58 +0800)]
support build-in function "rotate"

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: He Junyan <junyan.he@inbox.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
11 years agoRefine the get_local_id/... builtins.
Zhigang Gong [Tue, 25 Jun 2013 06:15:09 +0000 (14:15 +0800)]
Refine the get_local_id/... builtins.

As we could prepare correct value on runtime library side and give
a correct value in the payload for dim 0, 1 and 2. So for these 3
dim argument, we don't need to check it whether in the valid range,
we just read the payload's value.

This way, we can avoid any unecessary branching for normal usage of
these builtin functions. And could avoid a known bool related bug.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Sun, Yi <yi.sun@intel.com>
11 years agoutests: Add basic arithmetic test case
Ruiling Song [Wed, 26 Jun 2013 07:52:13 +0000 (15:52 +0800)]
utests: Add basic arithmetic test case

test case for + - * / % of data type (u)int8/16/32
remove duplicated cases.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoGBE: Add more support of char and short arithmetic
Ruiling Song [Wed, 26 Jun 2013 07:52:12 +0000 (15:52 +0800)]
GBE: Add more support of char and short arithmetic

add * / % support of char and short type.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agotest case for function "clz"
Homer Hsing [Wed, 26 Jun 2013 05:11:54 +0000 (13:11 +0800)]
test case for function "clz"

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Junyan He <junyan.he@linux.intel.com>
11 years agosupport zero bit counting
Homer Hsing [Wed, 26 Jun 2013 05:11:53 +0000 (13:11 +0800)]
support zero bit counting

support OpenCL built-in function "clz", returning number of leading zero bits

add GEN GPU "fbh", "fbl" instructions, for counting zero bits

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Junyan He <junyan.he@linux.intel.com>
11 years agoAdd the test case for clGetProgramInfo API
Junyan He [Wed, 26 Jun 2013 03:34:20 +0000 (11:34 +0800)]
Add the test case for clGetProgramInfo API

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd the test case for clGetCommandQueueInfo API
Junyan He [Tue, 25 Jun 2013 10:15:32 +0000 (18:15 +0800)]
Add the test case for clGetCommandQueueInfo API

Because all the get clGetXXXInfo API have similar
structure in function type, we will integrate them
together, and rename the get_program_info.cpp to get_cl_info.cpp

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd the support of the API: clGetCommandQueueInfo
Junyan He [Tue, 25 Jun 2013 10:15:25 +0000 (18:15 +0800)]
Add the support of the API: clGetCommandQueueInfo

Though we support get the CL_QUEUE_PROPERTIES,
but because the CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
and CL_QUEUE_PROFILING_ENABLE will never be set when
create the queue, it just return a all 0 bitfield now.

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Song, Ruiling <ruiling.song@intel.com>
11 years agoDocs: Rearrange documents according to wiki page on fd.o.
Zhigang Gong [Tue, 25 Jun 2013 10:04:53 +0000 (18:04 +0800)]
Docs: Rearrange documents according to wiki page on fd.o.

We use fd.o wiki to host our documents. To make the maintainence easier,
I change the direcotry and some links in the markdown files. And make
them the same as fd.o's.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoutests: Enable the local memory barrier test case.
Zhigang Gong [Tue, 25 Jun 2013 06:24:44 +0000 (14:24 +0800)]
utests: Enable the local memory barrier test case.

And fix one typo. Without the local memory barrier work around,
this test case will fail.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoGBE: work around the local memory barrier fence issue.
Zhigang Gong [Tue, 25 Jun 2013 06:22:43 +0000 (14:22 +0800)]
GBE: work around the local memory barrier fence issue.

According to the spec, we do not need to add a memory fence
for SLM's access. But, in pratice, we have to add it to make
sure get correct result. Don't know the root cause yet, just
use this work around now.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoGBE: Fix a bug in encoding MATH instruction
Ruiling Song [Mon, 24 Jun 2013 02:53:57 +0000 (10:53 +0800)]
GBE: Fix a bug in encoding MATH instruction

For std::vector, a push_back may cause memory relocation if no enough
memory in the vector pool. And iterator or pointer got previously will
become invalid after relocation.

Here in GenEncoder::next(), which will call push_back(), memory
relocation may occur. Then relocation will make 'insn' point to
invalid memory that does not belong to GenEncoder::store anymore.
Then, the setting of execution_width will fail.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoutests: Fix a typo.
Zhigang Gong [Mon, 24 Jun 2013 06:57:45 +0000 (14:57 +0800)]
utests: Fix a typo.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd one test case for clGetProgramInfo
Junyan He [Thu, 20 Jun 2013 05:58:14 +0000 (13:58 +0800)]
Add one test case for clGetProgramInfo

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoFix several typos in unit test.
Ruiling Song [Wed, 19 Jun 2013 02:04:54 +0000 (10:04 +0800)]
Fix several typos in unit test.

compiler_sub_bytes and compiler_sub_shorts

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoEnable generation of convert_ and as_ functions for double
Simon Richter [Wed, 19 Jun 2013 10:17:58 +0000 (12:17 +0200)]
Enable generation of convert_ and as_ functions for double

Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoDefine double vector types
Simon Richter [Wed, 19 Jun 2013 10:17:57 +0000 (12:17 +0200)]
Define double vector types

Add the definition for the "doubleN" vector types

Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoEnable cl_khr_fp64 extension for OpenCL stdlib header
Simon Richter [Wed, 19 Jun 2013 10:17:56 +0000 (12:17 +0200)]
Enable cl_khr_fp64 extension for OpenCL stdlib header

This allows the stdlib header to define overloads for doubles.

Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd long and ulong types to generated functions.
Simon Richter [Thu, 13 Jun 2013 11:15:52 +0000 (13:15 +0200)]
Add long and ulong types to generated functions.

This enables all generated functions for 64 bit integers.

Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: He Junyan <junyan.he@linux.intel.com>
11 years agoDefine all convert_* functions.
Simon Richter [Thu, 13 Jun 2013 11:15:51 +0000 (13:15 +0200)]
Define all convert_* functions.

These functions convert between vectors of the same length by casting each
member in turn.

Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: He Junyan <junyan.he@linux.intel.com>
11 years agoGenerate all supported as_* functions
Simon Richter [Thu, 13 Jun 2013 11:15:50 +0000 (13:15 +0200)]
Generate all supported as_* functions

This adds support for all reinterpreting type conversions currently
possible.

The conversion functions can be updated by invoking the
update_as.sh script.

Signed-off-by: Simon Richter <Simon.Richter@hogyros.de>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: He Junyan <junyan.he@linux.intel.com>
11 years agoadd a lost special register name
Homer Hsing [Fri, 21 Jun 2013 01:48:32 +0000 (09:48 +0800)]
add a lost special register name

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agotest cases for 64-bit float
Homer Hsing [Fri, 21 Jun 2013 04:26:32 +0000 (12:26 +0800)]
test cases for 64-bit float

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Yang Rong <rong.r.yang@intel.com>
11 years agoSupport 64-bit float
Homer Hsing [Fri, 21 Jun 2013 04:26:31 +0000 (12:26 +0800)]
Support 64-bit float

support:
  arithmetic(+ - *)
  store load
  immediate_value
  if else
  select

other change:
  add "nib control" field in machine instruction format
  support "nib control"
  fix "directly store float-64 after load float-64".
  change hard coded store size (4) to flexible size (4 or 8)
  when using float-64 load(store), change to SIMD8

example:

/* support arithmetic store load immediate_value */
kernel void f(global double *src, global double *dst) {
  int i = get_global_id(0);
  double d = 1.234567890123456789;
  dst[i] = d * (src[i] + d);
}

/* support if else */
kernel void f(global float *src, global double *dst) {
  int i = get_global_id(0);
  float d = 1.234567890123456789f;
  if (i < 14)
    dst[i] = d * (d + src[i]);
  else
    dst[i] = 14;
}

/* support select */
kernel void f(global float *src, global double *dst) {
  int i = get_global_id(0);
  float d = 1.234567890123456789f;
  dst[i] = i < 14 ? d : 14;
}

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Yang Rong <rong.r.yang@intel.com>
11 years agoGBE: Fixed one bug in scalarize pass
Zhigang Gong [Wed, 19 Jun 2013 10:30:47 +0000 (18:30 +0800)]
GBE: Fixed one bug in scalarize pass

I met segfault at void Scalarize::dce() randomly when I integrate
a openCL kernel to the chromium's GPU process. After discuss with
Yang Rong, I found one bug in this function. As it use two loops
to erase the dead instructions, but it doesn't set the pointer to
NULL at the first loop when it already erased the instruction. Thus
at the second loop, when it call (*i)->getParent, the (*i) may
already be deleted then it may refer a freed region and may cause
segfault.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>
11 years agoFix some piglit constant buffer tests fail.
Yang Rong [Wed, 19 Jun 2013 07:36:34 +0000 (15:36 +0800)]
Fix some piglit constant buffer tests fail.

If indirect move's source is scalrar reg, such as using cb[0] in kernel,
should not unpack.

Change test case compiler_function_constant0 to trigger this bug.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoupdate to OpenCL 1.1 header
Homer Hsing [Wed, 19 Jun 2013 06:49:51 +0000 (14:49 +0800)]
update to OpenCL 1.1 header

Based on our current progress, we plan to implement OpenCL 1.1 rather
than OpenCL 1.2 for the next release. Thus we downgrade the header file
to 1.1 in this commit.

put OpenCL 1.1 header in include/CL/

add OpenCL 1.2 defs, used by some code, by
 "#ifndef CL_VERSION_1.2"
 " some OpenCL 1.2 defs"
 "#endif"

add OCL_CREATE_IMAGE2D, OCL_CREATE_IMAGE3D,
    OCL_CREATE_GL_IMAGE2D, OCL_CREATE_GL_IMAGE3D

update test cases

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoadd the support of clGetProgramBuildInfo and clGetProgramInfo
Junyan He [Tue, 18 Jun 2013 08:44:49 +0000 (16:44 +0800)]
add the support of clGetProgramBuildInfo and clGetProgramInfo

For clGetProgramBuildInfo,
CL_BUILD_IN_PROGRESS not support now
and CL_PROGRAM_BUILD_LOG need do add the info collection
logic in backend and not support too, just return null
string now.
clGetProgramInfo all are fully supported.

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Ruiling Song <ruiling.song@intel.com>
11 years agoSupport multi-source CL program
Ruiling Song [Wed, 19 Jun 2013 02:30:21 +0000 (10:30 +0800)]
Support multi-source CL program

Also fix several errors on clBuildProgram.

v2: minor refine of cl_program_release_sources()

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: He Junyan <junyan.he@linux.intel.com>
11 years agoutests: Add a new local memory barrier case
Yang Rong [Tue, 18 Jun 2013 09:31:10 +0000 (17:31 +0800)]
utests: Add a new local memory barrier case

We fail this case right now, disalbe it. And need more work
to check the root cause. Change the local memory barrier can
pass it, but it doesn't comply with BSPEC which says SLM doesn't
need a memory fence.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoUtest: Add a test case for validating built-in function get_global_size()
Yi Sun [Thu, 13 Jun 2013 07:12:44 +0000 (15:12 +0800)]
Utest: Add a test case for validating built-in function get_global_size()

v1:
According to the OpenCL v1.1 & v1.2 chapter 6.11, the behavior of function get_global_size should be as following:

get_global_size(-1) = 1 (dimension:1)
get_global_size(0) = 3 (dimension:1)
get_global_size(1) = 1 (dimension:1)
get_global_size(2) = 1 (dimension:1)

get_global_size(-1) = 1 (dimension:2)
get_global_size(0) = 3 (dimension:2)
get_global_size(1) = 4 (dimension:2)
get_global_size(2) = 1 (dimension:2)
get_global_size(3) = 1 (dimension:2)

get_global_size(-1) = 1 (dimension:3)
get_global_size(0) = 3 (dimension:3)
get_global_size(1) = 4 (dimension:3)
get_global_size(2) = 5 (dimension:3)
get_global_size(3) = 1 (dimension:3)
get_global_size(4) = 1 (dimension:3)

if defined:
  globals[0] = 3;
  globals[1] = 4;
  globals[2] = 5;

v2:
    Handle different version.

    Add #if and #elif to make the test case be suitable to different version.

Signed-off-by: Yi Sun <yi.sun@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoGBE: Fixed GBE: Fix some builtin functions' return value.
Zhigang Gong [Mon, 17 Jun 2013 09:50:35 +0000 (17:50 +0800)]
GBE: Fixed GBE: Fix some builtin functions' return value.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Sun, Yi <yi.sun@intel.com>
11 years agoutests: enable test case for global memory barrier.
Zhigang Gong [Mon, 17 Jun 2013 09:35:44 +0000 (17:35 +0800)]
utests: enable test case for global memory barrier.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
11 years agoAdd memory fence before barrier to support global memory barrier.
Yang Rong [Mon, 17 Jun 2013 07:13:05 +0000 (15:13 +0800)]
Add memory fence before barrier to support global memory barrier.

Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>