[Pkg-opencl-devel] [beignet] 22/47: Imported Debian patch 0.1+git20130625+97c3a9b-1
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 21:45:48 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 bda8a6e8e61b0bc42489c2eeedb73f948d3b9f4c
Author: Simon Richter <sjr at debian.org>
Date: Tue Jun 25 13:52:00 2013 +0200
Imported Debian patch 0.1+git20130625+97c3a9b-1
---
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 +
10 files changed, 1731 insertions(+), 15 deletions(-)
diff --git a/debian/changelog b/debian/changelog
index 38fba13..f7f7c49 100644
--- a/debian/changelog
+++ b/debian/changelog
@@ -1,3 +1,9 @@
+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
diff --git a/debian/patches/0001-Add-vector-argument-test-case.patch b/debian/patches/0001-Add-vector-argument-test-case.patch
index 34a125f..eac26c2 100644
--- a/debian/patches/0001-Add-vector-argument-test-case.patch
+++ b/debian/patches/0001-Add-vector-argument-test-case.patch
@@ -1,7 +1,7 @@
-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>
@@ -13,10 +13,11 @@ Signed-off-by: Yang Rong <rong.r.yang at intel.com>
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)
@@ -24,11 +25,11 @@ Index: beignet-0.1+git20130621+30586bf/kernels/compiler_function_argument2.cl
+ 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
@@ -36,10 +37,11 @@ Index: beignet-0.1+git20130621+30586bf/utests/CMakeLists.txt
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"
+
@@ -67,3 +69,6 @@ Index: beignet-0.1+git20130621+30586bf/utests/compiler_function_argument2.cpp
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_function_argument2);
+--
+1.7.10.4
+
diff --git a/debian/patches/0002-Readd-OpenCL-1.2-definitions-required-for-ICD.patch b/debian/patches/0002-Readd-OpenCL-1.2-definitions-required-for-ICD.patch
new file mode 100644
index 0000000..c385dcd
--- /dev/null
+++ b/debian/patches/0002-Readd-OpenCL-1.2-definitions-required-for-ICD.patch
@@ -0,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 --git a/debian/patches/0003-GBE-Add-more-support-of-char-and-short-arithmetic.patch b/debian/patches/0003-GBE-Add-more-support-of-char-and-short-arithmetic.patch
new file mode 100644
index 0000000..22a0eec
--- /dev/null
+++ b/debian/patches/0003-GBE-Add-more-support-of-char-and-short-arithmetic.patch
@@ -0,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 --git a/debian/patches/0004-utests-Add-basic-arithmetic-test-case.patch b/debian/patches/0004-utests-Add-basic-arithmetic-test-case.patch
new file mode 100644
index 0000000..a3ae637
--- /dev/null
+++ b/debian/patches/0004-utests-Add-basic-arithmetic-test-case.patch
@@ -0,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 --git a/debian/patches/0005-Add-the-builtin-function-abs-and-the-according-test-.patch b/debian/patches/0005-Add-the-builtin-function-abs-and-the-according-test-.patch
new file mode 100644
index 0000000..0f0490e
--- /dev/null
+++ b/debian/patches/0005-Add-the-builtin-function-abs-and-the-according-test-.patch
@@ -0,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 --git a/debian/patches/0006-PATCH-Refine-the-get_local_id-.-builtins.patch b/debian/patches/0006-PATCH-Refine-the-get_local_id-.-builtins.patch
new file mode 100644
index 0000000..25f2f39
--- /dev/null
+++ b/debian/patches/0006-PATCH-Refine-the-get_local_id-.-builtins.patch
@@ -0,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 --git a/debian/patches/0007-Add-the-support-of-the-API-clGetCommandQueueInfo.patch b/debian/patches/0007-Add-the-support-of-the-API-clGetCommandQueueInfo.patch
new file mode 100644
index 0000000..e43e1cc
--- /dev/null
+++ b/debian/patches/0007-Add-the-support-of-the-API-clGetCommandQueueInfo.patch
@@ -0,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 --git a/debian/patches/0008-Add-the-test-case-for-clGetCommandQueueInfo-API.patch b/debian/patches/0008-Add-the-test-case-for-clGetCommandQueueInfo-API.patch
new file mode 100644
index 0000000..5bfcf07
--- /dev/null
+++ b/debian/patches/0008-Add-the-test-case-for-clGetCommandQueueInfo-API.patch
@@ -0,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 --git a/debian/patches/series b/debian/patches/series
index 036a483..6902ed0 100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@ -4,3 +4,9 @@ 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