add a lost special register name

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

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>

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>

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>

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>

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>

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>

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>

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>

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>

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>

utests: enable test case for global memory barrier.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

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>

add test case for null kernel arg of global/constant buffer

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Refine error check in SetKernelArg() and support NULL buffer argument

1. refine error check in clSetKernelArg() to follow spec.
2. add support NULL buffer as argument, so user could write like below:
    __kernel void func(__global int * p1, constant int* p2) {
      if(p1) {
        //do some thing if p1 is not NULL.
      } else {
        //do other things if p1 is NULL
      }
    }
    Then calling clSetKernelArg(k, 0, sizeof(cl_mem), NULL);
3. as there maybe small possibility that we get a 0 starting address for bo,
   we add an assert() to make sure we do not get a 0 addressed bo.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Refine error check in clCreateProgramFromSource() and clCreateKernel()

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Refine error check in clCreateCommandQueue()

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

utests: Add test case for global memory barrier

Disable this test case currently, as we don't support
global memory barrier now.

Signed-off-by: Dag Lem <dag@nimrod.no>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

utests: Added one test case for clFinish().

Reviewed-by: Zou, Nanhai <nanhai.zou@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

CL: Fix the bug in clfinish.

The previous implementation forget to set the last batch
buffer, so it always does nothing at clFinish(). Now we
move the gpu sync to driver side, and set the last batch
to proper buffer object and make clFinish work as expected.

Reported and tested by: Edward Ching <edward.k.ching@gmail.com>

Signed-off-by: Zou, Nanhai <nanhai.zou@intel.com>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Fix two tests fail when OCL_SIMD_WIDTH=8.

Add barrier for compiler_local_memory and compiler_local_memory_two_ptr,
otherwise tests may fail if work group size bigger than thread's simd size.

After add barrier, the test compiler_local_memory is same as
compiler_local_memory_barrier, so delete test compiler_local_memory.

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

Add mem_fence built-in.

According openCL spec, mem_fence is used to orders loads and stores of a work-item.
This is guarantee by hardware, so add empty mem_fence functions.

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

GBE: Fixed a bug in register expieration.

Some special register which are preallocated at g0 and not
managed by our register allocator. So we need to ignore
them when process the register expieration. Otherwise,
it will trigger an assert latter.

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

Refine error check in clCreateContext()

v2: check is also needed in clCreateContextFromType()

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Refine error check in clGetDeviceIDs()

move error check to api level. correctly handle mixed device type.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Refine error check in clGetPlatformIDs() and clGetPlatformInfo()

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

improve disassembling GPU binary code

support 64-bit float data type
not print prefix zero for unsigned hex immediate constant
display SIMD mode, category, message type in data port message

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Check for exhaustion of local memory

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Pad instruction stream with 8 nops;

   According to Bspec, EU may access 8 instrunctions beyond the
   kernel program because of prefetch.
   Pad the instruction stream with 8 nop to avoid access into an
   invalide page.

Signed-off-by: Zou Nanhai <nanhai.zou@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Adaptions for LLVM 3.3 / SPIR

Handle the fact that several include files are moved from llvm/ to
llvm/IR/ in LLVM 3.3.

"__attribute__((always_inline)) inline" no longer works as intended,
and is replaced by "inline __attribute__((always_inline))".

For LLVM 3.3, the target is changed from "nvptx" to "spir", and
built-in address space qualifiers are used. For now, the built-in
types image2d_t, image3d_t, sampler_t, and event_t are overridden by
defines.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Xing, Homer <homer.xing@intel.com>
Tested-by: Yang, Rong R <rong.r.yang@intel.com>

utests: change all kernels to unix style text.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Xing, Homer <homer.xing@intel.com>
Tested-by: Yang, Rong R <rong.r.yang@intel.com>

utest: Add test case for build-in function get_work_dim

v2: Refine the case, verifying the result of function get_work_dim.

