[Pkg-opencl-devel] [beignet] 35/66: Imported Debian patch 0.1+git20130625+97c3a9b-1
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 07:27:05 UTC 2014
This is an automated email from the git hooks/post-receive script.
anbe pushed a commit to branch master
in repository beignet.
commit e06914c4ba508466a2674589a499dcad34ecf36d
Merge: 123ba56 b961e6b
Author: Simon Richter <sjr at debian.org>
Date: Tue Jun 25 13:52:00 2013 +0200
Imported Debian patch 0.1+git20130625+97c3a9b-1
README.html | 134 -----
README.md | 140 +----
backend/README.html | 52 --
backend/doc/TODO.html | 94 ---
backend/doc/compiler_backend.html | 107 ----
backend/doc/flat_address_space.html | 93 ---
backend/doc/gen_ir.html | 248 --------
backend/doc/unstructured_branches.html | 280 ---------
backend/src/backend/gen_encoder.cpp | 5 +-
backend/src/backend/gen_insn_selection.cpp | 4 +-
debian/changelog | 6 +
.../0001-Add-vector-argument-test-case.patch | 35 +-
...d-OpenCL-1.2-definitions-required-for-ICD.patch | 95 ++++
...more-support-of-char-and-short-arithmetic.patch | 139 +++++
...004-utests-Add-basic-arithmetic-test-case.patch | 376 ++++++++++++
...ltin-function-abs-and-the-according-test-.patch | 218 +++++++
...-PATCH-Refine-the-get_local_id-.-builtins.patch | 55 ++
...-support-of-the-API-clGetCommandQueueInfo.patch | 185 ++++++
...e-test-case-for-clGetCommandQueueInfo-API.patch | 631 +++++++++++++++++++++
debian/patches/series | 6 +
README.md => docs/Beignet.mdwn | 8 +-
backend/README.md => docs/Beignet/Backend.mdwn | 13 +-
.../doc/TODO.md => docs/Beignet/Backend/TODO.mdwn | 9 +-
.../Beignet/Backend/compiler_backend.mdwn | 4 +-
.../Beignet/Backend/flat_address_space.mdwn | 3 -
.../gen_ir.md => docs/Beignet/Backend/gen_ir.mdwn | 4 +-
.../Beignet/Backend/unstructured_branches.mdwn | 5 +-
kernels/compiler_local_memory_barrier_2.cl | 2 +-
utests/CMakeLists.txt | 2 +-
utests/compiler_local_memory_barrier_2.cpp | 4 +-
30 files changed, 1759 insertions(+), 1198 deletions(-)
diff --cc debian/changelog
index 38fba13,0000000..f7f7c49
mode 100644,000000..100644
--- a/debian/changelog
+++ b/debian/changelog
@@@ -1,133 -1,0 +1,139 @@@
++beignet (0.1+git20130625+97c3a9b-1) unstable; urgency=low
++
++ * New upstream release
++
++ -- Simon Richter <sjr at debian.org> Tue, 25 Jun 2013 13:52:00 +0200
++
+beignet (0.1+git20130621+30586bf-1) unstable; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Fri, 21 Jun 2013 12:08:49 +0200
+
+beignet (0.1+git20130619+42967d2-2) unstable; urgency=low
+
+ * Drop EGL support for kFreeBSD for the time being
+
+ -- Simon Richter <sjr at debian.org> Thu, 20 Jun 2013 11:18:59 +0200
+
+beignet (0.1+git20130619+42967d2-1) unstable; urgency=low
+
+ * New upstream release
+ * Build against Mesa 9
+ * Enable GL sharing extension
+
+ -- Simon Richter <sjr at debian.org> Wed, 19 Jun 2013 20:48:03 +0200
+
+beignet (0.1+git20130614+89b5e40-2) unstable; urgency=low
+
+ * Add Ubuntu support
+ * Upload to unstable to get an ICD capable package there
+
+ -- Simon Richter <sjr at debian.org> Fri, 14 Jun 2013 17:40:45 +0200
+
+beignet (0.1+git20130614+89b5e40-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Fri, 14 Jun 2013 15:22:18 +0200
+
+beignet (0.1+git20130521+a7ea35c-1) experimental; urgency=low
+
+ * Rename binary package
+
+ -- Simon Richter <sjr at debian.org> Tue, 21 May 2013 10:48:39 +0200
+
+beignet (0.1+git20130521+a7ea35c-1~prerename) experimental; urgency=low
+
+ * New upstream release
+ * Move libraries to /usr/lib/beignet (should not be used directly)
+
+ -- Simon Richter <sjr at debian.org> Tue, 21 May 2013 09:17:45 +0200
+
+beignet (0.1+git20130514+19e9c58-1) experimental; urgency=low
+
+ * New upstream release
+ * Added a number of tentative patches
+
+ -- Simon Richter <sjr at debian.org> Tue, 14 May 2013 20:04:29 +0200
+
+beignet (0.1+git20130502+63e60ed-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Mon, 06 May 2013 06:30:32 +0200
+
+beignet (0.1+git20130426+0c8f6fe-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Fri, 26 Apr 2013 14:42:21 +0200
+
+beignet (0.1+git20130422+003fac5-2) experimental; urgency=low
+
+ * Add patch for select()
+ * Add patch for fmin() / fmax()
+
+ -- Simon Richter <sjr at debian.org> Mon, 22 Apr 2013 18:26:01 +0200
+
+beignet (0.1+git20130422+003fac5-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Mon, 22 Apr 2013 15:10:54 +0200
+
+beignet (0.1+git20130419+9c11c18-1) experimental; urgency=low
+
+ * Add more functionality patches
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Fri, 19 Apr 2013 14:14:39 +0200
+
+beignet (0.1+git20130418+0546d2e-2) experimental; urgency=low
+
+ * Add functionality patches
+ * Use clang 3.0 command line syntax
+
+ -- Simon Richter <sjr at debian.org> Fri, 19 Apr 2013 09:53:23 +0200
+
+beignet (0.1+git20130418+0546d2e-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Thu, 18 Apr 2013 11:51:37 +0200
+
+beignet (0.1-1) unstable; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Tue, 16 Apr 2013 17:16:18 +0200
+
+beignet (0.0.0+git2013.04.11+e6b503e-1) unstable; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Mon, 15 Apr 2013 18:22:45 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-4) unstable; urgency=low
+
+ * Build fix for kfreebsd-*
+
+ -- Simon Richter <sjr at debian.org> Fri, 12 Apr 2013 11:22:36 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-3) unstable; urgency=low
+
+ * Adjust Build-Depends, Architecture list
+
+ -- Simon Richter <sjr at debian.org> Fri, 12 Apr 2013 10:32:36 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-2) unstable; urgency=low
+
+ * Add patch to support size queries in device info
+
+ -- Simon Richter <sjr at debian.org> Thu, 11 Apr 2013 14:00:59 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-1) unstable; urgency=low
+
+ * Initial release.
+
+ -- Simon Richter <sjr at debian.org> Tue, 09 Apr 2013 17:14:00 +0200
diff --cc debian/patches/0001-Add-vector-argument-test-case.patch
index 34a125f,0000000..eac26c2
mode 100644,000000..100644
--- a/debian/patches/0001-Add-vector-argument-test-case.patch
+++ b/debian/patches/0001-Add-vector-argument-test-case.patch
@@@ -1,69 -1,0 +1,74 @@@
- From ee47f1b7f325f6e8b1c54c81a16f2480c968513b Mon Sep 17 00:00:00 2001
++From 0ee7f97ae6e740ea80766d1126e0520583780d40 Mon Sep 17 00:00:00 2001
+From: Yang Rong <rong.r.yang at intel.com>
+Date: Thu, 16 May 2013 12:36:35 +0800
- Subject: [PATCH 1/2] Add vector argument test case.
++Subject: [PATCH 1/8] Add vector argument test case.
+To: beignet at lists.freedesktop.org
+
+Signed-off-by: Yang Rong <rong.r.yang at intel.com>
+---
+ kernels/compiler_function_argument2.cl | 6 ++++++
+ utests/CMakeLists.txt | 1 +
+ utests/compiler_function_argument2.cpp | 26 ++++++++++++++++++++++++++
+ 3 files changed, 33 insertions(+)
+ create mode 100644 kernels/compiler_function_argument2.cl
+ create mode 100644 utests/compiler_function_argument2.cpp
+
- Index: beignet-0.1+git20130621+30586bf/kernels/compiler_function_argument2.cl
- ===================================================================
- --- /dev/null 1970-01-01 00:00:00.000000000 +0000
- +++ beignet-0.1+git20130621+30586bf/kernels/compiler_function_argument2.cl 2013-06-21 12:09:33.002557680 +0200
++diff --git a/kernels/compiler_function_argument2.cl b/kernels/compiler_function_argument2.cl
++new file mode 100644
++index 0000000..0985dbd
++--- /dev/null
+++++ b/kernels/compiler_function_argument2.cl
+@@ -0,0 +1,6 @@
++__kernel void
++compiler_function_argument2(__global int *dst, int4 value)
++{
++ int id = (int)get_global_id(0);
++ dst[id] = value.w;
++}
- Index: beignet-0.1+git20130621+30586bf/utests/CMakeLists.txt
- ===================================================================
- --- beignet-0.1+git20130621+30586bf.orig/utests/CMakeLists.txt 2013-06-21 10:13:03.000000000 +0200
- +++ beignet-0.1+git20130621+30586bf/utests/CMakeLists.txt 2013-06-21 12:09:33.002557680 +0200
- @@ -37,6 +37,7 @@
++diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
++index c009d99..ea23f31 100644
++--- a/utests/CMakeLists.txt
+++++ b/utests/CMakeLists.txt
++@@ -37,6 +37,7 @@ set (utests_sources
+ compiler_fill_image_3d_2.cpp
+ compiler_function_argument0.cpp
+ compiler_function_argument1.cpp
++ compiler_function_argument2.cpp
+ compiler_function_argument.cpp
+ compiler_function_constant0.cpp
+ compiler_function_constant1.cpp
- Index: beignet-0.1+git20130621+30586bf/utests/compiler_function_argument2.cpp
- ===================================================================
- --- /dev/null 1970-01-01 00:00:00.000000000 +0000
- +++ beignet-0.1+git20130621+30586bf/utests/compiler_function_argument2.cpp 2013-06-21 12:09:33.002557680 +0200
++diff --git a/utests/compiler_function_argument2.cpp b/utests/compiler_function_argument2.cpp
++new file mode 100644
++index 0000000..1e398a9
++--- /dev/null
+++++ b/utests/compiler_function_argument2.cpp
+@@ -0,0 +1,26 @@
++#include "utest_helper.hpp"
++
++struct int4 {int x,y,z,w;};
++void compiler_function_argument2(void)
++{
++ const size_t n = 2048;
++ const int4 value = {31, 32, 33, 34};
++
++ // Setup kernel and buffers
++ OCL_CREATE_KERNEL("compiler_function_argument2");
++ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
++ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
++ OCL_SET_ARG(1, sizeof(int4), &value);
++
++ // Run the kernel
++ globals[0] = n;
++ locals[0] = 16;
++ OCL_NDRANGE(1);
++ OCL_MAP_BUFFER(0);
++
++ // Check results
++ for (uint32_t i = 0; i < n; ++i)
++ OCL_ASSERT(((int*)buf_data[0])[i] == value.w);
++}
++
++MAKE_UTEST_FROM_FUNCTION(compiler_function_argument2);
++--
++1.7.10.4
++
diff --cc debian/patches/0002-Readd-OpenCL-1.2-definitions-required-for-ICD.patch
index 0000000,0000000..c385dcd
new file mode 100644
--- /dev/null
+++ b/debian/patches/0002-Readd-OpenCL-1.2-definitions-required-for-ICD.patch
@@@ -1,0 -1,0 +1,95 @@@
++From 3d63c833d4ebcafe8e46fd498c28d08cb3046cd8 Mon Sep 17 00:00:00 2001
++From: Simon Richter <Simon.Richter at hogyros.de>
++Date: Wed, 19 Jun 2013 11:30:36 +0200
++Subject: [PATCH 2/8] Readd OpenCL 1.2 definitions required for ICD
++To: beignet at lists.freedesktop.org
++
++The definition for the ICD dispatch table requires a few additional
++definitions from OpenCL 1.2.
++---
++ include/CL/cl.h | 15 +++++++++++++++
++ include/CL/cl_platform.h | 2 ++
++ src/cl_mem.h | 12 ------------
++ 3 files changed, 17 insertions(+), 12 deletions(-)
++
++diff --git a/include/CL/cl.h b/include/CL/cl.h
++index 4355e74..a7f25d1 100644
++--- a/include/CL/cl.h
+++++ b/include/CL/cl.h
++@@ -67,6 +67,7 @@ typedef cl_uint cl_channel_type;
++ typedef cl_bitfield cl_mem_flags;
++ typedef cl_uint cl_mem_object_type;
++ typedef cl_uint cl_mem_info;
+++typedef cl_bitfield cl_mem_migration_flags;
++ typedef cl_uint cl_image_info;
++ typedef cl_uint cl_buffer_create_type;
++ typedef cl_uint cl_addressing_mode;
++@@ -75,8 +76,10 @@ typedef cl_uint cl_sampler_info;
++ typedef cl_bitfield cl_map_flags;
++ typedef cl_uint cl_program_info;
++ typedef cl_uint cl_program_build_info;
+++typedef intptr_t cl_device_partition_property;
++ typedef cl_int cl_build_status;
++ typedef cl_uint cl_kernel_info;
+++typedef cl_uint cl_kernel_arg_info;
++ typedef cl_uint cl_kernel_work_group_info;
++ typedef cl_uint cl_event_info;
++ typedef cl_uint cl_command_type;
++@@ -87,6 +90,18 @@ typedef struct _cl_image_format {
++ cl_channel_type image_channel_data_type;
++ } cl_image_format;
++
+++typedef struct _cl_image_desc {
+++ cl_mem_object_type image_type;
+++ size_t image_width;
+++ size_t image_height;
+++ size_t image_depth;
+++ size_t image_array_size;
+++ size_t image_row_pitch;
+++ size_t image_slice_pitch;
+++ cl_uint num_mip_levels;
+++ cl_uint num_samples;
+++ cl_mem buffer;
+++} cl_image_desc;
++
++ typedef struct _cl_buffer_region {
++ size_t origin;
++diff --git a/include/CL/cl_platform.h b/include/CL/cl_platform.h
++index 043b048..9a2f17a 100644
++--- a/include/CL/cl_platform.h
+++++ b/include/CL/cl_platform.h
++@@ -58,6 +58,8 @@ extern "C" {
++ #define CL_EXT_SUFFIX__VERSION_1_0
++ #define CL_API_SUFFIX__VERSION_1_1
++ #define CL_EXT_SUFFIX__VERSION_1_1
+++ #define CL_API_SUFFIX__VERSION_1_2
+++ #define CL_EXT_SUFFIX__VERSION_1_2
++ #define CL_EXT_SUFFIX__VERSION_1_0_DEPRECATED
++ #endif
++
++diff --git a/src/cl_mem.h b/src/cl_mem.h
++index 33ad174..66815fe 100644
++--- a/src/cl_mem.h
+++++ b/src/cl_mem.h
++@@ -29,18 +29,6 @@
++ #define CL_MEM_OBJECT_IMAGE1D_ARRAY 0x10F5
++ #define CL_MEM_OBJECT_IMAGE1D_BUFFER 0x10F6
++ #define CL_MEM_OBJECT_IMAGE2D_ARRAY 0x10F3
++-typedef struct _cl_image_desc {
++- cl_mem_object_type image_type;
++- size_t image_width;
++- size_t image_height;
++- size_t image_depth;
++- size_t image_array_size;
++- size_t image_row_pitch;
++- size_t image_slice_pitch;
++- cl_uint num_mip_levels;
++- cl_uint num_samples;
++- cl_mem buffer;
++-} cl_image_desc;
++ #endif
++
++ typedef enum cl_image_tiling {
++--
++1.7.10.4
++
diff --cc debian/patches/0003-GBE-Add-more-support-of-char-and-short-arithmetic.patch
index 0000000,0000000..22a0eec
new file mode 100644
--- /dev/null
+++ b/debian/patches/0003-GBE-Add-more-support-of-char-and-short-arithmetic.patch
@@@ -1,0 -1,0 +1,139 @@@
++From 1525f53083d7623659e51a9d6f1e4835a83a6caf Mon Sep 17 00:00:00 2001
++From: Ruiling Song <ruiling.song at intel.com>
++Date: Tue, 25 Jun 2013 15:38:48 +0800
++Subject: [PATCH 3/8] GBE: Add more support of char and short arithmetic
++To: beignet at lists.freedesktop.org
++
++add * / % support of char and short type.
++
++Signed-off-by: Ruiling Song <ruiling.song at intel.com>
++---
++ backend/src/backend/gen_insn_selection.cpp | 67 +++++++++++++++++++++++-----
++ backend/src/llvm/llvm_gen_backend.cpp | 4 +-
++ 2 files changed, 57 insertions(+), 14 deletions(-)
++
++diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
++index 1e5f514..b1c6093 100644
++--- a/backend/src/backend/gen_insn_selection.cpp
+++++ b/backend/src/backend/gen_insn_selection.cpp
++@@ -1260,30 +1260,73 @@ namespace gbe
++ const Opcode opcode = insn.getOpcode();
++ const Type type = insn.getType();
++ GenRegister dst = sel.selReg(insn.getDst(0), type);
+++ const uint32_t simdWidth = sel.curr.execWidth;
+++ const RegisterFamily family = getFamily(type);
+++
+++ //bytes and shorts must be converted to int for DIV and REM per GEN restriction
+++ if((opcode == OP_DIV || opcode == OP_REM)
+++ && (family == FAMILY_WORD || family == FAMILY_BYTE)) {
+++ GenRegister src0 = sel.selReg(insn.getSrc(0), type);
+++ GenRegister src1 = sel.selReg(insn.getSrc(1), type);
+++ uint32_t function = (opcode == OP_DIV)?
+++ GEN_MATH_FUNCTION_INT_DIV_QUOTIENT :
+++ GEN_MATH_FUNCTION_INT_DIV_REMAINDER;
+++ GenRegister tmp0 = src0;
+++ GenRegister tmp1 = src1;
+++ GenRegister tmp2 = dst;
+++ tmp0 = GenRegister::udxgrf(simdWidth, sel.reg(FAMILY_DWORD));
+++ tmp0 = GenRegister::retype(tmp0, GEN_TYPE_D);
+++ sel.MOV(tmp0, src0);
+++
+++ tmp1 = GenRegister::udxgrf(simdWidth, sel.reg(FAMILY_DWORD));
+++ tmp1 = GenRegister::retype(tmp1, GEN_TYPE_D);
+++ sel.MOV(tmp1, src1);
+++
+++ tmp2 = GenRegister::udxgrf(simdWidth, sel.reg(FAMILY_DWORD));
+++ tmp2 = GenRegister::retype(tmp2, GEN_TYPE_D);
+++
+++ sel.MATH(tmp2, function, tmp0, tmp1);
+++ GenRegister unpacked;
+++ if(family == FAMILY_WORD) {
+++ unpacked = GenRegister::unpacked_uw(sel.reg(FAMILY_DWORD));
+++ } else {
+++ unpacked = GenRegister::unpacked_ub(sel.reg(FAMILY_DWORD));
+++ }
+++ unpacked = GenRegister::retype(unpacked, getGenType(type));
+++ sel.MOV(unpacked, tmp2);
+++ sel.MOV(dst, unpacked);
++
+++ markAllChildren(dag);
+++ return true;
+++ }
++ // Immediates not supported
++ if (opcode == OP_DIV || opcode == OP_POW) {
++ GenRegister src0 = sel.selReg(insn.getSrc(0), type);
++ GenRegister src1 = sel.selReg(insn.getSrc(1), type);
++ uint32_t function;
++- if (type == TYPE_S32 || type == TYPE_U32)
+++ if (type == TYPE_S32 || type == TYPE_U32 ) {
++ function = GEN_MATH_FUNCTION_INT_DIV_QUOTIENT;
++- else
+++ sel.MATH(dst, function, src0, src1);
+++ } else if(type == TYPE_FLOAT) {
++ function = opcode == OP_DIV ?
++ GEN_MATH_FUNCTION_FDIV :
++ GEN_MATH_FUNCTION_POW;
++- sel.MATH(dst, function, src0, src1);
+++ sel.MATH(dst, function, src0, src1);
+++ } else {
+++ NOT_IMPLEMENTED;
+++ }
++ markAllChildren(dag);
++ return true;
++ }
++ if (opcode == OP_REM) {
++ GenRegister src0 = sel.selReg(insn.getSrc(0), type);
++ GenRegister src1 = sel.selReg(insn.getSrc(1), type);
++- if (type == TYPE_U32 || type == TYPE_S32) {
+++ if(type == TYPE_S32 || type == TYPE_U32) {
++ sel.MATH(dst, GEN_MATH_FUNCTION_INT_DIV_REMAINDER, src0, src1);
++- markAllChildren(dag);
++- } else
++- NOT_IMPLEMENTED;
+++ } else {
+++ GBE_ASSERTM(0, "Unsupported type in remainder operation!");
+++ }
+++ markAllChildren(dag);
++ return true;
++ }
++
++@@ -1345,14 +1388,14 @@ namespace gbe
++ case OP_SHR: sel.SHR(dst, src0, src1); break;
++ case OP_ASR: sel.ASR(dst, src0, src1); break;
++ case OP_MUL:
++- if (type == TYPE_FLOAT || type == TYPE_DOUBLE)
++- sel.MUL(dst, src0, src1);
++- else if (type == TYPE_U32 || type == TYPE_S32) {
+++ if (type == TYPE_U32 || type == TYPE_S32) {
++ sel.pop();
++ return false;
++ }
++- else
++- NOT_IMPLEMENTED;
+++ else {
+++ GBE_ASSERTM((type != TYPE_S64 && type != TYPE_U64), "64bit integer not supported yet!" );
+++ sel.MUL(dst, src0, src1);
+++ }
++ break;
++ default: NOT_IMPLEMENTED;
++ }
++diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
++index 5b7754c..b0e8c6c 100644
++--- a/backend/src/llvm/llvm_gen_backend.cpp
+++++ b/backend/src/llvm/llvm_gen_backend.cpp
++@@ -1276,10 +1276,10 @@ namespace gbe
++ case Instruction::FSub: ctx.SUB(type, dst, src0, src1); break;
++ case Instruction::Mul:
++ case Instruction::FMul: ctx.MUL(type, dst, src0, src1); break;
++- case Instruction::URem:
+++ case Instruction::URem: ctx.REM(getUnsignedType(ctx, I.getType()), dst, src0, src1); break;
++ case Instruction::SRem:
++ case Instruction::FRem: ctx.REM(type, dst, src0, src1); break;
++- case Instruction::UDiv:
+++ case Instruction::UDiv: ctx.DIV(getUnsignedType(ctx, I.getType()), dst, src0, src1); break;
++ case Instruction::SDiv:
++ case Instruction::FDiv: ctx.DIV(type, dst, src0, src1); break;
++ case Instruction::And: ctx.AND(type, dst, src0, src1); break;
++--
++1.7.10.4
++
diff --cc debian/patches/0004-utests-Add-basic-arithmetic-test-case.patch
index 0000000,0000000..a3ae637
new file mode 100644
--- /dev/null
+++ b/debian/patches/0004-utests-Add-basic-arithmetic-test-case.patch
@@@ -1,0 -1,0 +1,376 @@@
++From c89dbb34332c104df22c8ea8c22bac0bcb0b5221 Mon Sep 17 00:00:00 2001
++From: Ruiling Song <ruiling.song at intel.com>
++Date: Tue, 25 Jun 2013 15:38:49 +0800
++Subject: [PATCH 4/8] utests: Add basic arithmetic test case
++To: beignet at lists.freedesktop.org
++
++test case for + - * / % of data type (u)int8/16/32
++remove duplicated cases.
++
++Signed-off-by: Ruiling Song <ruiling.song at intel.com>
++---
++ kernels/compiler_basic_arithmetic.cl | 73 +++++++++++++++++++
++ kernels/compiler_sub_bytes.cl | 7 --
++ kernels/compiler_sub_shorts.cl | 7 --
++ utests/CMakeLists.txt | 3 +-
++ utests/compiler_basic_arithmetic.cpp | 132 ++++++++++++++++++++++++++++++++++
++ utests/compiler_sub_bytes.cpp | 35 ---------
++ utests/compiler_sub_shorts.cpp | 36 ----------
++ 7 files changed, 206 insertions(+), 87 deletions(-)
++ create mode 100644 kernels/compiler_basic_arithmetic.cl
++ delete mode 100644 kernels/compiler_sub_bytes.cl
++ delete mode 100644 kernels/compiler_sub_shorts.cl
++ create mode 100644 utests/compiler_basic_arithmetic.cpp
++ delete mode 100644 utests/compiler_sub_bytes.cpp
++ delete mode 100644 utests/compiler_sub_shorts.cpp
++
++diff --git a/kernels/compiler_basic_arithmetic.cl b/kernels/compiler_basic_arithmetic.cl
++new file mode 100644
++index 0000000..2bc2c27
++--- /dev/null
+++++ b/kernels/compiler_basic_arithmetic.cl
++@@ -0,0 +1,73 @@
+++#define DECL_KERNEL_SUB(type)\
+++__kernel void \
+++compiler_sub_##type(__global type *src0, __global type *src1, __global type *dst) \
+++{ \
+++ int id = (int)get_global_id(0); \
+++ dst[id] = src0[id] - src1[id]; \
+++}
+++
+++#define DECL_KERNEL_ADD(type)\
+++__kernel void \
+++compiler_add_##type(__global type *src0, __global type *src1, __global type *dst) \
+++{ \
+++ int id = (int)get_global_id(0); \
+++ dst[id] = src0[id] + src1[id]; \
+++}
+++
+++#define DECL_KERNEL_MUL(type)\
+++__kernel void \
+++compiler_mul_##type(__global type *src0, __global type *src1, __global type *dst) \
+++{ \
+++ int id = (int)get_global_id(0); \
+++ dst[id] = src0[id] * src1[id]; \
+++}
+++
+++#define DECL_KERNEL_DIV(type)\
+++__kernel void \
+++compiler_div_##type(__global type *src0, __global type *src1, __global type *dst) \
+++{ \
+++ int id = (int)get_global_id(0); \
+++ dst[id] = src0[id] / src1[id]; \
+++}
+++
+++#define DECL_KERNEL_REM(type)\
+++__kernel void \
+++compiler_rem_##type(__global type *src0, __global type *src1, __global type *dst) \
+++{ \
+++ int id = (int)get_global_id(0); \
+++ dst[id] = src0[id] % src1[id]; \
+++}
+++DECL_KERNEL_SUB(char)
+++DECL_KERNEL_SUB(uchar)
+++DECL_KERNEL_SUB(short)
+++DECL_KERNEL_SUB(ushort)
+++DECL_KERNEL_SUB(int)
+++DECL_KERNEL_SUB(uint)
+++
+++DECL_KERNEL_ADD(char)
+++DECL_KERNEL_ADD(uchar)
+++DECL_KERNEL_ADD(short)
+++DECL_KERNEL_ADD(ushort)
+++DECL_KERNEL_ADD(int)
+++DECL_KERNEL_ADD(uint)
+++
+++DECL_KERNEL_MUL(char)
+++DECL_KERNEL_MUL(uchar)
+++DECL_KERNEL_MUL(short)
+++DECL_KERNEL_MUL(ushort)
+++DECL_KERNEL_MUL(int)
+++DECL_KERNEL_MUL(uint)
+++
+++DECL_KERNEL_DIV(char)
+++DECL_KERNEL_DIV(uchar)
+++DECL_KERNEL_DIV(short)
+++DECL_KERNEL_DIV(ushort)
+++DECL_KERNEL_DIV(int)
+++DECL_KERNEL_DIV(uint)
+++
+++DECL_KERNEL_REM(char)
+++DECL_KERNEL_REM(uchar)
+++DECL_KERNEL_REM(short)
+++DECL_KERNEL_REM(ushort)
+++DECL_KERNEL_REM(int)
+++DECL_KERNEL_REM(uint)
++diff --git a/kernels/compiler_sub_bytes.cl b/kernels/compiler_sub_bytes.cl
++deleted file mode 100644
++index f058561..0000000
++--- a/kernels/compiler_sub_bytes.cl
+++++ /dev/null
++@@ -1,7 +0,0 @@
++-__kernel void
++-compiler_sub_bytes(__global char *src0, __global char *src1, __global char *dst)
++-{
++- int id = (int)get_global_id(0);
++- dst[id] = src0[id] - src1[id];
++-}
++-
++diff --git a/kernels/compiler_sub_shorts.cl b/kernels/compiler_sub_shorts.cl
++deleted file mode 100644
++index d26de7f..0000000
++--- a/kernels/compiler_sub_shorts.cl
+++++ /dev/null
++@@ -1,7 +0,0 @@
++-__kernel void
++-compiler_sub_shorts(__global short *src0, __global short *src1, __global short *dst)
++-{
++- int id = (int)get_global_id(0);
++- dst[id] = src0[id] - src1[id];
++-}
++-
++diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
++index ea23f31..edfbda1 100644
++--- a/utests/CMakeLists.txt
+++++ b/utests/CMakeLists.txt
++@@ -5,6 +5,7 @@ link_directories (${LLVM_LIBRARY_DIR})
++ set (utests_sources
++ cl_create_kernel.cpp
++ utest_error.c
+++ compiler_basic_arithmetic.cpp
++ compiler_displacement_map_element.cpp
++ compiler_shader_toy.cpp
++ compiler_mandelbrot.cpp
++@@ -56,8 +57,6 @@ set (utests_sources
++ compiler_saturate_sub.cpp
++ compiler_shift_right.cpp
++ compiler_short_scatter.cpp
++- compiler_sub_bytes.cpp
++- compiler_sub_shorts.cpp
++ compiler_uint2_copy.cpp
++ compiler_uint3_copy.cpp
++ compiler_uint8_copy.cpp
++diff --git a/utests/compiler_basic_arithmetic.cpp b/utests/compiler_basic_arithmetic.cpp
++new file mode 100644
++index 0000000..5ab5f44
++--- /dev/null
+++++ b/utests/compiler_basic_arithmetic.cpp
++@@ -0,0 +1,132 @@
+++#include "utest_helper.hpp"
+++
+++enum eTestOP {
+++ TEST_OP_ADD =0,
+++ TEST_OP_SUB,
+++ TEST_OP_MUL,
+++ TEST_OP_DIV,
+++ TEST_OP_REM
+++};
+++
+++template <typename T, eTestOP op>
+++static void test_exec(const char* kernel_name)
+++{
+++ const size_t n = 160;
+++
+++ // Setup kernel and buffers
+++ OCL_CREATE_KERNEL_FROM_FILE("compiler_basic_arithmetic", kernel_name);
+++std::cout <<"kernel name: " << kernel_name << std::endl;
+++ buf_data[0] = (T*) malloc(sizeof(T) * n);
+++ buf_data[1] = (T*) malloc(sizeof(T) * n);
+++ for (uint32_t i = 0; i < n; ++i) ((T*)buf_data[0])[i] = (T) rand();
+++ for (uint32_t i = 0; i < n; ++i) ((T*)buf_data[1])[i] = (T) rand();
+++ if(op == TEST_OP_DIV || op == TEST_OP_REM) {
+++ for (uint32_t i = 0; i < n; ++i) {
+++ if(((T*)buf_data[1])[i] == 0)
+++ ((T*)buf_data[1])[i] = (T) 1;
+++ }
+++ }
+++ OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[0]);
+++ OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(T), buf_data[1]);
+++ OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(T), NULL);
+++
+++ // Run the kernel
+++ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+++ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+++ OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
+++ globals[0] = n;
+++ locals[0] = 16;
+++ OCL_NDRANGE(1);
+++
+++ // Check result
+++ OCL_MAP_BUFFER(2);
+++ if(op == TEST_OP_SUB) {
+++ for (uint32_t i = 0; i < n; ++i)
+++ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] - ((T*)buf_data[1])[i]));
+++ } else if(op == TEST_OP_ADD) {
+++ for (uint32_t i = 0; i < n; ++i)
+++ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] + ((T*)buf_data[1])[i]));
+++ } else if(op == TEST_OP_MUL) {
+++ for (uint32_t i = 0; i < n; ++i)
+++ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] * ((T*)buf_data[1])[i]));
+++ } else if(op == TEST_OP_DIV) {
+++ for (uint32_t i = 0; i < n; ++i)
+++ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] / ((T*)buf_data[1])[i]));
+++ } else {
+++ for (uint32_t i = 0; i < n; ++i)
+++ OCL_ASSERT(((T*)buf_data[2])[i] == (T)(((T*)buf_data[0])[i] % ((T*)buf_data[1])[i]));
+++ }
+++ free(buf_data[0]);
+++ free(buf_data[1]);
+++ buf_data[0] = buf_data[1] = NULL;
+++}
+++
+++#define DECL_TEST_SUB(type, alias) \
+++static void compiler_sub_ ##alias(void)\
+++{\
+++ test_exec<type, TEST_OP_SUB>("compiler_sub_" # alias);\
+++}\
+++MAKE_UTEST_FROM_FUNCTION(compiler_sub_ ## alias)
+++
+++#define DECL_TEST_ADD(type, alias) \
+++static void compiler_add_ ##alias(void)\
+++{\
+++ test_exec<type, TEST_OP_ADD>("compiler_add_" # alias);\
+++}\
+++MAKE_UTEST_FROM_FUNCTION(compiler_add_ ## alias)
+++
+++#define DECL_TEST_MUL(type, alias) \
+++static void compiler_mul_ ##alias(void)\
+++{\
+++ test_exec<type, TEST_OP_MUL>("compiler_mul_" # alias);\
+++}\
+++MAKE_UTEST_FROM_FUNCTION(compiler_mul_ ## alias)
+++
+++#define DECL_TEST_DIV(type, alias) \
+++static void compiler_div_ ##alias(void)\
+++{\
+++ test_exec<type, TEST_OP_DIV>("compiler_div_" # alias);\
+++}\
+++MAKE_UTEST_FROM_FUNCTION(compiler_div_ ## alias)
+++
+++#define DECL_TEST_REM(type, alias) \
+++static void compiler_rem_ ##alias(void)\
+++{\
+++ test_exec<type, TEST_OP_REM>("compiler_rem_" # alias);\
+++}\
+++MAKE_UTEST_FROM_FUNCTION(compiler_rem_ ## alias)
+++
+++DECL_TEST_SUB(int8_t, char);
+++DECL_TEST_SUB(uint8_t, uchar);
+++DECL_TEST_SUB(int16_t, short);
+++DECL_TEST_SUB(uint16_t, ushort);
+++DECL_TEST_SUB(int32_t, int);
+++DECL_TEST_SUB(uint32_t, uint);
+++
+++DECL_TEST_ADD(int8_t, char);
+++DECL_TEST_ADD(uint8_t, uchar);
+++DECL_TEST_ADD(int16_t, short);
+++DECL_TEST_ADD(uint16_t, ushort);
+++DECL_TEST_ADD(int32_t, int);
+++DECL_TEST_ADD(uint32_t, uint);
+++
+++DECL_TEST_MUL(int8_t, char);
+++DECL_TEST_MUL(uint8_t, uchar);
+++DECL_TEST_MUL(int16_t, short);
+++DECL_TEST_MUL(uint16_t, ushort);
+++DECL_TEST_MUL(int32_t, int);
+++DECL_TEST_MUL(uint32_t, uint);
+++
+++DECL_TEST_DIV(int8_t, char);
+++DECL_TEST_DIV(uint8_t, uchar);
+++DECL_TEST_DIV(int16_t, short);
+++DECL_TEST_DIV(uint16_t, ushort);
+++DECL_TEST_DIV(int32_t, int);
+++DECL_TEST_DIV(uint32_t, uint);
+++
+++DECL_TEST_REM(int8_t, char);
+++DECL_TEST_REM(uint8_t, uchar);
+++DECL_TEST_REM(int16_t, short);
+++DECL_TEST_REM(uint16_t, ushort);
+++DECL_TEST_REM(int32_t, int);
+++DECL_TEST_REM(uint32_t, uint);
++diff --git a/utests/compiler_sub_bytes.cpp b/utests/compiler_sub_bytes.cpp
++deleted file mode 100644
++index 740a8fd..0000000
++--- a/utests/compiler_sub_bytes.cpp
+++++ /dev/null
++@@ -1,35 +0,0 @@
++-#include "utest_helper.hpp"
++-
++-static void compiler_sub_bytes(void)
++-{
++- const size_t n = 16;
++-
++- // Setup kernel and buffers
++- OCL_CREATE_KERNEL("compiler_sub_bytes");
++- buf_data[0] = (int8_t*) malloc(sizeof(int8_t) * n);
++- buf_data[1] = (int8_t*) malloc(sizeof(int8_t) * n);
++- for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[0])[i] = (int8_t) rand();
++- for (uint32_t i = 0; i < n; ++i) ((int8_t*)buf_data[1])[i] = (int8_t) rand();
++- OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[0]);
++- OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int8_t), buf_data[1]);
++- OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int8_t), NULL);
++-
++- // Run the kernel
++- OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
++- OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
++- OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
++- globals[0] = n;
++- locals[0] = 16;
++- OCL_NDRANGE(1);
++-
++- // Check result
++- OCL_MAP_BUFFER(2);
++- for (uint32_t i = 0; i < n; ++i)
++- OCL_ASSERT(((int8_t*)buf_data[2])[i] == (int8_t)(((int8_t*)buf_data[0])[i] - ((int8_t*)buf_data[1])[i]));
++- free(buf_data[0]);
++- free(buf_data[1]);
++- buf_data[0] = buf_data[1] = NULL;
++-}
++-
++-MAKE_UTEST_FROM_FUNCTION(compiler_sub_bytes);
++-
++diff --git a/utests/compiler_sub_shorts.cpp b/utests/compiler_sub_shorts.cpp
++deleted file mode 100644
++index 7c24a56..0000000
++--- a/utests/compiler_sub_shorts.cpp
+++++ /dev/null
++@@ -1,36 +0,0 @@
++-#include "utest_helper.hpp"
++-
++-static void compiler_sub_shorts(void)
++-{
++- const size_t n = 16;
++-
++- // Setup kernel and buffers
++- OCL_CREATE_KERNEL("compiler_sub_shorts");
++- buf_data[0] = (int16_t*) malloc(sizeof(int16_t) * n);
++- buf_data[1] = (int16_t*) malloc(sizeof(int16_t) * n);
++- for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[0])[i] = (int16_t) rand();
++- for (uint32_t i = 0; i < n; ++i) ((int16_t*)buf_data[1])[i] = (int16_t) rand();
++- OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[0]);
++- OCL_CREATE_BUFFER(buf[1], CL_MEM_COPY_HOST_PTR, n * sizeof(int16_t), buf_data[1]);
++- OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int16_t), NULL);
++-
++- // Run the kernel
++- OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
++- OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
++- OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
++- globals[0] = n;
++- locals[0] = 16;
++- OCL_NDRANGE(1);
++-
++- // Check result
++- OCL_MAP_BUFFER(2);
++- for (uint32_t i = 0; i < n; ++i)
++- OCL_ASSERT(((int16_t*)buf_data[2])[i] == (int16_t)(((int16_t*)buf_data[0])[i] - ((int16_t*)buf_data[1])[i]));
++- free(buf_data[0]);
++- free(buf_data[1]);
++- buf_data[0] = buf_data[1] = NULL;
++-}
++-
++-MAKE_UTEST_FROM_FUNCTION(compiler_sub_shorts);
++-
++-
++--
++1.7.10.4
++
diff --cc debian/patches/0005-Add-the-builtin-function-abs-and-the-according-test-.patch
index 0000000,0000000..0f0490e
new file mode 100644
--- /dev/null
+++ b/debian/patches/0005-Add-the-builtin-function-abs-and-the-according-test-.patch
@@@ -1,0 -1,0 +1,218 @@@
++From 1e18f92e0de249a403d4a50842d778d61a84d053 Mon Sep 17 00:00:00 2001
++From: Junyan He <junyan.he at linux.intel.com>
++Date: Tue, 25 Jun 2013 15:50:54 +0800
++Subject: [PATCH 5/8] Add the builtin function abs() and the according test
++ case
++To: beignet at lists.freedesktop.org
++
++Signed-off-by: Junyan He <junyan.he at linux.intel.com>
++---
++ backend/src/backend/gen_insn_selection.cpp | 10 ++++-
++ backend/src/llvm/llvm_gen_backend.cpp | 10 ++++-
++ backend/src/llvm/llvm_gen_ocl_function.hxx | 3 +-
++ backend/src/ocl_stdlib.h | 4 ++
++ kernels/compiler_abs.cl | 5 +++
++ utests/CMakeLists.txt | 1 +
++ utests/compiler_abs.cpp | 61 ++++++++++++++++++++++++++++
++ 7 files changed, 91 insertions(+), 3 deletions(-)
++ create mode 100644 kernels/compiler_abs.cl
++ create mode 100644 utests/compiler_abs.cpp
++
++diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
++index b1c6093..4f5b337 100644
++--- a/backend/src/backend/gen_insn_selection.cpp
+++++ b/backend/src/backend/gen_insn_selection.cpp
++@@ -1216,7 +1216,15 @@ namespace gbe
++ const GenRegister dst = sel.selReg(insn.getDst(0));
++ const GenRegister src = sel.selReg(insn.getSrc(0));
++ switch (opcode) {
++- case ir::OP_ABS: sel.MOV(dst, GenRegister::abs(src)); break;
+++ case ir::OP_ABS:
+++ if (insn.getType() == ir::TYPE_S32 || insn.getType() == ir::TYPE_U32) {
+++ const GenRegister src_ = GenRegister::retype(src, GEN_TYPE_D);
+++ const GenRegister dst_ = GenRegister::retype(dst, GEN_TYPE_D);
+++ sel.MOV(dst_, GenRegister::abs(src_));
+++ } else {
+++ sel.MOV(dst, GenRegister::abs(src));
+++ }
+++ break;
++ case ir::OP_MOV:
++ if (dst.isdf()) {
++ ir::Register r = sel.reg(ir::RegisterFamily::FAMILY_QWORD);
++diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
++index b0e8c6c..71f41b4 100644
++--- a/backend/src/llvm/llvm_gen_backend.cpp
+++++ b/backend/src/llvm/llvm_gen_backend.cpp
++@@ -1686,6 +1686,7 @@ namespace gbe
++ case GEN_OCL_POW:
++ case GEN_OCL_RCP:
++ case GEN_OCL_ABS:
+++ case GEN_OCL_FABS:
++ case GEN_OCL_RNDZ:
++ case GEN_OCL_RNDE:
++ case GEN_OCL_RNDU:
++@@ -1842,13 +1843,20 @@ namespace gbe
++ ctx.POW(ir::TYPE_FLOAT, dst, src0, src1);
++ break;
++ }
+++ case GEN_OCL_ABS:
+++ {
+++ const ir::Register src = this->getRegister(*AI);
+++ const ir::Register dst = this->getRegister(&I);
+++ ctx.ALU1(ir::OP_ABS, ir::TYPE_S32, dst, src);
+++ break;
+++ }
++ case GEN_OCL_COS: this->emitUnaryCallInst(I,CS,ir::OP_COS); break;
++ case GEN_OCL_SIN: this->emitUnaryCallInst(I,CS,ir::OP_SIN); break;
++ case GEN_OCL_LOG: this->emitUnaryCallInst(I,CS,ir::OP_LOG); break;
++ case GEN_OCL_SQR: this->emitUnaryCallInst(I,CS,ir::OP_SQR); break;
++ case GEN_OCL_RSQ: this->emitUnaryCallInst(I,CS,ir::OP_RSQ); break;
++ case GEN_OCL_RCP: this->emitUnaryCallInst(I,CS,ir::OP_RCP); break;
++- case GEN_OCL_ABS: this->emitUnaryCallInst(I,CS,ir::OP_ABS); break;
+++ case GEN_OCL_FABS: this->emitUnaryCallInst(I,CS,ir::OP_ABS); break;
++ case GEN_OCL_RNDZ: this->emitUnaryCallInst(I,CS,ir::OP_RNDZ); break;
++ case GEN_OCL_RNDE: this->emitUnaryCallInst(I,CS,ir::OP_RNDE); break;
++ case GEN_OCL_RNDU: this->emitUnaryCallInst(I,CS,ir::OP_RNDU); break;
++diff --git a/backend/src/llvm/llvm_gen_ocl_function.hxx b/backend/src/llvm/llvm_gen_ocl_function.hxx
++index 6cd7298..9cfad78 100644
++--- a/backend/src/llvm/llvm_gen_ocl_function.hxx
+++++ b/backend/src/llvm/llvm_gen_ocl_function.hxx
++@@ -19,7 +19,8 @@ DECL_LLVM_GEN_FUNCTION(GET_GLOBAL_OFFSET2, __gen_ocl_get_global_offset2)
++ DECL_LLVM_GEN_FUNCTION(GET_WORK_DIM, __gen_ocl_get_work_dim)
++
++ // Math function
++-DECL_LLVM_GEN_FUNCTION(ABS, __gen_ocl_fabs)
+++DECL_LLVM_GEN_FUNCTION(ABS, __gen_ocl_abs)
+++DECL_LLVM_GEN_FUNCTION(FABS, __gen_ocl_fabs)
++ DECL_LLVM_GEN_FUNCTION(COS, __gen_ocl_cos)
++ DECL_LLVM_GEN_FUNCTION(SIN, __gen_ocl_sin)
++ DECL_LLVM_GEN_FUNCTION(SQR, __gen_ocl_sqrt)
++diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
++index 81a0193..eaf8f21 100644
++--- a/backend/src/ocl_stdlib.h
+++++ b/backend/src/ocl_stdlib.h
++@@ -4337,6 +4337,7 @@ INLINE uint get_global_id(uint dim) {
++ /////////////////////////////////////////////////////////////////////////////
++ // Math Functions (see 6.11.2 of OCL 1.1 spec)
++ /////////////////////////////////////////////////////////////////////////////
+++PURE CONST int __gen_ocl_abs(int x);
++ PURE CONST float __gen_ocl_fabs(float x);
++ PURE CONST float __gen_ocl_sin(float x);
++ PURE CONST float __gen_ocl_cos(float x);
++@@ -4515,6 +4516,7 @@ INLINE_OVERLOADABLE float __gen_ocl_internal_erfc(float x) {
++ // XXX work-around PTX profile
++ #define sqrt native_sqrt
++ INLINE_OVERLOADABLE float rsqrt(float x) { return native_rsqrt(x); }
+++INLINE_OVERLOADABLE int __gen_ocl_internal_abs(int x) { return __gen_ocl_abs(x); }
++ INLINE_OVERLOADABLE float __gen_ocl_internal_fabs(float x) { return __gen_ocl_fabs(x); }
++ INLINE_OVERLOADABLE float __gen_ocl_internal_trunc(float x) { return __gen_ocl_rndz(x); }
++ INLINE_OVERLOADABLE float __gen_ocl_internal_round(float x) { return __gen_ocl_rnde(x); }
++@@ -4860,6 +4862,7 @@ DECL_VECTOR_1OP(native_exp10, float);
++ DECL_VECTOR_1OP(__gen_ocl_internal_expm1, float);
++ DECL_VECTOR_1OP(__gen_ocl_internal_cbrt, float);
++ DECL_VECTOR_1OP(__gen_ocl_internal_fabs, float);
+++DECL_VECTOR_1OP(__gen_ocl_internal_abs, int);
++ DECL_VECTOR_1OP(__gen_ocl_internal_trunc, float);
++ DECL_VECTOR_1OP(__gen_ocl_internal_round, float);
++ DECL_VECTOR_1OP(__gen_ocl_internal_floor, float);
++@@ -4987,6 +4990,7 @@ INLINE_OVERLOADABLE float8 mix(float8 x, float8 y, float a) { return mix(x,y,(fl
++ INLINE_OVERLOADABLE float16 mix(float16 x, float16 y, float a) { return mix(x,y,(float16)(a));}
++
++ // XXX workaround ptx profile
+++#define abs __gen_ocl_internal_abs
++ #define fabs __gen_ocl_internal_fabs
++ #define trunc __gen_ocl_internal_trunc
++ #define round __gen_ocl_internal_round
++diff --git a/kernels/compiler_abs.cl b/kernels/compiler_abs.cl
++new file mode 100644
++index 0000000..7030a26
++--- /dev/null
+++++ b/kernels/compiler_abs.cl
++@@ -0,0 +1,5 @@
+++kernel void compiler_abs(global int *src, global int *dst) {
+++ int i = get_global_id(0);
+++ dst[i] = abs(src[i]);
+++}
+++
++diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
++index edfbda1..d63b31f 100644
++--- a/utests/CMakeLists.txt
+++++ b/utests/CMakeLists.txt
++@@ -32,6 +32,7 @@ set (utests_sources
++ compiler_double_2.cpp
++ compiler_double_3.cpp
++ compiler_fabs.cpp
+++ compiler_abs.cpp
++ compiler_fill_image.cpp
++ compiler_fill_image0.cpp
++ compiler_fill_image_3d.cpp
++diff --git a/utests/compiler_abs.cpp b/utests/compiler_abs.cpp
++new file mode 100644
++index 0000000..908a32a
++--- /dev/null
+++++ b/utests/compiler_abs.cpp
++@@ -0,0 +1,61 @@
+++#include "utest_helper.hpp"
+++
+++static void cpu(int global_id, int *src, int *dst) {
+++ int f = src[global_id];
+++ f = f < 0 ? -f : f;
+++ dst[global_id] = f;
+++}
+++
+++void compiler_abs(void)
+++{
+++ const size_t n = 16;
+++ int cpu_dst[16], cpu_src[16];
+++
+++ // Setup kernel and buffers
+++ OCL_CREATE_KERNEL("compiler_abs");
+++ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int), NULL);
+++ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int), NULL);
+++ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+++ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+++ globals[0] = 16;
+++ locals[0] = 16;
+++
+++ // Run random tests
+++ for (uint32_t pass = 0; pass < 8; ++pass) {
+++ OCL_MAP_BUFFER(0);
+++ for (int32_t i = 0; i < (int32_t) n; ++i)
+++ cpu_src[i] = ((int*)buf_data[0])[i] = (rand() & 15) - 7;
+++
+++ // Run the kernel on GPU
+++ OCL_NDRANGE(1);
+++
+++ // Run on CPU
+++ for (int32_t i = 0; i < (int32_t) n; ++i) cpu(i, cpu_src, cpu_dst);
+++
+++ // Compare
+++ OCL_MAP_BUFFER(1);
+++
+++#if 0
+++ printf("Raw DATA: \n");
+++ for (int32_t i = 0; i < (int32_t) n; ++i) {
+++ printf(" %d", ((int *)buf_data[0])[i]);
+++ }
+++
+++ printf("\nCPU: \n");
+++ for (int32_t i = 0; i < (int32_t) n; ++i) {
+++ printf(" %d", cpu_dst[i]);
+++ }
+++ printf("\nGPU: \n");
+++ for (int32_t i = 0; i < (int32_t) n; ++i) {
+++ printf(" %d", ((int *)buf_data[1])[i]);
+++ }
+++#endif
+++
+++ for (int32_t i = 0; i < (int32_t) n; ++i)
+++ OCL_ASSERT(((int *)buf_data[1])[i] == cpu_dst[i]);
+++ OCL_UNMAP_BUFFER(1);
+++ OCL_UNMAP_BUFFER(0);
+++ }
+++}
+++
+++MAKE_UTEST_FROM_FUNCTION(compiler_abs);
++--
++1.7.10.4
++
diff --cc debian/patches/0006-PATCH-Refine-the-get_local_id-.-builtins.patch
index 0000000,0000000..25f2f39
new file mode 100644
--- /dev/null
+++ b/debian/patches/0006-PATCH-Refine-the-get_local_id-.-builtins.patch
@@@ -1,0 -1,0 +1,55 @@@
++From 8a2826a9a7b2085935f5396ae633645d35255b17 Mon Sep 17 00:00:00 2001
++From: Zhigang Gong <zhigang.gong at linux.intel.com>
++Date: Tue, 25 Jun 2013 14:15:09 +0800
++Subject: [PATCH 6/8] [PATCH] Refine the get_local_id/... builtins.
++To: beignet at lists.freedesktop.org
++
++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 at linux.intel.com>
++---
++ backend/src/ocl_stdlib.h | 8 +++-----
++ src/cl_api.c | 2 +-
++ 2 files changed, 4 insertions(+), 6 deletions(-)
++
++diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
++index eaf8f21..dd70da9 100644
++--- a/backend/src/ocl_stdlib.h
+++++ b/backend/src/ocl_stdlib.h
++@@ -4315,11 +4315,9 @@ DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)
++ #define DECL_PUBLIC_WORK_ITEM_FN(NAME, OTHER_RET) \
++ INLINE unsigned NAME(unsigned int dim) { \
++ if (dim == 0) return __gen_ocl_##NAME##0(); \
++- else if (dim > 0 && dim < get_work_dim()) { \
++- if (dim == 1) return __gen_ocl_##NAME##1(); \
++- else if (dim == 2) return __gen_ocl_##NAME##2(); \
++- } \
++- return OTHER_RET; \
+++ else if (dim == 1) return __gen_ocl_##NAME##1(); \
+++ else if (dim == 2) return __gen_ocl_##NAME##2(); \
+++ else return OTHER_RET; \
++ }
++
++ DECL_PUBLIC_WORK_ITEM_FN(get_group_id, 0)
++diff --git a/src/cl_api.c b/src/cl_api.c
++index 3c78243..ebca294 100644
++--- a/src/cl_api.c
+++++ b/src/cl_api.c
++@@ -1570,7 +1570,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
++ {
++ size_t fixed_global_off[] = {0,0,0};
++ size_t fixed_global_sz[] = {1,1,1};
++- size_t fixed_local_sz[] = {16,1,1};
+++ size_t fixed_local_sz[] = {1,1,1};
++ cl_int err = CL_SUCCESS;
++ cl_uint i;
++
++--
++1.7.10.4
++
diff --cc debian/patches/0007-Add-the-support-of-the-API-clGetCommandQueueInfo.patch
index 0000000,0000000..e43e1cc
new file mode 100644
--- /dev/null
+++ b/debian/patches/0007-Add-the-support-of-the-API-clGetCommandQueueInfo.patch
@@@ -1,0 -1,0 +1,185 @@@
++From 1e1556a65ff1779dca6a5b30c204985ad14f0e4f Mon Sep 17 00:00:00 2001
++From: Junyan He <junyan.he at linux.intel.com>
++Date: Tue, 25 Jun 2013 18:15:25 +0800
++Subject: [PATCH 7/8] Add the support of the API: clGetCommandQueueInfo
++To: beignet at lists.freedesktop.org
++
++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 at linux.intel.com>
++---
++ src/cl_api.c | 63 ++++++++++++++++++++++++++++--------------------
++ src/cl_command_queue.h | 15 ++++++------
++ src/cl_context.c | 1 +
++ 3 files changed, 46 insertions(+), 33 deletions(-)
++
++diff --git a/src/cl_api.c b/src/cl_api.c
++index ebca294..bb09c07 100644
++--- a/src/cl_api.c
+++++ b/src/cl_api.c
++@@ -46,6 +46,19 @@
++ typedef intptr_t cl_device_partition_property;
++ #endif
++
+++#define FILL_GETINFO_RET(TYPE, ELT, VAL, RET) \
+++ do { \
+++ if (param_value && param_value_size < sizeof(TYPE)*ELT) \
+++ return CL_INVALID_VALUE; \
+++ if (param_value) { \
+++ memcpy(param_value, (VAL), sizeof(TYPE)*ELT); \
+++ } \
+++ \
+++ if (param_value_size_ret) \
+++ *param_value_size_ret = sizeof(TYPE)*ELT; \
+++ return RET; \
+++ } while(0)
+++
++ static cl_int
++ cl_check_device_type(cl_device_type device_type)
++ {
++@@ -341,7 +354,20 @@ clGetCommandQueueInfo(cl_command_queue command_queue,
++ {
++ cl_int err = CL_SUCCESS;
++ CHECK_QUEUE (command_queue);
++- NOT_IMPLEMENTED;
+++
+++ if (param_name == CL_QUEUE_CONTEXT) {
+++ FILL_GETINFO_RET (cl_context, 1, &command_queue->ctx, CL_SUCCESS);
+++ } else if (param_name == CL_QUEUE_DEVICE) {
+++ FILL_GETINFO_RET (cl_device_id, 1, &command_queue->ctx->device, CL_SUCCESS);
+++ } else if (param_name == CL_QUEUE_REFERENCE_COUNT) {
+++ cl_uint ref = command_queue->ref_n;
+++ FILL_GETINFO_RET (cl_uint, 1, &ref, CL_SUCCESS);
+++ } else if (param_name == CL_QUEUE_PROPERTIES) {
+++ FILL_GETINFO_RET (cl_command_queue_properties, 1, &command_queue->props, CL_SUCCESS);
+++ } else {
+++ return CL_INVALID_VALUE;
+++ }
+++
++ error:
++ return err;
++ }
++@@ -734,19 +760,6 @@ clUnloadCompiler(void)
++ return 0;
++ }
++
++-#define FILL_AND_RET(TYPE, ELT, VAL, RET) \
++- do { \
++- if (param_value && param_value_size < sizeof(TYPE)*ELT) \
++- return CL_INVALID_VALUE; \
++- if (param_value) { \
++- memcpy(param_value, (VAL), sizeof(TYPE)*ELT); \
++- } \
++- \
++- if (param_value_size_ret) \
++- *param_value_size_ret = sizeof(TYPE)*ELT; \
++- return RET; \
++- } while(0)
++-
++ cl_int
++ clGetProgramInfo(cl_program program,
++ cl_program_info param_name,
++@@ -761,24 +774,24 @@ clGetProgramInfo(cl_program program,
++
++ if (param_name == CL_PROGRAM_REFERENCE_COUNT) {
++ cl_uint ref = program->ref_n;
++- FILL_AND_RET (cl_uint, 1, (&ref), CL_SUCCESS);
+++ FILL_GETINFO_RET (cl_uint, 1, (&ref), CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_CONTEXT) {
++ cl_context context = program->ctx;
++- FILL_AND_RET (cl_context, 1, &context, CL_SUCCESS);
+++ FILL_GETINFO_RET (cl_context, 1, &context, CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_NUM_DEVICES) {
++ cl_uint num_dev = 1; // Just 1 dev now.
++- FILL_AND_RET (cl_uint, 1, &num_dev, CL_SUCCESS);
+++ FILL_GETINFO_RET (cl_uint, 1, &num_dev, CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_DEVICES) {
++ cl_device_id dev_id = program->ctx->device;
++- FILL_AND_RET (cl_device_id, 1, &dev_id, CL_SUCCESS);
+++ FILL_GETINFO_RET (cl_device_id, 1, &dev_id, CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_SOURCE) {
++
++ if (!program->source)
++- FILL_AND_RET (char, 1, &ret_str, CL_SUCCESS);
++- FILL_AND_RET (char, (strlen(program->source) + 1),
+++ FILL_GETINFO_RET (char, 1, &ret_str, CL_SUCCESS);
+++ FILL_GETINFO_RET (char, (strlen(program->source) + 1),
++ program->source, CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_BINARY_SIZES) {
++- FILL_AND_RET (size_t, 1, (&program->bin_sz), CL_SUCCESS);
+++ FILL_GETINFO_RET (size_t, 1, (&program->bin_sz), CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_BINARIES) {
++ if (!param_value)
++ return CL_SUCCESS;
++@@ -825,15 +838,15 @@ clGetProgramBuildInfo(cl_program program,
++ status = CL_BUILD_ERROR;
++ // TODO: Support CL_BUILD_IN_PROGRESS ?
++
++- FILL_AND_RET (cl_build_status, 1, &status, CL_SUCCESS);
+++ FILL_GETINFO_RET (cl_build_status, 1, &status, CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_BUILD_OPTIONS) {
++ if (program->is_built && program->build_opts)
++ ret_str = program->build_opts;
++
++- FILL_AND_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
+++ FILL_GETINFO_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
++ } else if (param_name == CL_PROGRAM_BUILD_LOG) {
++ // TODO: need to add logs in backend when compiling.
++- FILL_AND_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
+++ FILL_GETINFO_RET (char, (strlen(ret_str)+1), ret_str, CL_SUCCESS);
++ } else {
++ return CL_INVALID_VALUE;
++ }
++@@ -842,8 +855,6 @@ error:
++ return err;
++ }
++
++-#undef FILL_AND_RET
++-
++ cl_kernel
++ clCreateKernel(cl_program program,
++ const char * kernel_name,
++diff --git a/src/cl_command_queue.h b/src/cl_command_queue.h
++index 0e04ff3..4f6f987 100644
++--- a/src/cl_command_queue.h
+++++ b/src/cl_command_queue.h
++@@ -30,13 +30,14 @@ struct intel_gpgpu;
++ /* Basically, this is a (kind-of) batch buffer */
++ struct _cl_command_queue {
++ DEFINE_ICD(dispatch)
++- uint64_t magic; /* To identify it as a command queue */
++- volatile int ref_n; /* We reference count this object */
++- cl_context ctx; /* Its parent context */
++- cl_command_queue prev, next; /* We chain the command queues together */
++- cl_gpgpu gpgpu; /* Setup all GEN commands */
++- cl_mem perf; /* Where to put the perf counters */
++- cl_mem fulsim_out; /* Fulsim will output this buffer */
+++ uint64_t magic; /* To identify it as a command queue */
+++ volatile int ref_n; /* We reference count this object */
+++ cl_context ctx; /* Its parent context */
+++ cl_command_queue_properties props; /* Queue properties */
+++ cl_command_queue prev, next; /* We chain the command queues together */
+++ cl_gpgpu gpgpu; /* Setup all GEN commands */
+++ cl_mem perf; /* Where to put the perf counters */
+++ cl_mem fulsim_out; /* Fulsim will output this buffer */
++ };
++
++ /* Allocate and initialize a new command queue. Also insert it in the list of
++diff --git a/src/cl_context.c b/src/cl_context.c
++index fa4c7e0..0331151 100644
++--- a/src/cl_context.c
+++++ b/src/cl_context.c
++@@ -196,6 +196,7 @@ cl_context_create_queue(cl_context ctx,
++
++ /* We create the command queue and store it in the context list of queues */
++ TRY_ALLOC (queue, cl_command_queue_new(ctx));
+++ queue->props = properties;
++
++ exit:
++ if (errcode_ret)
++--
++1.7.10.4
++
diff --cc debian/patches/0008-Add-the-test-case-for-clGetCommandQueueInfo-API.patch
index 0000000,0000000..5bfcf07
new file mode 100644
--- /dev/null
+++ b/debian/patches/0008-Add-the-test-case-for-clGetCommandQueueInfo-API.patch
@@@ -1,0 -1,0 +1,631 @@@
++From bb741a356647c2959135844373c6d7287cedfd2f Mon Sep 17 00:00:00 2001
++From: Junyan He <junyan.he at linux.intel.com>
++Date: Tue, 25 Jun 2013 18:15:32 +0800
++Subject: [PATCH 8/8] Add the test case for clGetCommandQueueInfo API
++To: beignet at lists.freedesktop.org
++
++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 at linux.intel.com>
++---
++ kernels/get_program_info.cl | 10 --
++ utests/CMakeLists.txt | 2 +-
++ utests/get_cl_info.cpp | 319 +++++++++++++++++++++++++++++++++++++++++++
++ utests/get_program_info.cpp | 247 ---------------------------------
++ 4 files changed, 320 insertions(+), 258 deletions(-)
++ delete mode 100644 kernels/get_program_info.cl
++ create mode 100644 utests/get_cl_info.cpp
++ delete mode 100644 utests/get_program_info.cpp
++
++diff --git a/kernels/get_program_info.cl b/kernels/get_program_info.cl
++deleted file mode 100644
++index 8e0dd94..0000000
++--- a/kernels/get_program_info.cl
+++++ /dev/null
++@@ -1,10 +0,0 @@
++-kernel void get_program_info( __global int *ret ) {
++- uint x = get_work_dim();
++- size_t y = get_global_size(0);
++- y = get_global_id(0);
++- y = get_local_size(0);
++- y = get_local_id(0);
++- y = get_num_groups(0);
++- y = get_group_id(0);
++- y = get_global_offset(0);
++-}
++diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
++index d63b31f..e12069c 100644
++--- a/utests/CMakeLists.txt
+++++ b/utests/CMakeLists.txt
++@@ -86,7 +86,7 @@ set (utests_sources
++ compiler_get_image_info.cpp
++ compiler_vector_load_store.cpp
++ compiler_cl_finish.cpp
++- get_program_info.cpp
+++ get_cl_info.cpp
++ buildin_work_dim.cpp
++ builtin_global_size.cpp
++ runtime_createcontext.cpp
++diff --git a/utests/get_cl_info.cpp b/utests/get_cl_info.cpp
++new file mode 100644
++index 0000000..6d5e7bb
++--- /dev/null
+++++ b/utests/get_cl_info.cpp
++@@ -0,0 +1,319 @@
+++#include <string.h>
+++#include <string>
+++#include <map>
+++#include <iostream>
+++#include <fstream>
+++#include <algorithm>
+++#include "utest_helper.hpp"
+++
+++using namespace std;
+++
+++/* ***************************************************** *
+++ * This file to test all the API like: clGetXXXXInfo *
+++ * ***************************************************** */
+++#define NO_STANDARD_REF 0xFFFFF
+++
+++template <typename T = cl_uint>
+++struct Info_Result {
+++ T ret;
+++ T refer;
+++ int size;
+++ typedef T type_value;
+++
+++ void * get_ret(void) {
+++ return (void *)&ret;
+++ }
+++
+++ Info_Result(T other) {
+++ refer = other;
+++ size = sizeof(T);
+++ }
+++
+++ bool check_result (void) {
+++ if (ret != refer && refer != (T)NO_STANDARD_REF)
+++ return false;
+++
+++ return true;
+++ }
+++};
+++
+++template <>
+++struct Info_Result<char *> {
+++ char * ret;
+++ char * refer;
+++ int size;
+++ typedef char* type_value;
+++
+++ Info_Result(char *other, int sz) {
+++ size = sz;
+++ ret = (char *)malloc(sizeof(char) * sz);
+++ if (other) {
+++ refer = (char *)malloc(sizeof(char) * sz);
+++ memcpy(refer, other, sz);
+++ }
+++ }
+++
+++ ~Info_Result(void) {
+++ free(refer);
+++ free(ret);
+++ }
+++
+++ void * get_ret(void) {
+++ return (void *)ret;
+++ }
+++
+++ bool check_result (void) {
+++ if (refer && ::memcmp(ret, refer, size))
+++ return false;
+++
+++ return true;
+++ }
+++};
+++
+++template <> //Used for such as CL_PROGRAM_BINARIES
+++struct Info_Result<char **> {
+++ char ** ret;
+++ char ** refer;
+++ int *elt_size;
+++ int size;
+++ typedef char** type_value;
+++
+++ Info_Result(char **other, int *sz, int elt_num) {
+++ size = elt_num;
+++
+++ ret = (char **)malloc(elt_num * sizeof(char *));
+++ memset(ret, 0, (elt_num * sizeof(char *)));
+++ refer = (char **)malloc(elt_num * sizeof(char *));
+++ memset(refer, 0, (elt_num * sizeof(char *)));
+++ elt_size = (int *)malloc(elt_num * sizeof(int));
+++ memset(elt_size, 0, (elt_num * sizeof(int)));
+++ if (sz) {
+++ int i = 0;
+++ for (; i < elt_num; i++) {
+++ elt_size[i] = sz[i];
+++ ret[i] = (char *)malloc(sz[i] * sizeof(char));
+++
+++ if (other[i] && elt_size[i] > 0) {
+++ refer[i] = (char *)malloc(sz[i] * sizeof(char));
+++ memcpy(&refer[i], &other[i], sz[i]);
+++ }
+++ else
+++ refer[i] = NULL;
+++ }
+++ }
+++ }
+++
+++ ~Info_Result(void) {
+++ int i = 0;
+++ for (; i < size; i++) {
+++ if (refer[i])
+++ free(refer[i]);
+++ free(ret[i]);
+++ }
+++ free(ret);
+++ free(refer);
+++ free(elt_size);
+++ }
+++
+++ void * get_ret(void) {
+++ return (void *)ret;
+++ }
+++
+++ bool check_result (void) {
+++ int i = 0;
+++ for (; i < size; i++) {
+++ if (refer[i] && ::memcmp(ret[i], refer[i], elt_size[i]))
+++ return false;
+++ }
+++
+++ return true;
+++ }
+++};
+++
+++template <typename T1, typename T2>
+++struct Traits {
+++ static bool Is_Same(void) {
+++ return false;
+++ };
+++};
+++
+++template <typename T1>
+++struct Traits<T1, T1> {
+++ static bool Is_Same(void) {
+++ return true;
+++ };
+++};
+++
+++template <typename T>
+++Info_Result<T>* cast_as(void *info)
+++{
+++ Info_Result<T>* ret;
+++ ret = reinterpret_cast<Info_Result<T>*>(info);
+++ OCL_ASSERT((Traits<T, typename Info_Result<T>::type_value>::Is_Same()));
+++ return ret;
+++}
+++
+++
+++#define CALL_INFO_AND_RET(TYPE, FUNC, OBJ) \
+++ do { \
+++ cl_int ret; \
+++ size_t ret_size; \
+++ \
+++ Info_Result<TYPE>* info = cast_as<TYPE>(x.second); \
+++ ret = FUNC (OBJ, x.first, \
+++ info->size, info->get_ret(), &ret_size); \
+++ OCL_ASSERT((!ret)); \
+++ OCL_ASSERT((info->check_result())); \
+++ delete info; \
+++ } while(0)
+++
+++/* ***************************************************** *
+++ * clGetProgramInfo *
+++ * ***************************************************** */
+++#define CALL_PROGINFO_AND_RET(TYPE) CALL_INFO_AND_RET(TYPE, clGetProgramInfo, program)
+++
+++void get_program_info(void)
+++{
+++ map<cl_program_info, void *> maps;
+++ int expect_value;
+++ char * expect_source;
+++ int sz;
+++ char *ker_path = (char *)malloc(4096 * sizeof(char));
+++ const char *kiss_path = getenv("OCL_KERNEL_PATH");
+++ string line;
+++ string source_code;
+++
+++ sprintf(ker_path, "%s/%s", kiss_path, "compiler_if_else.cl");
+++
+++ ifstream in(ker_path);
+++ while (getline(in,line)) {
+++ source_code = (source_code == "") ?
+++ source_code + line : source_code + "\n" + line;
+++ }
+++ free(ker_path);
+++ //cout<< source_code;
+++ source_code = source_code + "\n";
+++
+++ expect_source = (char *)source_code.c_str();
+++
+++ OCL_CREATE_KERNEL("compiler_if_else");
+++
+++ /* First test for clGetProgramInfo. We just have 1 devices now */
+++ expect_value = 2;//One program, one kernel.
+++ maps.insert(make_pair(CL_PROGRAM_REFERENCE_COUNT,
+++ (void *)(new Info_Result<>(((cl_uint)expect_value)))));
+++ maps.insert(make_pair(CL_PROGRAM_CONTEXT,
+++ (void *)(new Info_Result<cl_context>(ctx))));
+++ expect_value = 1;
+++ maps.insert(make_pair(CL_PROGRAM_NUM_DEVICES,
+++ (void *)(new Info_Result<>(((cl_uint)expect_value)))));
+++ maps.insert(make_pair(CL_PROGRAM_DEVICES,
+++ (void *)(new Info_Result<cl_device_id>(device))));
+++ sz = (strlen(expect_source) + 1);
+++ maps.insert(make_pair(CL_PROGRAM_SOURCE,
+++ (void *)(new Info_Result<char *>(expect_source, sz))));
+++ expect_value = NO_STANDARD_REF;
+++ maps.insert(make_pair(CL_PROGRAM_BINARY_SIZES,
+++ (void *)(new Info_Result<size_t>((size_t)expect_value))));
+++ sz = 4096; //big enough?
+++ expect_source = NULL;
+++ maps.insert(make_pair(CL_PROGRAM_BINARIES,
+++ (void *)(new Info_Result<char **>(&expect_source, &sz, 1))));
+++
+++ std::for_each(maps.begin(), maps.end(), [](pair<cl_program_info, void *> x) {
+++ switch (x.first) {
+++ case CL_PROGRAM_REFERENCE_COUNT:
+++ case CL_PROGRAM_NUM_DEVICES:
+++ CALL_PROGINFO_AND_RET(cl_uint);
+++ break;
+++ case CL_PROGRAM_CONTEXT:
+++ CALL_PROGINFO_AND_RET(cl_context);
+++ break;
+++ case CL_PROGRAM_DEVICES:
+++ CALL_PROGINFO_AND_RET(cl_device_id);
+++ break;
+++ case CL_PROGRAM_SOURCE:
+++ CALL_PROGINFO_AND_RET(char *);
+++ break;
+++ case CL_PROGRAM_BINARY_SIZES:
+++ CALL_PROGINFO_AND_RET(size_t);
+++ break;
+++ case CL_PROGRAM_BINARIES:
+++ CALL_PROGINFO_AND_RET(char **);
+++ break;
+++ default:
+++ break;
+++ }
+++ });
+++}
+++
+++MAKE_UTEST_FROM_FUNCTION(get_program_info);
+++
+++/* ***************************************************** *
+++ * clGetCommandQueueInfo *
+++ * ***************************************************** */
+++#define CALL_QUEUEINFO_AND_RET(TYPE) CALL_INFO_AND_RET(TYPE, clGetCommandQueueInfo, queue)
+++
+++void get_queue_info(void)
+++{
+++ /* use the compiler_fabs case to test us. */
+++ const size_t n = 16;
+++ map<cl_program_info, void *> maps;
+++ int expect_ref;
+++ cl_command_queue_properties prop;
+++
+++ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
+++ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL);
+++ OCL_CREATE_KERNEL("compiler_fabs");
+++
+++ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+++ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+++
+++ globals[0] = 16;
+++ locals[0] = 16;
+++
+++ OCL_MAP_BUFFER(0);
+++ for (int32_t i = 0; i < (int32_t) n; ++i)
+++ ((float*)buf_data[0])[i] = .1f * (rand() & 15) - .75f;
+++ OCL_UNMAP_BUFFER(0);
+++
+++ // Run the kernel on GPU
+++ OCL_NDRANGE(1);
+++
+++ /* Do our test.*/
+++ maps.insert(make_pair(CL_QUEUE_CONTEXT,
+++ (void *)(new Info_Result<cl_context>(ctx))));
+++ maps.insert(make_pair(CL_QUEUE_DEVICE,
+++ (void *)(new Info_Result<cl_device_id>(device))));
+++
+++ expect_ref = 1;
+++ maps.insert(make_pair(CL_QUEUE_REFERENCE_COUNT,
+++ (void *)(new Info_Result<>(((cl_uint)expect_ref)))));
+++
+++ prop = 0;
+++ maps.insert(make_pair(CL_QUEUE_PROPERTIES,
+++ (void *)(new Info_Result<cl_command_queue_properties>(
+++ ((cl_command_queue_properties)prop)))));
+++
+++ std::for_each(maps.begin(), maps.end(), [](pair<cl_program_info, void *> x) {
+++ switch (x.first) {
+++ case CL_QUEUE_CONTEXT:
+++ CALL_QUEUEINFO_AND_RET(cl_context);
+++ break;
+++ case CL_QUEUE_DEVICE:
+++ CALL_QUEUEINFO_AND_RET(cl_device_id);
+++ break;
+++ case CL_QUEUE_REFERENCE_COUNT:
+++ CALL_QUEUEINFO_AND_RET(cl_uint);
+++ break;
+++ case CL_QUEUE_PROPERTIES:
+++ CALL_QUEUEINFO_AND_RET(cl_command_queue_properties);
+++ break;
+++ default:
+++ break;
+++ }
+++ });
+++}
+++
+++MAKE_UTEST_FROM_FUNCTION(get_queue_info);
+++
++diff --git a/utests/get_program_info.cpp b/utests/get_program_info.cpp
++deleted file mode 100644
++index 20248e8..0000000
++--- a/utests/get_program_info.cpp
+++++ /dev/null
++@@ -1,247 +0,0 @@
++-#include <string.h>
++-#include <string>
++-#include <map>
++-#include <iostream>
++-#include <fstream>
++-#include <algorithm>
++-#include "utest_helper.hpp"
++-
++-using namespace std;
++-
++-/* ********************************************** *
++- * This file to test the API of: *
++- * clGetProgramInfo *
++- * ********************************************** */
++-#define NO_STANDARD_REF 0xFFFFF
++-
++-template <typename T = cl_uint>
++-struct Info_Result {
++- T ret;
++- T refer;
++- int size;
++- typedef T type_value;
++-
++- void * get_ret(void) {
++- return (void *)&ret;
++- }
++-
++- Info_Result(T other) {
++- refer = other;
++- size = sizeof(T);
++- }
++-
++- bool check_result (void) {
++- if (ret != refer && refer != (T)NO_STANDARD_REF)
++- return false;
++-
++- return true;
++- }
++-};
++-
++-template <>
++-struct Info_Result<char *> {
++- char * ret;
++- char * refer;
++- int size;
++- typedef char* type_value;
++-
++- Info_Result(char *other, int sz) {
++- size = sz;
++- ret = (char *)malloc(sizeof(char) * sz);
++- if (other) {
++- refer = (char *)malloc(sizeof(char) * sz);
++- memcpy(refer, other, sz);
++- }
++- }
++-
++- ~Info_Result(void) {
++- free(refer);
++- free(ret);
++- }
++-
++- void * get_ret(void) {
++- return (void *)ret;
++- }
++-
++- bool check_result (void) {
++- if (refer && ::memcmp(ret, refer, size))
++- return false;
++-
++- return true;
++- }
++-};
++-
++-template <> //Used for such as CL_PROGRAM_BINARIES
++-struct Info_Result<char **> {
++- char ** ret;
++- char ** refer;
++- int *elt_size;
++- int size;
++- typedef char** type_value;
++-
++- Info_Result(char **other, int *sz, int elt_num) {
++- size = elt_num;
++-
++- ret = (char **)malloc(elt_num * sizeof(char *));
++- memset(ret, 0, (elt_num * sizeof(char *)));
++- refer = (char **)malloc(elt_num * sizeof(char *));
++- memset(refer, 0, (elt_num * sizeof(char *)));
++- elt_size = (int *)malloc(elt_num * sizeof(int));
++- memset(elt_size, 0, (elt_num * sizeof(int)));
++- if (sz) {
++- int i = 0;
++- for (; i < elt_num; i++) {
++- elt_size[i] = sz[i];
++- ret[i] = (char *)malloc(sz[i] * sizeof(char));
++-
++- if (other[i] && elt_size[i] > 0) {
++- refer[i] = (char *)malloc(sz[i] * sizeof(char));
++- memcpy(&refer[i], &other[i], sz[i]);
++- }
++- else
++- refer[i] = NULL;
++- }
++- }
++- }
++-
++- ~Info_Result(void) {
++- int i = 0;
++- for (; i < size; i++) {
++- if (refer[i])
++- free(refer[i]);
++- free(ret[i]);
++- }
++- free(ret);
++- free(refer);
++- free(elt_size);
++- }
++-
++- void * get_ret(void) {
++- return (void *)ret;
++- }
++-
++- bool check_result (void) {
++- int i = 0;
++- for (; i < size; i++) {
++- if (refer[i] && ::memcmp(ret[i], refer[i], elt_size[i]))
++- return false;
++- }
++-
++- return true;
++- }
++-};
++-
++-template <typename T1, typename T2>
++-struct Traits {
++- static bool Is_Same(void) {
++- return false;
++- };
++-};
++-
++-template <typename T1>
++-struct Traits<T1, T1> {
++- static bool Is_Same(void) {
++- return true;
++- };
++-};
++-
++-template <typename T>
++-Info_Result<T>* cast_as(void *info)
++-{
++- Info_Result<T>* ret;
++- ret = reinterpret_cast<Info_Result<T>*>(info);
++- OCL_ASSERT((Traits<T, typename Info_Result<T>::type_value>::Is_Same()));
++- return ret;
++-}
++-
++-
++-#define CALL_PROGINFO_AND_RET(TYPE) \
++- do { \
++- cl_int ret; \
++- size_t ret_size; \
++- \
++- Info_Result<TYPE>* info = cast_as<TYPE>(x.second); \
++- ret = clGetProgramInfo(program, x.first, \
++- info->size, info->get_ret(), &ret_size); \
++- OCL_ASSERT((!ret)); \
++- OCL_ASSERT((info->check_result())); \
++- delete info; \
++- } while(0)
++-
++-void get_program_info(void)
++-{
++- map<cl_program_info, void *> maps;
++- int expect_value;
++- char * expect_source;
++- int sz;
++- char *ker_path = (char *)malloc(4096 * sizeof(char));
++- const char *kiss_path = getenv("OCL_KERNEL_PATH");
++- string line;
++- string source_code;
++-
++- sprintf(ker_path, "%s/%s", kiss_path, "get_program_info.cl");
++-
++- ifstream in(ker_path);
++- while (getline(in,line)) {
++- source_code = (source_code == "") ?
++- source_code + line : source_code + "\n" + line;
++- }
++- free(ker_path);
++- //cout<< source_code;
++- source_code = source_code + "\n";
++-
++- expect_source = (char *)source_code.c_str();
++-
++- OCL_CREATE_KERNEL("get_program_info");
++-
++- /* First test for clGetProgramInfo. We just have 1 devices now */
++- expect_value = 2;//One program, one kernel.
++- maps.insert(make_pair(CL_PROGRAM_REFERENCE_COUNT,
++- (void *)(new Info_Result<>(((cl_uint)expect_value)))));
++- maps.insert(make_pair(CL_PROGRAM_CONTEXT,
++- (void *)(new Info_Result<cl_context>(ctx))));
++- expect_value = 1;
++- maps.insert(make_pair(CL_PROGRAM_NUM_DEVICES,
++- (void *)(new Info_Result<>(((cl_uint)expect_value)))));
++- maps.insert(make_pair(CL_PROGRAM_DEVICES,
++- (void *)(new Info_Result<cl_device_id>(device))));
++- sz = (strlen(expect_source) + 1);
++- maps.insert(make_pair(CL_PROGRAM_SOURCE,
++- (void *)(new Info_Result<char *>(expect_source, sz))));
++- expect_value = NO_STANDARD_REF;
++- maps.insert(make_pair(CL_PROGRAM_BINARY_SIZES,
++- (void *)(new Info_Result<size_t>((size_t)expect_value))));
++- sz = 4096; //big enough?
++- expect_source = NULL;
++- maps.insert(make_pair(CL_PROGRAM_BINARIES,
++- (void *)(new Info_Result<char **>(&expect_source, &sz, 1))));
++-
++- std::for_each(maps.begin(), maps.end(), [](pair<cl_program_info, void *> x) {
++- switch (x.first) {
++- case CL_PROGRAM_REFERENCE_COUNT:
++- case CL_PROGRAM_NUM_DEVICES:
++- CALL_PROGINFO_AND_RET(cl_uint);
++- break;
++- case CL_PROGRAM_CONTEXT:
++- CALL_PROGINFO_AND_RET(cl_context);
++- break;
++- case CL_PROGRAM_DEVICES:
++- CALL_PROGINFO_AND_RET(cl_device_id);
++- break;
++- case CL_PROGRAM_SOURCE:
++- CALL_PROGINFO_AND_RET(char *);
++- break;
++- case CL_PROGRAM_BINARY_SIZES:
++- CALL_PROGINFO_AND_RET(size_t);
++- break;
++- case CL_PROGRAM_BINARIES:
++- CALL_PROGINFO_AND_RET(char **);
++- break;
++- default:
++- break;
++- }
++- });
++-}
++-
++-MAKE_UTEST_FROM_FUNCTION(get_program_info);
++-
++--
++1.7.10.4
++
diff --cc debian/patches/series
index 036a483,0000000..6902ed0
mode 100644,000000..100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@@ -1,6 -1,0 +1,12 @@@
+debug
+flags
+khronos
+deprecated-in-utest
+private
+0001-Add-vector-argument-test-case.patch
++0003-GBE-Add-more-support-of-char-and-short-arithmetic.patch
++0004-utests-Add-basic-arithmetic-test-case.patch
++0005-Add-the-builtin-function-abs-and-the-according-test-.patch
++0006-PATCH-Refine-the-get_local_id-.-builtins.patch
++0007-Add-the-support-of-the-API-clGetCommandQueueInfo.patch
++0008-Add-the-test-case-for-clGetCommandQueueInfo-API.patch
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/beignet.git
More information about the Pkg-opencl-devel
mailing list