v3: Since the 16 work group size limitation is fixed, re-side the global size and local size with 1.

Signed-off-by: Yi Sun <yi.sun@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

GBE: Add two builtin functions get_work_dim / get_global_offset.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Yi Sun <yi.sun@intel.com>

test case for DW multiplication

This case tests whether a predication bug of DW multiplication
had been fixed.

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Enable built-in sub_sat/add_sat vector type.

also fix a typo in DECL_VECTOR_20P, DECL_VECTOR_30P macro.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

GBE: fixed a predication bug for DW multiplication.

Per bspec:
  mul (8) acc0:d r2.0<8;8,1>:d r3.0<8;8,1>:d //All channels must be enabled
  mach (8) rTemp<1>:d r2.0<8;8,1>:d r3.0<8;8,1>:d //All channels must be enabled
  mov (8) r5.0<1>:d rTemp<8;8,1>:d // High 32 bits
  mov (8) r6.0<1>:d acc0:d // Low 32 bits

  The mul and mach instructions must have all channels enabled.
  The first mov should have channel enable from the destHI of IMUL,
  the second mov should have the channel enable from the destLO of IMUL.

We need to disable the predication and the mask rather than only set noMask to 1.
The strange thing here is for the first quarter, it seems we don't need to do so.
As change both quarter to this style will waste some registers which cause some
kernels fail to compile (compiler_box_blur.cl), I just change the second quater
to fully comply with bspec here. And in practice, it works fine with all unit
test cases and Homer's specific test case.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Xing, Homer <homer.xing@intel.com>

Add test case for group size

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Support non-16 multiple group size.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

GBE: Fixed a 3 elements vector load/store bug.

Per OpenCL spec, for 3-component vector data types,the
size of the data type is 4 * sizeof(component). And llvm
FE really cast a type3 data to type4 data for load/store
instruction, thus break our implementation. We need to
fixup it to the actual element size.

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

Fix a random assert caused by scalarize pass.

Revome the dead values in unit.valueMap at each begin of pass to avoid
the new value have some address.
Also fix a typo

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

utests: Add test case for box blur on image buffer

This test demonstrates box blur on an image buffer, using an extremely
simple kernel.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Correct sampler address clamping for read image functions.

Integer coordinates were converted to unsigned integers, yielding
incorrect clamping of negative coordinates.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

GBE: Fixed a hang issue on 64bit system.

As the npos is a (size_t)-1, we have to change the start and
end to size_t type. Otherwise, the comparison will always
be true, and it will hang there forever.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Boqun Feng <boqun.feng@intel.com>

Pass user options to clang.

clang do not support all options in opencl.
clang unsupport options:
       -cl-denorms-are-zero, -cl-strict-aliasing
       -cl-no-signed-zeros, -cl-fp32-correctly-rounded-divide-sqrt
all support options, refer to clang/include/clang/Driver/Options.inc
Can ignore these options to avoid warning.

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

Fix int div/rem assert in SIMD8 mode.

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

utests: test vector load and store.

Add float4/short4/char4 test case.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

GBE: support load/store of char/short vector.

We just load/store those vector elements one by one.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

Update documents.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

GBE: use the simd width environment variable if set.

If the user set OCL_SIMD_WIDTH to 8 or 16, we force the backend
to use it. Without this patch, the OCL_SIMD_WIDTH never have a
chance to take effect, as we always set the functions simdWidth
before we create a gen context.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

GBE: fixed a bug on simd8 mode for typed_write instruction.

On simd8 mode, the total src length equal to 17 thus, we have
to adjust the max src num here, and adjust the bit width of
the srcNum from 4 to 5. After this fix, all the image/gl
related test case could pass on both simd8 and simd16 mode.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

GBE: Fixed a bug in byte gather/scatter.

We can't just use the alignment to determine whether use
the gather/scatter or not. As for a short type, the compiler
may also generate a 4 alignment, thus it will trigger this
bug.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

Gracefully handle unsupported systems

Gracefully handle cases with unsupported or unreachable GPUs.
Improved error handling in cl_get_device_ids.

Signed-off-by: Mario Kicherer <dev@kicherer.org>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Correct clEnqueueReadBuffer, clEnqueueWriteBuffer and clEnqueueMapBuffer

This implements handling of the offset parameter, and adds sanity
checks according to spec.

A bug is fixed in clEnqueueReadBuffer, where the buffer was not
unmapped after copying.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

utests: Correct box blur

The box blur test kernel incorrectly calculates the bottom extents of
the image. This yields visible blocking artifacts in the top of the
test image (since BMP images are stored bottom to top).

These calculations are corrected, other extent calculations are
simplified, and some dead code is removed.

The reference image is corrected accordingly, and is now identical to the
reference image for the box blur float buffer test.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Change CMake to support LLVM 3.1.

Also fix LLVM 3.1 build errors caused by my vector scalarize commit.

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

Change clang system call to libclang api call.

The original call clang command directly as frontend. The implement is not very flexible.
I change to call libclang apis, support both clang 3.1 and clang 3.2.
Now still write the intermediate to the file, for code simply.

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

Add ICD dispatch table in cloned kernels

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Fix several CL error code return bugs

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Fix a scalarize pass bug.

Need to clear vectors and maps after finish one function.

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

Correct clCreateImage(context, CL_MEM_COPY_HOST_PTR, ...)

The current implementation of clCreateImage initializes images from
host memory using one of the functions cl_mem_copy_data_linear,
cl_mem_copy_data_tilex and cl_mem_copy_data_tiley.

This yields garbled images on some platforms, since the tiled formats
do not always correspond to the formats assumed by the functions
above.

This is fixed by replacing the functions above with the new function
cl_mem_copy_image, whichs maps tiled images in GTT mode for copying.

cl_mem_copy_image also implements missing 3D image copy functionality
(image buffers should also be copied correctly, if/when they are
allowed).

Signed-off-by: Dag Lem <dag@nimrod.no>
Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

utests: extent get_image_size cases to other informations..

Extent it to test all the supported image informations.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Simon Richter <Simon.Richter@hogyros.de>

Add more get image info functions.

Add get image depth/channel data type/channel order/dim support.
Now, only those functions for the unsupported image type have not
been implemented. The unsupported image types are as below:
image1d_t,image1d_buffer_t,image1d_array_t,image2d_array_t.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Simon Richter <Simon.Richter@hogyros.de>

Corrected return of error code in clCreateImage.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Implement clEnqueueReadImage and clEnqueueWriteImage.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Implement clEnqueueMapImage.

For correct mapping and unmapping of tiled/untiled buffers, the new
functions cl_mem_map_auto and cl_mem_unmap_auto are used in
clEnqueueMapBuffer, clEnqueueMapImage, and clEnqueueUnmapMemObject.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

CL: Fixed a get image info bug on 64 bit system.

Forgot to merge it with previous patch. Now commit this one
separately to fix the bug on 64 bit system.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Simon Richter <simon.richter@hogyros.de>

Add CL/*.hpp to installing files

The orignal CMake config only installed CL/*.h. As C++ binding is added,
CL/*.hpp also need to be installing files.

Signed-off-by: Boqun Feng <boqun.feng@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Save depth and slice_pitch in cl_mem images.

Signed-off-by: Dag Lem <dag@nimrod.no>

utests: add one test case to test get_image_width/height.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

CL: complete get_image_width/height support at runtime side.

We get each image's information's curbe slot whcih is
actually used at kernel side, then fill the corresponding
value.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

GBE: Add support for get_image_width/height.

Implement the new GEN IR instruction GetImageInfo. This new instruction
will allocate curbe slot for a given image surface only if it really
referred by a get_image_xxx function. As one image information element
only takes 4 bytes, it's relatively cheap to keep it in the curbe region.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

GBE: preare for get_image_xxx functions support.

Only create the prototypes and the corresponding __gen_ocl_xxx
intrinsics. After this commit, the get_image_width/height could
pass the build, but will not generate code.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Remove useless vector check in GenWriter after scalarize pass.

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

Add a scalarize llvm pass.

In previous implementation expand vector ops in GenWrite,
it is hard to optimize. Now, I add new llvm pass to scalarize.
This pass will expand all normal vector ops to scalar ops,
except store/load, image read/write and function's argument.
Add fake ExtractElement/InsertElement instructions to avoid
dead instruction elimination, and unit valueMap hold the
relationship between these fake instructions and real
load/store instructions.

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

utests: Refine the fill image0 test case to use map gtt.

Now we don't fill the whole image to a const color. we
fill it according to the coords. Then we can use map gtt
to get the result and verify the result easily on cpu
side.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Simon Richter <Simon.Richter@hogyros.de>

CL: Tell the kernel an image bo's tiling mode.

For an image bo allocation, we need to set its tiling mode thus
latter when we use map gtt to map its to a linear address space, the
kernel can do correct mapping. Otherwise, kernel will treat it as
not a tiled suface.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Simon Richter <Simon.Richter@hogyros.de>

Stubs for C++ Bindings

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Updated cl.hpp from khronos.org/registry/cl/api/1.2/cl.hpp

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Add samplerless read image functions for 2D and 3D images.

Signed-off-by: Dag Lem <dag@nimrod.no>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Add missing include of <limits>

For std::numeric_limits<T>, the <limits> header is required.

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Fix a negative number alignment bug in RegisterFilePartitioner allocate.

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

Make libgbm optional without EGL support

If EGL or GBM cannot be found, the EGL support is disabled, and then
neither library is required.

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Update gitignore files

 - Ignore CMake built files
 - Ignore .so files only in the subdirectory
 - Ignore generated .bmp files
 - Ignore generated config headers
 - Ignore generated source for OCL

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Correct type of device properties

 - CL_DEVICE_MAX_PARAMETER_SIZE is of type size_t
 - CL_DEVICE_MAX_WORK_GROUP_SIZE is of type size_t

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Add clGetDeviceInfo(..., CL_BUILT_IN_KERNELS, ...)

Currently, there are no built-in kernels, so this function returns an empty
string.

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Define clamp(value, lower, upper)

The clamp(value, lower, upper) function is part of the standard library.

 - Define the function, using min() and max() on the lower level
 - Remove private definitions from kernels

Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Fix out-of-date math macros

Gen math function "tan" was replaced by "fdiv", and "sincos" was removed.

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Zhigang, Gong <zhigang.gong@linux.intel.com>

Test new math built-in functions

Use random test data.
Test 1000 times.

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Feng, Boqun <boqun.feng@intel.com>

Enable 39 math built-in functions

Enable math built-in functions, including:
  cospi, cbrt, tan, ilogb, nan, sincos, asin, asinpi, acos, acospi,
  atan, atanpi, copysign, erf, erfc, maxmag, minmag, hypot, sinpi,
  log1p, logb, tanpi, sinh, cosh, tanh, asinh, acosh, atanh, remainder,
  rint, fdim, fract, native_divide, pown, rootn, ldexp, exp2, exp10, expm1

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Reviewed-by: Feng, Boqun <boqun.feng@intel.com>

Fix instruction scheduler ScheduleDAGNode

When we add dependency for ScheduleDAGNode, we use
ScheduleDAGNode(node1)->children.push_back(node0) to express
node0 depends on node1.

To check this dependency, we should check whether node0 is inside
node1->children.

Original code checks whether node1 is inside node0->children.

So we fix this.

Signed-off-by: Homer Hsing <homer.xing@intel.com>
Tested-by: Zhigang Gong <zhigang.gong@linux.intel.com>

Rename CBMove to IndirectMove

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

utests: Fix a bug in movforphi test case.

This test case is to trigger a old MovForPHI bug, and it
just use read/write_image. But it has a bug in itself.
As in the kernel, the write image will only write the
first lane data not all the 16 lanes. As the previous
patch fix the write_image bug, thus now the write_image
work correctly and thus it only touch the first data element
thus it trigger the bug in this test case. Now fix it.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

GBE: fixed a prediction bug in typed write instruction.

We need to put the header initialization and the LOD initialization
to no mask and no predication state. For all the other parts, we need
to enable mask and use the current predication state and need to set
the quater control properly. Otherwise, when write_image is called in
a condition-branch block, it will trigger this bug and doesn't write
data out correctly.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>

GBE/Runtime: Optimize Sample/TypedWrite instruction.

This commit does two major things as below:
1. Allocate image surface at compile time, and add new gbe interfaces to let runtime know
how many image surfaces we have, and the image allocation informations. Thus the runtime
library know how to bind those image surfaces.

2. As now for both image and sampler, at compile time, we know the eaxct binding table
index. We no longer need to get those index from the input argument(curbe) and prepare
the desc to the architecture register. We can use imm as the desc thus we can save
4 out of 4 instructions for SampleInstruction and save 2 out of 12 instructions for
the TypedWriteInstruction.

This patch is also a major prepartion for the get_image_width/height/... functions.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>

GBE: concentrate all samplers' allocation at compile time.

This is the first step to do image/sampler allocation fully
at compile time. Thus we can determine all the sampler id and image
bti index at compile time. So it can make the following things
easier or faster:

1. After we finish both image/sampler, we can treat all image bti and sampler
   as constant and can get their value when we encode the Sampler and TypedWrite
   instructions. Then we don't need to compute the message header at runtime which
   cost 3 instructions each call.

2. get image width/height/depth. As we know the surface bti at compile time,
   we can put those data at specified curbe entry and generate correct indirect
   register access to get those information at compile time.

This is the first step. And just finish the sampler part. Now all the
samplers including those defeined in kernel arguments will be allocated
at compile time. At runtime, it just need to fill in the sampler value
into the proper slot which map to the specified input argument. Then the
driver will create and bind the sampler to the correct slot.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>

GBE: add scalar register support in loadImmInstruction.

There is a slight possibility that the destination register
is a scalar register. We need to check it here.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>

GBE: remove sampler address space.

As now sampler_t is a normal integer data type, we don't
need the sampler address space any more.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>

utests: Add one test cases for sampler support.

This new case tests define sampler in kernel side and in the
kernel argument.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>

CL: Support kernel side defined samplers.

We changed the way to handle samplers. We gather all the kernel side
defined samplers and those sampler in kernel argument into one samplers
array. And don't allocate one single sampler each time.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>

GBE: refine the sampler implementation to comply with spec.

The previous implementation is to use a new address space pointer to
represent a sampler. The reason is that there is no specified data
type for sampler_t in LLVM front end thus we can't determine the
sampler argument type if we use a normal interger to represnet the
sampler. But that breaks the OCL spec, the spec allows the kernel
to define and initialize sampler variables in kernel side.

Now I use a little tricky way to fix this problem. First, I decide
to use normal unsigned interger to represent sampler_t in kernel side.
Then at compile time, I check read_imagexxx function's sampler
arguments. If the argument is a constant value, then it should be a
kernel side defined sampler, then I insert the sampler type into a
global sampler set for the current kernel function. If the argument
is not a constant value, then I will check whether it's a kernel
argument, if it is, then I fix up the corresponding kernel arg type
to SAMPLER there.

To unify the kernel side defined sampler and kernel argument sampler,
I add two new gbe API. To export all the kernel side defined sampler
data and size to the runtime library. Then latter, the runtime library
can use this information to append new sampler to the unified sampler
buffer and bind all the sampler at one time.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Lv, Meng <meng.lv@intel.com>