[Pkg-opencl-devel] [beignet] 37/66: Imported Debian patch 0.1+git20130626+41005e0-1
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 07:27:06 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 22f7d4eb1a4bbe62caaab3a7173da69659daa69c
Merge: e06914c a60dfce
Author: Simon Richter <sjr at debian.org>
Date: Wed Jun 26 10:06:29 2013 +0200
Imported Debian patch 0.1+git20130626+41005e0-1
backend/src/backend/gen/gen_mesa_disasm.c | 2 +
backend/src/backend/gen_context.cpp | 2 +
backend/src/backend/gen_defs.hpp | 2 +
backend/src/backend/gen_encoder.cpp | 2 +
backend/src/backend/gen_encoder.hpp | 2 +
backend/src/backend/gen_insn_selection.cpp | 14 +-
backend/src/backend/gen_insn_selection.hxx | 2 +
backend/src/ir/instruction.cpp | 2 +
backend/src/ir/instruction.hpp | 4 +
backend/src/ir/instruction.hxx | 2 +
backend/src/llvm/llvm_gen_backend.cpp | 4 +
backend/src/llvm/llvm_gen_ocl_function.hxx | 4 +
backend/src/ocl_stdlib.h | 66 ++-
debian/changelog | 6 +
.../0001-Add-vector-argument-test-case.patch | 8 +-
...d-OpenCL-1.2-definitions-required-for-ICD.patch | 4 +-
...tin-function-abs-and-the-according-test-.patch} | 39 +-
...PATCH-Refine-the-get_local_id-.-builtins.patch} | 12 +-
...help-functions.-Support-global-and-local-.patch | 611 ++++++++++++++++++++
...Add-all-atomic-built-in-functions-support.patch | 289 ++++++++++
debian/patches/0007-Add-atomic-test-case.patch | 188 ++++++
...-support-of-the-API-clGetCommandQueueInfo.patch | 185 ------
...e-test-case-for-clGetCommandQueueInfo-API.patch | 631 ---------------------
.../0008-support-built-in-function-rotate.patch | 58 ++
.../0009-test-case-for-function-rotate.patch | 87 +++
...ore-support-of-char-and-short-arithmetic.patch} | 128 +++--
...11-utests-Add-basic-arithmetic-test-case.patch} | 122 ++--
debian/patches/series | 15 +-
kernels/compiler_clz_int.cl | 5 +
kernels/compiler_clz_short.cl | 5 +
kernels/get_program_info.cl | 10 -
src/cl_api.c | 63 +-
src/cl_command_queue.h | 15 +-
src/cl_context.c | 1 +
utests/CMakeLists.txt | 4 +-
utests/compiler_clz_int.cpp | 31 +
utests/compiler_clz_short.cpp | 31 +
utests/compiler_shader_toy.cpp | 2 +-
utests/{get_program_info.cpp => get_cl_info.cpp} | 133 ++++-
utests/utest_helper.cpp | 6 +-
utests/utest_helper.hpp | 7 +-
41 files changed, 1751 insertions(+), 1053 deletions(-)
diff --cc debian/changelog
index f7f7c49,0000000..d5fb18e
mode 100644,000000..100644
--- a/debian/changelog
+++ b/debian/changelog
@@@ -1,139 -1,0 +1,145 @@@
++beignet (0.1+git20130626+41005e0-1) unstable; urgency=low
++
++ * New upstream release
++
++ -- Simon Richter <sjr at debian.org> Wed, 26 Jun 2013 10:06:29 +0200
++
+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 eac26c2,0000000..86a7636
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,74 -1,0 +1,74 @@@
- From 0ee7f97ae6e740ea80766d1126e0520583780d40 Mon Sep 17 00:00:00 2001
++From 5e262194e009622863d19e17c03cb44d0bd066cc 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/8] Add vector argument test case.
++Subject: [PATCH 01/11] 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
+
+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;
++}
+diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
- index c009d99..ea23f31 100644
++index df59feb..8a58ff4 100644
+--- a/utests/CMakeLists.txt
++++ b/utests/CMakeLists.txt
- @@ -37,6 +37,7 @@ set (utests_sources
++@@ -39,6 +39,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
+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 c385dcd,0000000..bbda3ac
mode 100644,000000..100644
--- 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
@@@ -1,95 -1,0 +1,95 @@@
- From 3d63c833d4ebcafe8e46fd498c28d08cb3046cd8 Mon Sep 17 00:00:00 2001
++From 116a8feff8b91cedcd91c087ba14214ee939fa72 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
++Subject: [PATCH 02/11] 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-Add-the-builtin-function-abs-and-the-according-test-.patch
index 0f0490e,0000000..cf69bf3
mode 100644,000000..100644
--- a/debian/patches/0003-Add-the-builtin-function-abs-and-the-according-test-.patch
+++ b/debian/patches/0003-Add-the-builtin-function-abs-and-the-according-test-.patch
@@@ -1,218 -1,0 +1,219 @@@
- From 1e18f92e0de249a403d4a50842d778d61a84d053 Mon Sep 17 00:00:00 2001
++From 8b61fc4c2644e6748577d30567e6e4bf70aef436 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
++Subject: [PATCH 03/11] 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
++index 8fb2a80..11efcb9 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));
++@@ -1224,7 +1224,15 @@ namespace gbe
++ const GenRegister dst = sel.selReg(insn.getDst(0), getType(opcode));
++ const GenRegister src = sel.selReg(insn.getSrc(0), getType(opcode));
+ 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
++index 08500ba..ba2192a 100644
+--- a/backend/src/llvm/llvm_gen_backend.cpp
++++ b/backend/src/llvm/llvm_gen_backend.cpp
- @@ -1686,6 +1686,7 @@ namespace gbe
++@@ -1688,6 +1688,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
++@@ -1844,6 +1845,13 @@ 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_FBH: this->emitUnaryCallInst(I,CS,ir::OP_FBH); break;
++ case GEN_OCL_FBL: this->emitUnaryCallInst(I,CS,ir::OP_FBL); 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;
++@@ -1852,7 +1860,7 @@ namespace gbe
+ 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
++index fe19844..89a04ea 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
++index 3b191ab..5ad829e 100644
+--- a/backend/src/ocl_stdlib.h
++++ b/backend/src/ocl_stdlib.h
- @@ -4337,6 +4337,7 @@ INLINE uint get_global_id(uint dim) {
++@@ -4401,6 +4401,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) {
++@@ -4579,6 +4580,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);
++@@ -4924,6 +4926,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
++@@ -5051,6 +5054,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
++index 8a58ff4..b75f3b4 100644
+--- a/utests/CMakeLists.txt
++++ b/utests/CMakeLists.txt
- @@ -32,6 +32,7 @@ set (utests_sources
++@@ -33,6 +33,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/0004-PATCH-Refine-the-get_local_id-.-builtins.patch
index 25f2f39,0000000..888ffc4
mode 100644,000000..100644
--- a/debian/patches/0004-PATCH-Refine-the-get_local_id-.-builtins.patch
+++ b/debian/patches/0004-PATCH-Refine-the-get_local_id-.-builtins.patch
@@@ -1,55 -1,0 +1,55 @@@
- From 8a2826a9a7b2085935f5396ae633645d35255b17 Mon Sep 17 00:00:00 2001
++From 3c0ee8cab035426e7e28425d2142e15051d83c90 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.
++Subject: [PATCH 04/11] [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
++index 5ad829e..227454d 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)
++@@ -4379,11 +4379,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
++index f7db4bc..bb09c07 100644
+--- a/src/cl_api.c
++++ b/src/cl_api.c
- @@ -1570,7 +1570,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
++@@ -1581,7 +1581,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/0005-Add-atomic-help-functions.-Support-global-and-local-.patch
index 0000000,0000000..f011afd
new file mode 100644
--- /dev/null
+++ b/debian/patches/0005-Add-atomic-help-functions.-Support-global-and-local-.patch
@@@ -1,0 -1,0 +1,611 @@@
++From 6e70caed68eb0ace38a8b42b105c0fd4c38e8f9e Mon Sep 17 00:00:00 2001
++From: Yang Rong <rong.r.yang at intel.com>
++Date: Wed, 26 Jun 2013 15:29:21 +0800
++Subject: [PATCH 05/11] Add atomic help functions. Support global and local
++ buffer.
++To: beignet at lists.freedesktop.org
++
++Signed-off-by: Yang Rong <rong.r.yang at intel.com>
++---
++ backend/src/backend/gen_context.cpp | 9 +++
++ backend/src/backend/gen_context.hpp | 1 +
++ backend/src/backend/gen_defs.hpp | 49 ++++++++++++--
++ backend/src/backend/gen_encoder.cpp | 35 ++++++++++
++ backend/src/backend/gen_encoder.hpp | 2 +
++ .../src/backend/gen_insn_gen7_schedule_info.hxx | 2 +-
++ backend/src/backend/gen_insn_selection.cpp | 61 +++++++++++++++++
++ backend/src/backend/gen_insn_selection.hpp | 4 +-
++ backend/src/backend/gen_insn_selection.hxx | 1 +
++ backend/src/ir/context.hpp | 6 ++
++ backend/src/ir/instruction.cpp | 70 +++++++++++++++++++-
++ backend/src/ir/instruction.hpp | 43 +++++++++++-
++ backend/src/ir/instruction.hxx | 13 ++++
++ 13 files changed, 283 insertions(+), 13 deletions(-)
++
++diff --git a/backend/src/backend/gen_context.cpp b/backend/src/backend/gen_context.cpp
++index 93d3932..41cab90 100644
++--- a/backend/src/backend/gen_context.cpp
+++++ b/backend/src/backend/gen_context.cpp
++@@ -226,6 +226,15 @@ namespace gbe
++ }
++ }
++
+++ void GenContext::emitAtomicInstruction(const SelectionInstruction &insn) {
+++ const GenRegister src = ra->genReg(insn.src(0));
+++ const GenRegister dst = ra->genReg(insn.dst(0));
+++ const uint32_t function = insn.extra.function;
+++ const uint32_t bti = insn.extra.elem;
+++
+++ p->ATOMIC(dst, function, src, bti, insn.srcNum);
+++ }
+++
++ void GenContext::emitIndirectMoveInstruction(const SelectionInstruction &insn) {
++ GenRegister src = ra->genReg(insn.src(0));
++ if(isScalarReg(src.reg()))
++diff --git a/backend/src/backend/gen_context.hpp b/backend/src/backend/gen_context.hpp
++index 804384d..5dfaef9 100644
++--- a/backend/src/backend/gen_context.hpp
+++++ b/backend/src/backend/gen_context.hpp
++@@ -91,6 +91,7 @@ namespace gbe
++ void emitWriteFloat64Instruction(const SelectionInstruction &insn);
++ void emitUntypedReadInstruction(const SelectionInstruction &insn);
++ void emitUntypedWriteInstruction(const SelectionInstruction &insn);
+++ void emitAtomicInstruction(const SelectionInstruction &insn);
++ void emitByteGatherInstruction(const SelectionInstruction &insn);
++ void emitByteScatterInstruction(const SelectionInstruction &insn);
++ void emitSampleInstruction(const SelectionInstruction &insn);
++diff --git a/backend/src/backend/gen_defs.hpp b/backend/src/backend/gen_defs.hpp
++index d1ce6b2..61412c4 100644
++--- a/backend/src/backend/gen_defs.hpp
+++++ b/backend/src/backend/gen_defs.hpp
++@@ -1,4 +1,4 @@
++-/*
+++/*
++ * Copyright © 2012 Intel Corporation
++ *
++ * This library is free software; you can redistribute it and/or
++@@ -21,7 +21,7 @@
++ Copyright (C) Intel Corp. 2006. All Rights Reserved.
++ Intel funded Tungsten Graphics (http://www.tungstengraphics.com) to
++ develop this 3D driver.
++-
+++
++ Permission is hereby granted, free of charge, to any person obtaining
++ a copy of this software and associated documentation files (the
++ "Software"), to deal in the Software without restriction, including
++@@ -29,11 +29,11 @@
++ distribute, sublicense, and/or sell copies of the Software, and to
++ permit persons to whom the Software is furnished to do so, subject to
++ the following conditions:
++-
+++
++ The above copyright notice and this permission notice (including the
++ next paragraph) shall be included in all copies or substantial
++ portions of the Software.
++-
+++
++ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
++ EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
++ MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
++@@ -41,7 +41,7 @@
++ LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION
++ OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION
++ WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
++-
+++
++ **********************************************************************/
++ /*
++ * Authors:
++@@ -169,6 +169,28 @@ enum opcode {
++ GEN_OPCODE_NOP = 126,
++ };
++
+++#define GEN_ATOMIC_SIMD16 0
+++#define GEN_ATOMIC_SIMD8 1
+++
+++enum GenAtomicOpCode {
+++ GEN_ATOMIC_OP_CMPWR8B = 0,
+++ GEN_ATOMIC_OP_AND = 1,
+++ GEN_ATOMIC_OP_OR = 2,
+++ GEN_ATOMIC_OP_XOR = 3,
+++ GEN_ATOMIC_OP_MOV = 4,
+++ GEN_ATOMIC_OP_INC = 5,
+++ GEN_ATOMIC_OP_DEC = 6,
+++ GEN_ATOMIC_OP_ADD = 7,
+++ GEN_ATOMIC_OP_SUB = 8,
+++ GEN_ATOMIC_OP_REVSUB = 9,
+++ GEN_ATOMIC_OP_IMAX = 10,
+++ GEN_ATOMIC_OP_IMIN = 11,
+++ GEN_ATOMIC_OP_UMAX = 12,
+++ GEN_ATOMIC_OP_UMIN = 13,
+++ GEN_ATOMIC_OP_CMPWR = 14,
+++ GEN_ATOMIC_OP_PREDEC = 15
+++};
+++
++ /*! Gen SFID */
++ enum GenMessageTarget {
++ GEN_SFID_NULL = 0,
++@@ -772,7 +794,7 @@ struct GenInstruction
++ /*! Memory fence */
++ struct {
++ uint32_t bti:8;
++- uint32_t ingored:5;
+++ uint32_t pad:5;
++ uint32_t commit_enable:1;
++ uint32_t msg_type:4;
++ uint32_t pad2:1;
++@@ -783,6 +805,21 @@ struct GenInstruction
++ uint32_t end_of_thread:1;
++ } gen7_memory_fence;
++
+++ /*! atomic messages */
+++ struct {
+++ uint32_t bti:8;
+++ uint32_t aop_type:4;
+++ uint32_t simd_mode:1;
+++ uint32_t return_data:1;
+++ uint32_t msg_type:4;
+++ uint32_t category:1;
+++ uint32_t header_present:1;
+++ uint32_t response_length:5;
+++ uint32_t msg_length:4;
+++ uint32_t pad3:2;
+++ uint32_t end_of_thread:1;
+++ } gen7_atomic_op;
+++
++ struct {
++ uint32_t src1_subreg_nr_high:1;
++ uint32_t src1_reg_nr:8;
++diff --git a/backend/src/backend/gen_encoder.cpp b/backend/src/backend/gen_encoder.cpp
++index e96678b..43658e8 100644
++--- a/backend/src/backend/gen_encoder.cpp
+++++ b/backend/src/backend/gen_encoder.cpp
++@@ -558,6 +558,41 @@ namespace gbe
++ response_length);
++ }
++
+++ void GenEncoder::ATOMIC(GenRegister dst, uint32_t function, GenRegister src, uint32_t bti, uint32_t srcNum) {
+++ GenInstruction *insn = this->next(GEN_OPCODE_SEND);
+++ uint32_t msg_length = 0;
+++ uint32_t response_length = 0;
+++
+++ if (this->curr.execWidth == 8) {
+++ msg_length = srcNum;
+++ response_length = 1;
+++ } else if (this->curr.execWidth == 16) {
+++ msg_length = 2*srcNum;
+++ response_length = 2;
+++ } else
+++ NOT_IMPLEMENTED;
+++
+++ this->setHeader(insn);
+++ this->setDst(insn, GenRegister::uw16grf(dst.nr, 0));
+++ this->setSrc0(insn, GenRegister::ud8grf(src.nr, 0));
+++ this->setSrc1(insn, GenRegister::immud(0));
+++
+++ const GenMessageTarget sfid = GEN_SFID_DATAPORT_DATA_CACHE;
+++ setMessageDescriptor(this, insn, sfid, msg_length, response_length);
+++ insn->bits3.gen7_atomic_op.msg_type = GEN_UNTYPED_ATOMIC_READ;
+++ insn->bits3.gen7_atomic_op.bti = bti;
+++ insn->bits3.gen7_atomic_op.return_data = 1;
+++ insn->bits3.gen7_atomic_op.aop_type = function;
+++
+++ if (this->curr.execWidth == 8)
+++ insn->bits3.gen7_atomic_op.simd_mode = GEN_ATOMIC_SIMD8;
+++ else if (this->curr.execWidth == 16)
+++ insn->bits3.gen7_atomic_op.simd_mode = GEN_ATOMIC_SIMD16;
+++ else
+++ NOT_SUPPORTED;
+++
+++ }
+++
++ GenInstruction *GenEncoder::next(uint32_t opcode) {
++ GenInstruction insn;
++ std::memset(&insn, 0, sizeof(GenInstruction));
++diff --git a/backend/src/backend/gen_encoder.hpp b/backend/src/backend/gen_encoder.hpp
++index 88a3e77..3ff8c97 100644
++--- a/backend/src/backend/gen_encoder.hpp
+++++ b/backend/src/backend/gen_encoder.hpp
++@@ -136,6 +136,8 @@ namespace gbe
++ void NOP(void);
++ /*! Wait instruction (used for the barrier) */
++ void WAIT(void);
+++ /*! Atomic instructions */
+++ void ATOMIC(GenRegister dst, uint32_t function, GenRegister src, uint32_t bti, uint32_t srcNum);
++ /*! Read 64-bits float arrays */
++ void READ_FLOAT64(GenRegister dst, GenRegister src, uint32_t bti, uint32_t elemNum);
++ /*! Write 64-bits float arrays */
++diff --git a/backend/src/backend/gen_insn_gen7_schedule_info.hxx b/backend/src/backend/gen_insn_gen7_schedule_info.hxx
++index a3b4621..f3f4a25 100644
++--- a/backend/src/backend/gen_insn_gen7_schedule_info.hxx
+++++ b/backend/src/backend/gen_insn_gen7_schedule_info.hxx
++@@ -21,4 +21,4 @@ DECL_GEN7_SCHEDULE(ByteScatter, 80, 1, 1)
++ DECL_GEN7_SCHEDULE(Sample, 80, 1, 1)
++ DECL_GEN7_SCHEDULE(TypedWrite, 80, 1, 1)
++ DECL_GEN7_SCHEDULE(GetImageInfo, 20, 4, 2)
++-
+++DECL_GEN7_SCHEDULE(Atomic, 80, 1, 1)
++diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
++index 11efcb9..c64afd9 100644
++--- a/backend/src/backend/gen_insn_selection.cpp
+++++ b/backend/src/backend/gen_insn_selection.cpp
++@@ -168,12 +168,14 @@ namespace gbe
++ bool SelectionInstruction::isRead(void) const {
++ return this->opcode == SEL_OP_UNTYPED_READ ||
++ this->opcode == SEL_OP_READ_FLOAT64 ||
+++ this->opcode == SEL_OP_ATOMIC ||
++ this->opcode == SEL_OP_BYTE_GATHER;
++ }
++
++ bool SelectionInstruction::isWrite(void) const {
++ return this->opcode == SEL_OP_UNTYPED_WRITE ||
++ this->opcode == SEL_OP_WRITE_FLOAT64 ||
+++ this->opcode == SEL_OP_ATOMIC ||
++ this->opcode == SEL_OP_BYTE_SCATTER;
++ }
++
++@@ -456,6 +458,8 @@ namespace gbe
++ void NOP(void);
++ /*! Wait instruction (used for the barrier) */
++ void WAIT(void);
+++ /*! Atomic instruction */
+++ void ATOMIC(Reg dst, uint32_t function, uint32_t srcNum, Reg src0, Reg src1, Reg src2, uint32_t bti);
++ /*! Read 64 bits float array */
++ void READ_FLOAT64(Reg addr, const GenRegister *dst, uint32_t elemNum, uint32_t bti);
++ /*! Write 64 bits float array */
++@@ -730,6 +734,23 @@ namespace gbe
++ insn->src(0) = src;
++ }
++
+++ void Selection::Opaque::ATOMIC(Reg dst, uint32_t function,
+++ uint32_t srcNum, Reg src0,
+++ Reg src1, Reg src2, uint32_t bti) {
+++ SelectionInstruction *insn = this->appendInsn(SEL_OP_ATOMIC, 1, srcNum);
+++ insn->dst(0) = dst;
+++ insn->src(0) = src0;
+++ if(srcNum > 1) insn->src(1) = src1;
+++ if(srcNum > 2) insn->src(2) = src2;
+++ insn->extra.function = function;
+++ insn->extra.elem = bti;
+++ SelectionVector *vector = this->appendVector();
+++
+++ vector->regNum = srcNum;
+++ vector->reg = &insn->src(0);
+++ vector->isSrc = 1;
+++ }
+++
++ void Selection::Opaque::EOT(void) { this->appendInsn(SEL_OP_EOT, 0, 0); }
++ void Selection::Opaque::NOP(void) { this->appendInsn(SEL_OP_NOP, 0, 0); }
++ void Selection::Opaque::WAIT(void) { this->appendInsn(SEL_OP_WAIT, 0, 0); }
++@@ -916,6 +937,7 @@ namespace gbe
++ bool Selection::Opaque::isRoot(const ir::Instruction &insn) const {
++ if (insn.getDstNum() > 1 ||
++ insn.hasSideEffect() ||
+++ insn.isMemberOf<ir::AtomicInstruction>() ||
++ insn.isMemberOf<ir::BranchInstruction>() ||
++ insn.isMemberOf<ir::LabelInstruction>())
++ return true;
++@@ -2033,6 +2055,44 @@ namespace gbe
++ DECL_CTOR(ConvertInstruction, 1, 1);
++ };
++
+++ /*! Convert instruction pattern */
+++ DECL_PATTERN(AtomicInstruction)
+++ {
+++ INLINE bool emitOne(Selection::Opaque &sel, const ir::AtomicInstruction &insn) const
+++ {
+++ using namespace ir;
+++ const Opcode opcode = insn.getOpcode();
+++ const AddressSpace space = insn.getAddressSpace();
+++ const uint32_t bti = space == MEM_LOCAL ? 0xfe : 0x01;
+++ const uint32_t srcNum = insn.getSrcNum();
+++ const GenRegister src0 = sel.selReg(insn.getSrc(0), TYPE_U32); //address
+++ GenRegister src1 = src0, src2 = src0;
+++ if(srcNum > 1) src1 = sel.selReg(insn.getSrc(1), TYPE_U32);
+++ if(srcNum > 2) src2 = sel.selReg(insn.getSrc(2), TYPE_U32);
+++ GenRegister dst = sel.selReg(insn.getDst(0), TYPE_U32);
+++ GenAtomicOpCode aop = GEN_ATOMIC_OP_CMPWR8B;
+++ switch (opcode) {
+++ case OP_ATOMIC_ADD: aop = GEN_ATOMIC_OP_ADD; break;
+++ case OP_ATOMIC_SUB: aop = GEN_ATOMIC_OP_SUB; break;
+++ case OP_ATOMIC_AND: aop = GEN_ATOMIC_OP_AND; break;
+++ case OP_ATOMIC_OR : aop = GEN_ATOMIC_OP_OR; break;
+++ case OP_ATOMIC_XOR: aop = GEN_ATOMIC_OP_XOR; break;
+++ case OP_ATOMIC_XCHG: aop = GEN_ATOMIC_OP_MOV; break;
+++ case OP_ATOMIC_UMIN: aop = GEN_ATOMIC_OP_UMIN; break;
+++ case OP_ATOMIC_UMAX: aop = GEN_ATOMIC_OP_UMAX; break;
+++ case OP_ATOMIC_IMIN: aop = GEN_ATOMIC_OP_IMIN; break;
+++ case OP_ATOMIC_IMAX: aop = GEN_ATOMIC_OP_IMAX; break;
+++ case OP_ATOMIC_INC: aop = GEN_ATOMIC_OP_INC; break;
+++ case OP_ATOMIC_DEC: aop = GEN_ATOMIC_OP_DEC; break;
+++ case OP_ATOMIC_CMPXCHG: aop = GEN_ATOMIC_OP_CMPWR; break;
+++ default: NOT_SUPPORTED;
+++ }
+++ sel.ATOMIC(dst, aop, srcNum, src0, src1, src2, bti);
+++ return true;
+++ }
+++ DECL_CTOR(AtomicInstruction, 1, 1);
+++ };
+++
++ /*! Select instruction pattern */
++ class SelectInstructionPattern : public SelectionPattern
++ {
++@@ -2379,6 +2439,7 @@ namespace gbe
++ this->insert<SelectInstructionPattern>();
++ this->insert<CompareInstructionPattern>();
++ this->insert<ConvertInstructionPattern>();
+++ this->insert<AtomicInstructionPattern>();
++ this->insert<LabelInstructionPattern>();
++ this->insert<BranchInstructionPattern>();
++ this->insert<Int32x32MulInstructionPattern>();
++diff --git a/backend/src/backend/gen_insn_selection.hpp b/backend/src/backend/gen_insn_selection.hpp
++index 778eb1f..5ae6e42 100644
++--- a/backend/src/backend/gen_insn_selection.hpp
+++++ b/backend/src/backend/gen_insn_selection.hpp
++@@ -92,9 +92,9 @@ namespace gbe
++ GenInstructionState state;
++ union {
++ struct {
++- /*! Store bti for loads/stores and function for math and compares */
+++ /*! Store bti for loads/stores and function for math, atomic and compares */
++ uint16_t function:8;
++- /*! elemSize for byte scatters / gathers, elemNum for untyped msg */
+++ /*! elemSize for byte scatters / gathers, elemNum for untyped msg, bti for atomic */
++ uint16_t elem:8;
++ };
++ struct {
++diff --git a/backend/src/backend/gen_insn_selection.hxx b/backend/src/backend/gen_insn_selection.hxx
++index cc2be08..f1a4701 100644
++--- a/backend/src/backend/gen_insn_selection.hxx
+++++ b/backend/src/backend/gen_insn_selection.hxx
++@@ -20,6 +20,7 @@ DECL_SELECTION_IR(RSL, BinaryInstruction)
++ DECL_SELECTION_IR(ASR, BinaryInstruction)
++ DECL_SELECTION_IR(ADD, BinaryInstruction)
++ DECL_SELECTION_IR(MUL, BinaryInstruction)
+++DECL_SELECTION_IR(ATOMIC, AtomicInstruction)
++ DECL_SELECTION_IR(MACH, BinaryInstruction)
++ DECL_SELECTION_IR(CMP, CompareInstruction)
++ DECL_SELECTION_IR(SEL_CMP, CompareInstruction)
++diff --git a/backend/src/ir/context.hpp b/backend/src/ir/context.hpp
++index c286f1d..55e76f2 100644
++--- a/backend/src/ir/context.hpp
+++++ b/backend/src/ir/context.hpp
++@@ -150,6 +150,12 @@ namespace ir {
++ this->append(insn);
++ }
++
+++ /*! For all atomic functions */
+++ void ATOMIC(Opcode opcode, Register dst, AddressSpace space, Tuple src) {
+++ const Instruction insn = gbe::ir::ATOMIC(opcode, dst, space, src);
+++ this->append(insn);
+++ }
+++
++ /*! LOAD with the destinations directly specified */
++ template <typename... Args>
++ void LOAD(Type type, Register offset, AddressSpace space, bool dwAligned, Args...values)
++diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp
++index 67a4c12..578f5d2 100644
++--- a/backend/src/ir/instruction.cpp
+++++ b/backend/src/ir/instruction.cpp
++@@ -1,4 +1,4 @@
++-/*
+++/*
++ * Copyright © 2012 Intel Corporation
++ *
++ * This library is free software; you can redistribute it and/or
++@@ -56,7 +56,7 @@ namespace ir {
++ };
++
++ /*! For regular n source instructions */
++- template <typename T, uint32_t srcNum>
+++ template <typename T, uint32_t srcNum>
++ struct NSrcPolicy {
++ INLINE uint32_t getSrcNum(void) const { return srcNum; }
++ INLINE Register getSrc(const Function &fn, uint32_t ID) const {
++@@ -246,6 +246,37 @@ namespace ir {
++ Type srcType; //!< Type to convert from
++ };
++
+++ class ALIGNED_INSTRUCTION AtomicInstruction :
+++ public BasePolicy,
+++ public TupleSrcPolicy<AtomicInstruction>,
+++ public NDstPolicy<AtomicInstruction, 1>
+++ {
+++ public:
+++ AtomicInstruction(Opcode opcode,
+++ Register dst,
+++ AddressSpace addrSpace,
+++ Tuple src)
+++ {
+++ this->opcode = opcode;
+++ this->dst[0] = dst;
+++ this->src = src;
+++ this->addrSpace = addrSpace;
+++ srcNum = 2;
+++ if((opcode == OP_ATOMIC_INC) ||
+++ (opcode == OP_ATOMIC_DEC))
+++ srcNum = 1;
+++ if(opcode == OP_ATOMIC_CMPXCHG)
+++ srcNum = 3;
+++ }
+++ INLINE AddressSpace getAddressSpace(void) const { return this->addrSpace; }
+++ INLINE bool wellFormed(const Function &fn, std::string &whyNot) const;
+++ INLINE void out(std::ostream &out, const Function &fn) const;
+++ Register dst[1];
+++ Tuple src;
+++ AddressSpace addrSpace; //!< Address space
+++ uint8_t srcNum; //!<Source Number
+++ };
+++
++ class ALIGNED_INSTRUCTION BranchInstruction :
++ public BasePolicy,
++ public NDstPolicy<BranchInstruction, 0>
++@@ -738,6 +769,20 @@ namespace ir {
++ return true;
++ }
++
+++ // We can convert anything to anything, but types and families must match
+++ INLINE bool AtomicInstruction::wellFormed(const Function &fn, std::string &whyNot) const
+++ {
+++ if (UNLIKELY(checkSpecialRegForWrite(dst[0], fn, whyNot) == false))
+++ return false;
+++ if (UNLIKELY(checkRegisterData(FAMILY_DWORD, dst[0], fn, whyNot) == false))
+++ return false;
+++ for (uint32_t srcID = 0; srcID < srcNum; ++srcID)
+++ if (UNLIKELY(checkRegisterData(FAMILY_DWORD, getSrc(fn, srcID), fn, whyNot) == false))
+++ return false;
+++
+++ return true;
+++ }
+++
++ /*! Loads and stores follow the same restrictions */
++ template <typename T>
++ INLINE bool wellFormedLoadStore(const T &insn, const Function &fn, std::string &whyNot)
++@@ -883,6 +928,15 @@ namespace ir {
++ ternaryOrSelectOut(*this, out, fn);
++ }
++
+++ INLINE void AtomicInstruction::out(std::ostream &out, const Function &fn) const {
+++ this->outOpcode(out);
+++ out << "." << addrSpace;
+++ out << " %" << this->getDst(fn, 0);
+++ out << " {" << "%" << this->getSrc(fn, 0) << "}";
+++ for (uint32_t i = 1; i < srcNum; ++i)
+++ out << " %" << this->getSrc(fn, i);
+++ }
+++
++ INLINE void ConvertInstruction::out(std::ostream &out, const Function &fn) const {
++ this->outOpcode(out);
++ out << "." << this->getDstType()
++@@ -1009,6 +1063,10 @@ START_INTROSPECTION(ConvertInstruction)
++ #include "ir/instruction.hxx"
++ END_INTROSPECTION(ConvertInstruction)
++
+++START_INTROSPECTION(AtomicInstruction)
+++#include "ir/instruction.hxx"
+++END_INTROSPECTION(AtomicInstruction)
+++
++ START_INTROSPECTION(SelectInstruction)
++ #include "ir/instruction.hxx"
++ END_INTROSPECTION(SelectInstruction)
++@@ -1180,7 +1238,7 @@ END_FUNCTION(Instruction, Register)
++ }
++
++ bool Instruction::hasSideEffect(void) const {
++- return opcode == OP_STORE ||
+++ return opcode == OP_STORE ||
++ opcode == OP_TYPED_WRITE ||
++ opcode == OP_SYNC;
++ }
++@@ -1197,6 +1255,7 @@ DECL_MEM_FN(SelectInstruction, Type, getType(void), getType())
++ DECL_MEM_FN(CompareInstruction, Type, getType(void), getType())
++ DECL_MEM_FN(ConvertInstruction, Type, getSrcType(void), getSrcType())
++ DECL_MEM_FN(ConvertInstruction, Type, getDstType(void), getDstType())
+++DECL_MEM_FN(AtomicInstruction, AddressSpace, getAddressSpace(void), getAddressSpace())
++ DECL_MEM_FN(StoreInstruction, Type, getValueType(void), getValueType())
++ DECL_MEM_FN(StoreInstruction, uint32_t, getValueNum(void), getValueNum())
++ DECL_MEM_FN(StoreInstruction, AddressSpace, getAddressSpace(void), getAddressSpace())
++@@ -1304,6 +1363,11 @@ DECL_MEM_FN(GetImageInfoInstruction, uint32_t, getInfoType(void), getInfoType())
++ return internal::ConvertInstruction(dstType, srcType, dst, src).convert();
++ }
++
+++ // For all unary functions with given opcode
+++ Instruction ATOMIC(Opcode opcode, Register dst, AddressSpace space, Tuple src) {
+++ return internal::AtomicInstruction(opcode, dst, space, src).convert();
+++ }
+++
++ // BRA
++ Instruction BRA(LabelIndex labelIndex) {
++ return internal::BranchInstruction(OP_BRA, labelIndex).convert();
++diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp
++index 0f3bd34..91d280d 100644
++--- a/backend/src/ir/instruction.hpp
+++++ b/backend/src/ir/instruction.hpp
++@@ -1,4 +1,4 @@
++-/*
+++/*
++ * Copyright © 2012 Intel Corporation
++ *
++ * This library is free software; you can redistribute it and/or
++@@ -228,6 +228,19 @@ namespace ir {
++ static bool isClassOf(const Instruction &insn);
++ };
++
+++ /*! Atomic instruction */
+++ class AtomicInstruction : public Instruction {
+++ public:
+++ /*! Where the address register goes */
+++ static const uint32_t addressIndex = 0;
+++ /*! Address space that is manipulated here */
+++ AddressSpace getAddressSpace(void) const;
+++ /*! Return the register that contains the addresses */
+++ INLINE Register getAddress(void) const { return this->getSrc(addressIndex); }
+++ /*! Return true if the given instruction is an instance of this class */
+++ static bool isClassOf(const Instruction &insn);
+++ };
+++
++ /*! Store instruction. First source is the address. Next sources are the
++ * values to store contiguously at the given address
++ */
++@@ -555,6 +568,34 @@ namespace ir {
++ Instruction GT(Type type, Register dst, Register src0, Register src1);
++ /*! cvt.{dstType <- srcType} dst src */
++ Instruction CVT(Type dstType, Type srcType, Register dst, Register src);
+++ /*! atomic dst addr.space {src1 {src2}} */
+++ Instruction ATOMIC(Opcode opcode, Register dst, AddressSpace space, Tuple src);
+++ /*! atomic_add dst addr.space src1 */
+++ Instruction ATOMIC_ADD(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_sub dst addr.space src1 */
+++ Instruction ATOMIC_SUB(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_and dst addr.space src1 */
+++ Instruction ATOMIC_AND(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_or dst addr.space src1 */
+++ Instruction ATOMIC_OR(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_xor dst addr.space src1 */
+++ Instruction ATOMIC_XOR(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_xchg dst addr.space src1 */
+++ Instruction ATOMIC_XCHG(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_imin dst addr.space src1 */
+++ Instruction ATOMIC_IMIN(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_imax dst addr.space src1 */
+++ Instruction ATOMIC_IMAX(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_umin dst addr.space src1 */
+++ Instruction ATOMIC_UMIN(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_umax dst addr.space src1 */
+++ Instruction ATOMIC_UMAX(Register dst, Register addr, AddressSpace space, Register src1);
+++ /*! atomic_inc dst addr.space */
+++ Instruction ATOMIC_INC(Register dst, Register addr, AddressSpace space);
+++ /*! atomic_dec dst addr.space */
+++ Instruction ATOMIC_DEC(Register dst, Register addr, AddressSpace space);
+++ /*! atomic_cmpxchg dst addr.space src1 src2 */
+++ Instruction ATOMIC_CMPXCHG(Register dst, Register addr, AddressSpace space, Register src1, Register src2);
++ /*! bra labelIndex */
++ Instruction BRA(LabelIndex labelIndex);
++ /*! (pred) bra labelIndex */
++diff --git a/backend/src/ir/instruction.hxx b/backend/src/ir/instruction.hxx
++index acfb45a..42c1e89 100644
++--- a/backend/src/ir/instruction.hxx
+++++ b/backend/src/ir/instruction.hxx
++@@ -61,6 +61,19 @@ DECL_INSN(LT, CompareInstruction)
++ DECL_INSN(GE, CompareInstruction)
++ DECL_INSN(GT, CompareInstruction)
++ DECL_INSN(CVT, ConvertInstruction)
+++DECL_INSN(ATOMIC_ADD, AtomicInstruction)
+++DECL_INSN(ATOMIC_SUB, AtomicInstruction)
+++DECL_INSN(ATOMIC_AND, AtomicInstruction)
+++DECL_INSN(ATOMIC_OR, AtomicInstruction)
+++DECL_INSN(ATOMIC_XOR, AtomicInstruction)
+++DECL_INSN(ATOMIC_XCHG, AtomicInstruction)
+++DECL_INSN(ATOMIC_UMIN, AtomicInstruction)
+++DECL_INSN(ATOMIC_UMAX, AtomicInstruction)
+++DECL_INSN(ATOMIC_IMIN, AtomicInstruction)
+++DECL_INSN(ATOMIC_IMAX, AtomicInstruction)
+++DECL_INSN(ATOMIC_INC, AtomicInstruction)
+++DECL_INSN(ATOMIC_DEC, AtomicInstruction)
+++DECL_INSN(ATOMIC_CMPXCHG, AtomicInstruction)
++ DECL_INSN(BRA, BranchInstruction)
++ DECL_INSN(RET, BranchInstruction)
++ DECL_INSN(LOADI, LoadImmInstruction)
++--
++1.7.10.4
++
diff --cc debian/patches/0006-Add-all-atomic-built-in-functions-support.patch
index 0000000,0000000..fac6999
new file mode 100644
--- /dev/null
+++ b/debian/patches/0006-Add-all-atomic-built-in-functions-support.patch
@@@ -1,0 -1,0 +1,289 @@@
++From 6d471350375328273070c9b73b22e9ab042c3313 Mon Sep 17 00:00:00 2001
++From: Yang Rong <rong.r.yang at intel.com>
++Date: Wed, 26 Jun 2013 15:29:22 +0800
++Subject: [PATCH 06/11] Add all atomic built-in functions support.
++To: beignet at lists.freedesktop.org
++
++Signed-off-by: Yang Rong <rong.r.yang at intel.com>
++---
++ backend/src/llvm/llvm_gen_backend.cpp | 76 +++++++++++++++++++++
++ backend/src/llvm/llvm_gen_ocl_function.hxx | 28 ++++++++
++ backend/src/ocl_stdlib.h | 100 +++++++++++++++++++++++++++-
++ 3 files changed, 203 insertions(+), 1 deletion(-)
++
++diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
++index ba2192a..5db3fbe 100644
++--- a/backend/src/llvm/llvm_gen_backend.cpp
+++++ b/backend/src/llvm/llvm_gen_backend.cpp
++@@ -534,6 +534,8 @@ namespace gbe
++
++ // Emit unary instructions from gen native function
++ void emitUnaryCallInst(CallInst &I, CallSite &CS, ir::Opcode opcode);
+++ // Emit unary instructions from gen native function
+++ void emitAtomicInst(CallInst &I, CallSite &CS, ir::Opcode opcode);
++
++ // These instructions are not supported at all
++ void visitVAArgInst(VAArgInst &I) {NOT_SUPPORTED;}
++@@ -693,10 +695,12 @@ namespace gbe
++ return doIt(uint64_t(0));
++ }
++ }
+++
++ // NULL pointers
++ if(isa<ConstantPointerNull>(CPV)) {
++ return doIt(uint32_t(0));
++ }
+++
++ // Floats and doubles
++ const Type::TypeID typeID = CPV->getType()->getTypeID();
++ switch (typeID) {
++@@ -1698,6 +1702,32 @@ namespace gbe
++ case GEN_OCL_GET_IMAGE_CHANNEL_DATA_TYPE:
++ case GEN_OCL_GET_IMAGE_CHANNEL_ORDER:
++ case GEN_OCL_GET_IMAGE_DEPTH:
+++ case GEN_OCL_ATOMIC_ADD0:
+++ case GEN_OCL_ATOMIC_ADD1:
+++ case GEN_OCL_ATOMIC_SUB0:
+++ case GEN_OCL_ATOMIC_SUB1:
+++ case GEN_OCL_ATOMIC_AND0:
+++ case GEN_OCL_ATOMIC_AND1:
+++ case GEN_OCL_ATOMIC_OR0:
+++ case GEN_OCL_ATOMIC_OR1:
+++ case GEN_OCL_ATOMIC_XOR0:
+++ case GEN_OCL_ATOMIC_XOR1:
+++ case GEN_OCL_ATOMIC_XCHG0:
+++ case GEN_OCL_ATOMIC_XCHG1:
+++ case GEN_OCL_ATOMIC_UMAX0:
+++ case GEN_OCL_ATOMIC_UMAX1:
+++ case GEN_OCL_ATOMIC_UMIN0:
+++ case GEN_OCL_ATOMIC_UMIN1:
+++ case GEN_OCL_ATOMIC_IMAX0:
+++ case GEN_OCL_ATOMIC_IMAX1:
+++ case GEN_OCL_ATOMIC_IMIN0:
+++ case GEN_OCL_ATOMIC_IMIN1:
+++ case GEN_OCL_ATOMIC_INC0:
+++ case GEN_OCL_ATOMIC_INC1:
+++ case GEN_OCL_ATOMIC_DEC0:
+++ case GEN_OCL_ATOMIC_DEC1:
+++ case GEN_OCL_ATOMIC_CMPXCHG0:
+++ case GEN_OCL_ATOMIC_CMPXCHG1:
++ // No structure can be returned
++ this->newRegister(&I);
++ break;
++@@ -1782,6 +1812,26 @@ namespace gbe
++ ctx.ALU1(opcode, ir::TYPE_FLOAT, dst, src);
++ }
++
+++ void GenWriter::emitAtomicInst(CallInst &I, CallSite &CS, ir::Opcode opcode) {
+++ CallSite::arg_iterator AI = CS.arg_begin();
+++#if GBE_DEBUG
+++ CallSite::arg_iterator AE = CS.arg_end();
+++#endif /* GBE_DEBUG */
+++ GBE_ASSERT(AI != AE);
+++ unsigned int llvmSpace = (*AI)->getType()->getPointerAddressSpace();
+++ const ir::AddressSpace addrSpace = addressSpaceLLVMToGen(llvmSpace);
+++ const ir::Register dst = this->getRegister(&I);
+++
+++ vector<ir::Register> src;
+++ uint32_t srcNum = 0;
+++ while(AI != AE) {
+++ src.push_back(this->getRegister(*(AI++)));
+++ srcNum++;
+++ }
+++ const ir::Tuple srcTuple = ctx.arrayTuple(&src[0], srcNum);
+++ ctx.ATOMIC(opcode, dst, addrSpace, srcTuple);
+++ }
+++
++ void GenWriter::emitCallInst(CallInst &I) {
++ if (Function *F = I.getCalledFunction()) {
++ if (F->getIntrinsicID() != 0) {
++@@ -1870,6 +1920,32 @@ namespace gbe
++ case GEN_OCL_LBARRIER: ctx.SYNC(ir::syncLocalBarrier); break;
++ case GEN_OCL_GBARRIER: ctx.SYNC(ir::syncGlobalBarrier); break;
++ case GEN_OCL_LGBARRIER: ctx.SYNC(ir::syncLocalBarrier | ir::syncGlobalBarrier); break;
+++ case GEN_OCL_ATOMIC_ADD0:
+++ case GEN_OCL_ATOMIC_ADD1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_ADD); break;
+++ case GEN_OCL_ATOMIC_SUB0:
+++ case GEN_OCL_ATOMIC_SUB1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_SUB); break;
+++ case GEN_OCL_ATOMIC_AND0:
+++ case GEN_OCL_ATOMIC_AND1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_AND); break;
+++ case GEN_OCL_ATOMIC_OR0:
+++ case GEN_OCL_ATOMIC_OR1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_OR); break;
+++ case GEN_OCL_ATOMIC_XOR0:
+++ case GEN_OCL_ATOMIC_XOR1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_XOR); break;
+++ case GEN_OCL_ATOMIC_XCHG0:
+++ case GEN_OCL_ATOMIC_XCHG1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_XCHG); break;
+++ case GEN_OCL_ATOMIC_INC0:
+++ case GEN_OCL_ATOMIC_INC1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_INC); break;
+++ case GEN_OCL_ATOMIC_DEC0:
+++ case GEN_OCL_ATOMIC_DEC1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_DEC); break;
+++ case GEN_OCL_ATOMIC_UMIN0:
+++ case GEN_OCL_ATOMIC_UMIN1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_UMIN); break;
+++ case GEN_OCL_ATOMIC_UMAX0:
+++ case GEN_OCL_ATOMIC_UMAX1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_UMAX); break;
+++ case GEN_OCL_ATOMIC_IMIN0:
+++ case GEN_OCL_ATOMIC_IMIN1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_IMIN); break;
+++ case GEN_OCL_ATOMIC_IMAX0:
+++ case GEN_OCL_ATOMIC_IMAX1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_IMAX); break;
+++ case GEN_OCL_ATOMIC_CMPXCHG0:
+++ case GEN_OCL_ATOMIC_CMPXCHG1: this->emitAtomicInst(I,CS,ir::OP_ATOMIC_CMPXCHG); break;
++ case GEN_OCL_GET_IMAGE_WIDTH:
++ case GEN_OCL_GET_IMAGE_HEIGHT:
++ case GEN_OCL_GET_IMAGE_DEPTH:
++diff --git a/backend/src/llvm/llvm_gen_ocl_function.hxx b/backend/src/llvm/llvm_gen_ocl_function.hxx
++index 89a04ea..2f79690 100644
++--- a/backend/src/llvm/llvm_gen_ocl_function.hxx
+++++ b/backend/src/llvm/llvm_gen_ocl_function.hxx
++@@ -79,6 +79,34 @@ DECL_LLVM_GEN_FUNCTION(GET_IMAGE_DEPTH, __gen_ocl_get_image_depth)
++ DECL_LLVM_GEN_FUNCTION(GET_IMAGE_CHANNEL_DATA_TYPE, __gen_ocl_get_image_channel_data_type)
++ DECL_LLVM_GEN_FUNCTION(GET_IMAGE_CHANNEL_ORDER, __gen_ocl_get_image_channel_order)
++
+++// atomic related functions.
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_ADD0, _Z20__gen_ocl_atomic_addPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_ADD1, _Z20__gen_ocl_atomic_addPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_SUB0, _Z20__gen_ocl_atomic_subPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_SUB1, _Z20__gen_ocl_atomic_subPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_AND0, _Z20__gen_ocl_atomic_andPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_AND1, _Z20__gen_ocl_atomic_andPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_OR0, _Z19__gen_ocl_atomic_orPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_OR1, _Z19__gen_ocl_atomic_orPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_XOR0, _Z20__gen_ocl_atomic_xorPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_XOR1, _Z20__gen_ocl_atomic_xorPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_UMIN0, _Z21__gen_ocl_atomic_uminPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_UMIN1, _Z21__gen_ocl_atomic_uminPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_UMAX0, _Z21__gen_ocl_atomic_umaxPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_UMAX1, _Z21__gen_ocl_atomic_umaxPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_IMIN0, _Z21__gen_ocl_atomic_iminPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_IMIN1, _Z21__gen_ocl_atomic_iminPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_IMAX0, _Z21__gen_ocl_atomic_imaxPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_IMAX1, _Z21__gen_ocl_atomic_imaxPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_XCHG0, _Z21__gen_ocl_atomic_xchgPU3AS1jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_XCHG1, _Z21__gen_ocl_atomic_xchgPU3AS3jj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_INC0, _Z20__gen_ocl_atomic_incPU3AS1j)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_INC1, _Z20__gen_ocl_atomic_incPU3AS3j)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_DEC0, _Z20__gen_ocl_atomic_decPU3AS1j)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_DEC1, _Z20__gen_ocl_atomic_decPU3AS3j)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_CMPXCHG0, _Z24__gen_ocl_atomic_cmpxchgPU3AS1jjj)
+++DECL_LLVM_GEN_FUNCTION(ATOMIC_CMPXCHG1, _Z24__gen_ocl_atomic_cmpxchgPU3AS3jjj)
+++
++ // saturation related functions.
++ DECL_LLVM_GEN_FUNCTION(SADD_SAT_CHAR, _Z12ocl_sadd_satcc)
++ DECL_LLVM_GEN_FUNCTION(SADD_SAT_SHORT, _Z12ocl_sadd_satss)
++diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
++index 227454d..7a98e04 100644
++--- a/backend/src/ocl_stdlib.h
+++++ b/backend/src/ocl_stdlib.h
++@@ -1,4 +1,4 @@
++-/*
+++/*
++ * Copyright © 2012 Intel Corporation
++ *
++ * This library is free software; you can redistribute it and/or
++@@ -5100,6 +5100,104 @@ INLINE void write_mem_fence(cl_mem_fence_flags flags) {
++ }
++
++ /////////////////////////////////////////////////////////////////////////////
+++// Atomic functions
+++/////////////////////////////////////////////////////////////////////////////
+++OVERLOADABLE uint __gen_ocl_atomic_add(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_add(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_sub(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_sub(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_and(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_and(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_or(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_or(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_xor(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_xor(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_xchg(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_xchg(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_inc(__global uint *p);
+++OVERLOADABLE uint __gen_ocl_atomic_inc(__local uint *p);
+++OVERLOADABLE uint __gen_ocl_atomic_dec(__global uint *p);
+++OVERLOADABLE uint __gen_ocl_atomic_dec(__local uint *p);
+++OVERLOADABLE uint __gen_ocl_atomic_cmpxchg(__global uint *p, uint cmp, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_cmpxchg(__local uint *p, uint cmp, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_imin(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_imin(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_imax(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_imax(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_umin(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_umin(__local uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_umax(__global uint *p, uint val);
+++OVERLOADABLE uint __gen_ocl_atomic_umax(__local uint *p, uint val);
+++
+++#define DECL_ATOMIC_OP_SPACE(NAME, TYPE, SPACE, PREFIX) \
+++ INLINE_OVERLOADABLE TYPE atomic_##NAME (volatile SPACE TYPE *p, TYPE val) { \
+++ return (TYPE)__gen_ocl_##PREFIX##NAME((SPACE uint *)p, val); \
+++ }
+++
+++#define DECL_ATOMIC_OP_TYPE(NAME, TYPE, PREFIX) \
+++ DECL_ATOMIC_OP_SPACE(NAME, TYPE, __global, PREFIX) \
+++ DECL_ATOMIC_OP_SPACE(NAME, TYPE, __local, PREFIX) \
+++
+++#define DECL_ATOMIC_OP(NAME) \
+++ DECL_ATOMIC_OP_TYPE(NAME, uint, atomic_) \
+++ DECL_ATOMIC_OP_TYPE(NAME, int, atomic_)
+++
+++DECL_ATOMIC_OP(add)
+++DECL_ATOMIC_OP(sub)
+++DECL_ATOMIC_OP(and)
+++DECL_ATOMIC_OP(or)
+++DECL_ATOMIC_OP(xor)
+++DECL_ATOMIC_OP(xchg)
+++DECL_ATOMIC_OP_TYPE(xchg, float, atomic_)
+++DECL_ATOMIC_OP_TYPE(min, int, atomic_i)
+++DECL_ATOMIC_OP_TYPE(max, int, atomic_i)
+++DECL_ATOMIC_OP_TYPE(min, uint, atomic_u)
+++DECL_ATOMIC_OP_TYPE(max, uint, atomic_u)
+++
+++#undef DECL_ATOMIC_OP
+++#undef DECL_ATOMIC_OP_TYPE
+++#undef DECL_ATOMIC_OP_SPACE
+++
+++#define DECL_ATOMIC_OP_SPACE(NAME, TYPE, SPACE) \
+++ INLINE_OVERLOADABLE TYPE atomic_##NAME (volatile SPACE TYPE *p) { \
+++ return (TYPE)__gen_ocl_atomic_##NAME((SPACE uint *)p); \
+++ }
+++
+++#define DECL_ATOMIC_OP_TYPE(NAME, TYPE) \
+++ DECL_ATOMIC_OP_SPACE(NAME, TYPE, __global) \
+++ DECL_ATOMIC_OP_SPACE(NAME, TYPE, __local)
+++
+++#define DECL_ATOMIC_OP(NAME) \
+++ DECL_ATOMIC_OP_TYPE(NAME, uint) \
+++ DECL_ATOMIC_OP_TYPE(NAME, int)
+++
+++DECL_ATOMIC_OP(inc)
+++DECL_ATOMIC_OP(dec)
+++
+++#undef DECL_ATOMIC_OP
+++#undef DECL_ATOMIC_OP_TYPE
+++#undef DECL_ATOMIC_OP_SPACE
+++
+++#define DECL_ATOMIC_OP_SPACE(NAME, TYPE, SPACE) \
+++ INLINE_OVERLOADABLE TYPE atomic_##NAME (volatile SPACE TYPE *p, TYPE cmp, TYPE val) { \
+++ return (TYPE)__gen_ocl_atomic_##NAME((SPACE uint *)p, (uint)cmp, (uint)val); \
+++ }
+++
+++#define DECL_ATOMIC_OP_TYPE(NAME, TYPE) \
+++ DECL_ATOMIC_OP_SPACE(NAME, TYPE, __global) \
+++ DECL_ATOMIC_OP_SPACE(NAME, TYPE, __local)
+++
+++#define DECL_ATOMIC_OP(NAME) \
+++ DECL_ATOMIC_OP_TYPE(NAME, uint) \
+++ DECL_ATOMIC_OP_TYPE(NAME, int)
+++
+++DECL_ATOMIC_OP(cmpxchg)
+++
+++#undef DECL_ATOMIC_OP
+++#undef DECL_ATOMIC_OP_TYPE
+++#undef DECL_ATOMIC_OP_SPACE
+++
+++/////////////////////////////////////////////////////////////////////////////
++ // Force the compilation to SIMD8 or SIMD16
++ /////////////////////////////////////////////////////////////////////////////
++
++--
++1.7.10.4
++
diff --cc debian/patches/0007-Add-atomic-test-case.patch
index 0000000,0000000..6a3ae8a
new file mode 100644
--- /dev/null
+++ b/debian/patches/0007-Add-atomic-test-case.patch
@@@ -1,0 -1,0 +1,188 @@@
++From d83c5424c918bba806f39c72f27db6a788600dc9 Mon Sep 17 00:00:00 2001
++From: Yang Rong <rong.r.yang at intel.com>
++Date: Wed, 26 Jun 2013 15:29:23 +0800
++Subject: [PATCH 07/11] Add atomic test case.
++To: beignet at lists.freedesktop.org
++
++Signed-off-by: Yang Rong <rong.r.yang at intel.com>
++---
++ kernels/compiler_atomic_functions.cl | 55 ++++++++++++++++-----
++ utests/CMakeLists.txt | 1 +
++ utests/compiler_atomic_functions.cpp | 87 ++++++++++++++++++++++++++++++++--
++ 3 files changed, 127 insertions(+), 16 deletions(-)
++
++diff --git a/kernels/compiler_atomic_functions.cl b/kernels/compiler_atomic_functions.cl
++index 23f3e73..dd9ec56 100644
++--- a/kernels/compiler_atomic_functions.cl
+++++ b/kernels/compiler_atomic_functions.cl
++@@ -1,14 +1,43 @@
++-/* test OpenCL 1.1 Atomic Functions (section 6.11.1, 9.4) */
++-__kernel void compiler_atomic_functions(global int *a, global int *b) {
++- atomic_add(a, *b);
++- atomic_sub(a, *b);
++- atomic_xchg(a, *b);
++- atomic_inc(a);
++- atomic_dec(a);
++- atomic_cmpxchg(a, b, 100);
++- atomic_min(a, *b);
++- atomic_max(a, *b);
++- atomic_and(a, *b);
++- atomic_or(a, *b);
++- atomic_xor(a, *b);
+++__kernel void compiler_atomic_functions(__global int *dst, __local int *tmp, __global int *src) {
+++ int lid = get_local_id(0);
+++ int i = lid % 12;
+++ atomic_xchg(&tmp[4], -1);
+++ switch(i) {
+++ case 0: atomic_inc(&tmp[i]); break;
+++ case 1: atomic_dec(&tmp[i]); break;
+++ case 2: atomic_add(&tmp[i], src[lid]); break;
+++ case 3: atomic_sub(&tmp[i], src[lid]); break;
+++ case 4: atomic_and(&tmp[i], ~(src[lid]<<(lid>>2))); break;
+++ case 5: atomic_or (&tmp[i], src[lid]<<(lid>>2)); break;
+++ case 6: atomic_xor(&tmp[i], src[lid]); break;
+++ case 7: atomic_min(&tmp[i], -src[lid]); break;
+++ case 8: atomic_max(&tmp[i], src[lid]); break;
+++ case 9: atomic_min((__local unsigned int *)&tmp[i], -src[lid]); break;
+++ case 10: atomic_max((__local unsigned int *)&tmp[i], src[lid]); break;
+++ case 11: atomic_cmpxchg(&(tmp[i]), 0, src[10]); break;
+++ default: break;
+++ }
+++
+++ switch(i) {
+++ case 0: atomic_inc(&dst[i]); break;
+++ case 1: atomic_dec(&dst[i]); break;
+++ case 2: atomic_add(&dst[i], src[lid]); break;
+++ case 3: atomic_sub(&dst[i], src[lid]); break;
+++ case 4: atomic_and(&dst[i], ~(src[lid]<<(lid>>2))); break;
+++ case 5: atomic_or (&dst[i], src[lid]<<(lid>>2)); break;
+++ case 6: atomic_xor(&dst[i], src[lid]); break;
+++ case 7: atomic_min(&dst[i], -src[lid]); break;
+++ case 8: atomic_max(&dst[i], src[lid]); break;
+++ case 9: atomic_min((__global unsigned int *)&dst[i], -src[lid]); break;
+++ case 10: atomic_max((__global unsigned int *)&dst[i], src[lid]); break;
+++ case 11: atomic_cmpxchg(&dst[i], 0, src[10]); break;
+++ default: break;
+++ }
+++
+++ barrier(CLK_GLOBAL_MEM_FENCE);
+++
+++ if(get_global_id(0) == 0) {
+++ for(i=0; i<12; i=i+1)
+++ atomic_add(&dst[i], tmp[i]);
+++ }
++ }
++diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
++index b75f3b4..c313acd 100644
++--- a/utests/CMakeLists.txt
+++++ b/utests/CMakeLists.txt
++@@ -75,6 +75,7 @@ set (utests_sources
++ compiler_write_only_shorts.cpp
++ compiler_switch.cpp
++ compiler_math.cpp
+++ compiler_atomic_functions.cpp
++ compiler_insn_selection_min.cpp
++ compiler_insn_selection_max.cpp
++ compiler_insn_selection_masked_min_max.cpp
++diff --git a/utests/compiler_atomic_functions.cpp b/utests/compiler_atomic_functions.cpp
++index 20202da..71e8384 100644
++--- a/utests/compiler_atomic_functions.cpp
+++++ b/utests/compiler_atomic_functions.cpp
++@@ -1,10 +1,91 @@
++ #include "utest_helper.hpp"
+++#include <cmath>
+++#include <algorithm>
++
++-void compiler_atomic_functions(void)
+++#define GROUP_NUM 16
+++#define LOCAL_SIZE 64
+++static void cpu_compiler_atomic(int *dst, int *src)
++ {
++- OCL_CREATE_KERNEL("compiler_atomic_functions");
+++ dst[4] = 0xffffffff;
+++ int tmp[16] = { 0 };
+++
+++ for(int j=0; j<LOCAL_SIZE; j++) {
+++ int i = j % 12;
+++
+++ switch(i) {
+++ case 0: tmp[i] += 1; break;
+++ case 1: tmp[i] -= 1; break;
+++ case 2: tmp[i] += src[j]; break;
+++ case 3: tmp[i] -= src[j]; break;
+++ case 4: tmp[i] &= ~(src[j]<<(j>>2)); break;
+++ case 5: tmp[i] |= src[j]<<(j>>2); break;
+++ case 6: tmp[i] ^= src[j]; break;
+++ case 7: tmp[i] = tmp[i] < -src[j] ? tmp[i] : -src[j]; break;
+++ case 8: tmp[i] = tmp[i] > src[j] ? tmp[i] : src[j]; break;
+++ case 9: tmp[i] = (unsigned int)tmp[i] < (unsigned int)(-src[j]) ? tmp[i] : -src[j]; break;
+++ case 10: tmp[i] = (unsigned int)tmp[i] > (unsigned int)(src[j]) ? tmp[i] : src[j]; break;
+++ case 11: tmp[i] = src[10]; break;
+++ default: break;
+++ }
+++ }
+++
+++ for(int k=0; k<GROUP_NUM; k++) {
+++ for(int j=0; j<LOCAL_SIZE; j++) {
+++ int i = j % 12;
+++
+++ switch(i) {
+++ case 0: dst[i] += 1; break;
+++ case 1: dst[i] -= 1; break;
+++ case 2: dst[i] += src[j]; break;
+++ case 3: dst[i] -= src[j]; break;
+++ case 4: dst[i] &= ~(src[j]<<(j>>2)); break;
+++ case 5: dst[i] |= src[j]<<(j>>2); break;
+++ case 6: dst[i] ^= src[j]; break;
+++ case 7: dst[i] = dst[i] < -src[j] ? dst[i] : -src[j]; break;
+++ case 8: dst[i] = dst[i] > src[j] ? dst[i] : src[j]; break;
+++ case 9: dst[i] = (unsigned int)dst[i] < (unsigned int)(-src[j]) ? dst[i] : -src[j]; break;
+++ case 10: dst[i] = (unsigned int)dst[i] > (unsigned int)(src[j]) ? dst[i] : src[j]; break;
+++ case 11: dst[i] = src[10]; break;
+++ default: break;
+++ }
+++ }
+++ }
+++
+++ for(int i=0; i<12; i++)
+++ dst[i] += tmp[i];
++ }
++
++-MAKE_UTEST_FROM_FUNCTION(compiler_atomic_functions);
+++static void compiler_atomic_functions(void)
+++{
+++ const size_t n = GROUP_NUM * LOCAL_SIZE;
+++ int cpu_dst[16] = {0}, cpu_src[256];
++
+++ globals[0] = n;
+++ locals[0] = LOCAL_SIZE;
+++
+++ // Setup kernel and buffers
+++ OCL_CREATE_KERNEL("compiler_atomic_functions");
+++ OCL_CREATE_BUFFER(buf[0], 0, 16 * sizeof(int), NULL);
+++ OCL_CREATE_BUFFER(buf[1], 0, locals[0] * sizeof(int), NULL);
+++ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+++ OCL_SET_ARG(1, 16 * sizeof(int), NULL);
+++ OCL_SET_ARG(2, sizeof(cl_mem), &buf[1]);
+++
+++ OCL_MAP_BUFFER(1);
+++ for (uint32_t i = 0; i < locals[0]; ++i)
+++ cpu_src[i] = ((int*)buf_data[1])[i] = rand() & 0xff;
+++ cpu_compiler_atomic(cpu_dst, cpu_src);
+++ OCL_UNMAP_BUFFER(1);
+++ OCL_NDRANGE(1);
+++
+++ OCL_MAP_BUFFER(0);
+++
+++ // Check results
+++ for(int i=0; i<12; i++) {
+++ //printf("The dst(%d) gpu(0x%x) cpu(0x%x)\n", i, ((uint32_t *)buf_data[0])[i], cpu_dst[i]);
+++ OCL_ASSERT(((int *)buf_data[0])[i] == cpu_dst[i]);
+++ }
+++ OCL_UNMAP_BUFFER(0);
+++}
++
+++MAKE_UTEST_FROM_FUNCTION(compiler_atomic_functions)
++--
++1.7.10.4
++
diff --cc debian/patches/0008-support-built-in-function-rotate.patch
index 0000000,0000000..9b917e8
new file mode 100644
--- /dev/null
+++ b/debian/patches/0008-support-built-in-function-rotate.patch
@@@ -1,0 -1,0 +1,58 @@@
++From 2e008d20fb703db8afb84476ad599dca92d9f763 Mon Sep 17 00:00:00 2001
++From: Homer Hsing <homer.xing at intel.com>
++Date: Wed, 26 Jun 2013 15:51:51 +0800
++Subject: [PATCH 08/11] support built-in function "rotate"
++To: beignet at lists.freedesktop.org
++
++Signed-off-by: Homer Hsing <homer.xing at intel.com>
++---
++ backend/src/ocl_stdlib.h | 33 +++++++++++++++++++++++++++++++++
++ 1 file changed, 33 insertions(+)
++
++diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
++index 7a98e04..133e995 100644
++--- a/backend/src/ocl_stdlib.h
+++++ b/backend/src/ocl_stdlib.h
++@@ -4355,6 +4355,39 @@ DEC(16)
++ #undef DEC4
++ #undef DEC8
++ #undef DEC16
+++
+++INLINE_OVERLOADABLE uchar __rotate_left(uchar x, uchar y) { return (x << y) | (x >> (8 - y)); }
+++INLINE_OVERLOADABLE char __rotate_left(char x, char y) { return __rotate_left((uchar)x, (uchar)y); }
+++INLINE_OVERLOADABLE ushort __rotate_left(ushort x, ushort y) { return (x << y) | (x >> (16 - y)); }
+++INLINE_OVERLOADABLE short __rotate_left(short x, short y) { return __rotate_left((ushort)x, (ushort)y); }
+++INLINE_OVERLOADABLE uint __rotate_left(uint x, uint y) { return (x << y) | (x >> (32 - y)); }
+++INLINE_OVERLOADABLE int __rotate_left(int x, int y) { return __rotate_left((uint)x, (uint)y); }
+++#define DEF(type, n, m) INLINE_OVERLOADABLE type rotate(type x, type y) { return __rotate_left(x, (type)(y < 0 ? n + y : y & m)); }
+++DEF(char, 8, 7)
+++DEF(uchar, 8, 7)
+++DEF(short, 16, 15)
+++DEF(ushort, 16, 15)
+++DEF(int, 32, 31)
+++DEF(uint, 32, 31)
+++#undef DEF
+++#define DEC2(type) INLINE_OVERLOADABLE type##2 rotate(type##2 a, type##2 b) { return (rotate(a.s0, b.s0), rotate(a.s1, b.s1)); }
+++#define DEC3(type) INLINE_OVERLOADABLE type##3 rotate(type##3 a, type##3 b) { return (rotate(a.s0, b.s0), rotate(a.s1, b.s1), rotate(a.s2, b.s2)); }
+++#define DEC4(type) INLINE_OVERLOADABLE type##4 rotate(type##4 a, type##4 b) { return (rotate(a.s0, b.s0), rotate(a.s1, b.s1), rotate(a.s2, b.s2), rotate(a.s3, b.s3)); }
+++#define DEC8(type) INLINE_OVERLOADABLE type##8 rotate(type##8 a, type##8 b) { return (rotate(a.s0, b.s0), rotate(a.s1, b.s1), rotate(a.s2, b.s2), rotate(a.s3, b.s3), rotate(a.s4, b.s4), rotate(a.s5, b.s5), rotate(a.s6, b.s6), rotate(a.s7, b.s7)); }
+++#define DEC16(type) INLINE_OVERLOADABLE type##16 rotate(type##16 a, type##16 b) { return (rotate(a.s0, b.s0), rotate(a.s1, b.s1), rotate(a.s2, b.s2), rotate(a.s3, b.s3), rotate(a.s4, b.s4), rotate(a.s5, b.s5), rotate(a.s6, b.s6), rotate(a.s7, b.s7), rotate(a.s8, b.s8), rotate(a.s9, b.s9), rotate(a.sa, b.sa), rotate(a.sb, b.sb), rotate(a.sc, b.sc), rotate(a.sd, b.sd), rotate(a.se, b.se), rotate(a.sf, b.sf)); }
+++#define DEF(n) DEC##n(char); DEC##n(uchar); DEC##n(short); DEC##n(ushort); DEC##n(int); DEC##n(uint)
+++DEF(2)
+++DEF(3)
+++DEF(4)
+++DEF(8)
+++DEF(16)
+++#undef DEF
+++#undef DEC2
+++#undef DEC3
+++#undef DEC4
+++#undef DEC8
+++#undef DEC16
+++
++ /////////////////////////////////////////////////////////////////////////////
++ // Work Items functions (see 6.11.1 of OCL 1.1 spec)
++ /////////////////////////////////////////////////////////////////////////////
++--
++1.7.10.4
++
diff --cc debian/patches/0009-test-case-for-function-rotate.patch
index 0000000,0000000..eef926a
new file mode 100644
--- /dev/null
+++ b/debian/patches/0009-test-case-for-function-rotate.patch
@@@ -1,0 -1,0 +1,87 @@@
++From ac4b81f5d44a0803f24d3fe2202ebf4c6bd7d523 Mon Sep 17 00:00:00 2001
++From: Homer Hsing <homer.xing at intel.com>
++Date: Wed, 26 Jun 2013 15:51:52 +0800
++Subject: [PATCH 09/11] test case for function "rotate"
++To: beignet at lists.freedesktop.org
++
++Signed-off-by: Homer Hsing <homer.xing at intel.com>
++---
++ kernels/compiler_rotate.cl | 5 +++++
++ utests/CMakeLists.txt | 1 +
++ utests/compiler_rotate.cpp | 40 ++++++++++++++++++++++++++++++++++++++++
++ 3 files changed, 46 insertions(+)
++ create mode 100644 kernels/compiler_rotate.cl
++ create mode 100644 utests/compiler_rotate.cpp
++
++diff --git a/kernels/compiler_rotate.cl b/kernels/compiler_rotate.cl
++new file mode 100644
++index 0000000..8d0dd0f
++--- /dev/null
+++++ b/kernels/compiler_rotate.cl
++@@ -0,0 +1,5 @@
+++kernel void compiler_rotate(global int *src, global int *dst, global int *y) {
+++ int i = get_global_id(0);
+++ dst[i] = rotate(src[i], y[i]);
+++}
+++
++diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
++index c313acd..3326064 100644
++--- a/utests/CMakeLists.txt
+++++ b/utests/CMakeLists.txt
++@@ -55,6 +55,7 @@ set (utests_sources
++ compiler_lower_return1.cpp
++ compiler_lower_return2.cpp
++ compiler_multiple_kernels.cpp
+++ compiler_rotate.cpp
++ compiler_saturate.cpp
++ compiler_saturate_sub.cpp
++ compiler_shift_right.cpp
++diff --git a/utests/compiler_rotate.cpp b/utests/compiler_rotate.cpp
++new file mode 100644
++index 0000000..bf52ca4
++--- /dev/null
+++++ b/utests/compiler_rotate.cpp
++@@ -0,0 +1,40 @@
+++#include "utest_helper.hpp"
+++
+++int cpu(int src, int y) {
+++ return (src << y) | (src >> (32 - y));
+++}
+++
+++void compiler_rotate(void)
+++{
+++ const int n = 32;
+++ int src[n], y[n];
+++
+++ // Setup kernel and buffers
+++ OCL_CREATE_KERNEL("compiler_rotate");
+++ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int), NULL);
+++ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int), NULL);
+++ OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(int), NULL);
+++ 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_MAP_BUFFER(0);
+++ OCL_MAP_BUFFER(2);
+++ for (int i = 0; i < n; ++i) {
+++ src[i] = ((int*)buf_data[0])[i] = rand();
+++ y[i] = ((int*)buf_data[2])[i] = rand() & 31;
+++ }
+++ OCL_UNMAP_BUFFER(0);
+++ OCL_UNMAP_BUFFER(2);
+++
+++ OCL_NDRANGE(1);
+++
+++ OCL_MAP_BUFFER(1);
+++ for (int i = 0; i < n; ++i)
+++ OCL_ASSERT(((int*)buf_data[1])[i] == cpu(src[i], y[i]));
+++ OCL_UNMAP_BUFFER(1);
+++}
+++
+++MAKE_UTEST_FROM_FUNCTION(compiler_rotate);
++--
++1.7.10.4
++
diff --cc debian/patches/0010-GBE-Add-more-support-of-char-and-short-arithmetic.patch
index 22a0eec,0000000..04dc32f
mode 100644,000000..100644
--- a/debian/patches/0010-GBE-Add-more-support-of-char-and-short-arithmetic.patch
+++ b/debian/patches/0010-GBE-Add-more-support-of-char-and-short-arithmetic.patch
@@@ -1,139 -1,0 +1,153 @@@
- From 1525f53083d7623659e51a9d6f1e4835a83a6caf Mon Sep 17 00:00:00 2001
++From eddef9eb62ed9ea04e657687c8ed28f43f1e0584 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
++Date: Wed, 26 Jun 2013 15:52:12 +0800
++Subject: [PATCH 10/11] 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/backend/gen_insn_selection.cpp | 87 ++++++++++++++++++++--------
+ backend/src/llvm/llvm_gen_backend.cpp | 4 +-
- 2 files changed, 57 insertions(+), 14 deletions(-)
++ 2 files changed, 65 insertions(+), 26 deletions(-)
+
+diff --git a/backend/src/backend/gen_insn_selection.cpp b/backend/src/backend/gen_insn_selection.cpp
- index 1e5f514..b1c6093 100644
++index c64afd9..e98be3e 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);
++@@ -1293,6 +1293,54 @@ namespace gbe
++ this->opcodes.push_back(ir::Opcode(op));
++ }
++
+++ bool emitDivRemInst(Selection::Opaque &sel, SelectionDAG &dag, ir::Opcode op) const
+++ {
+++ using namespace ir;
+++ const ir::BinaryInstruction &insn = cast<BinaryInstruction>(dag.insn);
+++ const Type type = insn.getType();
+++ GenRegister dst = sel.selReg(insn.getDst(0), type);
+++ GenRegister src0 = sel.selReg(insn.getSrc(0), type);
+++ GenRegister src1 = sel.selReg(insn.getSrc(1), type);
++ const uint32_t simdWidth = sel.curr.execWidth;
++ const RegisterFamily family = getFamily(type);
+++ uint32_t function = (op == OP_DIV)?
+++ GEN_MATH_FUNCTION_INT_DIV_QUOTIENT :
+++ GEN_MATH_FUNCTION_INT_DIV_REMAINDER;
++
++ //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));
+++ if((family == FAMILY_WORD || family == FAMILY_BYTE)) {
+++ GenRegister tmp0, tmp1;
+++ ir::Register reg = sel.reg(FAMILY_DWORD);
+++
+++ tmp0 = GenRegister::udxgrf(simdWidth, reg);
++ 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);
+++ sel.MATH(tmp0, function, tmp0, tmp1);
++ GenRegister unpacked;
++ if(family == FAMILY_WORD) {
- + unpacked = GenRegister::unpacked_uw(sel.reg(FAMILY_DWORD));
+++ unpacked = GenRegister::unpacked_uw(reg);
++ } else {
- + unpacked = GenRegister::unpacked_ub(sel.reg(FAMILY_DWORD));
+++ unpacked = GenRegister::unpacked_ub(reg);
++ }
++ unpacked = GenRegister::retype(unpacked, getGenType(type));
- + sel.MOV(unpacked, tmp2);
++ sel.MOV(dst, unpacked);
-
- + markAllChildren(dag);
- + return true;
+++ } else if (type == TYPE_S32 || type == TYPE_U32 ) {
+++ sel.MATH(dst, function, src0, src1);
+++ } else if(type == TYPE_FLOAT) {
+++ GBE_ASSERT(op != OP_REM);
+++ sel.MATH(dst, GEN_MATH_FUNCTION_FDIV, src0, src1);
+++ } else {
+++ NOT_IMPLEMENTED;
++ }
- // 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;
+++ markAllChildren(dag);
+++ return true;
+++ }
+++
++ INLINE bool emit(Selection::Opaque &sel, SelectionDAG &dag) const
++ {
++ using namespace ir;
++@@ -1301,29 +1349,20 @@ namespace gbe
++ const Type type = insn.getType();
++ GenRegister dst = sel.selReg(insn.getDst(0), type);
++
++- // 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;
++- 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;
++- 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;
++- markAllChildren(dag);
++- return true;
+++ if(opcode == OP_DIV || opcode == OP_REM) {
+++ return this->emitDivRemInst(sel, dag, opcode);
+ }
- if (opcode == OP_REM) {
++- if (opcode == OP_REM) {
+++ // Immediates not supported
+++ if (opcode == OP_POW) {
+ 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);
++- sel.MATH(dst, GEN_MATH_FUNCTION_INT_DIV_REMAINDER, src0, src1);
+- markAllChildren(dag);
+- } else
- - NOT_IMPLEMENTED;
+++
+++ if(type == TYPE_FLOAT) {
+++ sel.MATH(dst, GEN_MATH_FUNCTION_POW, src0, src1);
++ } else {
- + GBE_ASSERTM(0, "Unsupported type in remainder operation!");
++ NOT_IMPLEMENTED;
++ }
++ markAllChildren(dag);
+ return true;
+ }
+
- @@ -1345,14 +1388,14 @@ namespace gbe
++@@ -1385,14 +1424,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
++index 5db3fbe..fa052ce 100644
+--- a/backend/src/llvm/llvm_gen_backend.cpp
++++ b/backend/src/llvm/llvm_gen_backend.cpp
- @@ -1276,10 +1276,10 @@ namespace gbe
++@@ -1280,10 +1280,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/0011-utests-Add-basic-arithmetic-test-case.patch
index a3ae637,0000000..9a2e59e
mode 100644,000000..100644
--- a/debian/patches/0011-utests-Add-basic-arithmetic-test-case.patch
+++ b/debian/patches/0011-utests-Add-basic-arithmetic-test-case.patch
@@@ -1,376 -1,0 +1,336 @@@
- From c89dbb34332c104df22c8ea8c22bac0bcb0b5221 Mon Sep 17 00:00:00 2001
++From 1400415e754d8362ed6a628f5e77c3da2417adae 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
++Date: Wed, 26 Jun 2013 15:52:13 +0800
++Subject: [PATCH 11/11] 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 --
++ kernels/compiler_basic_arithmetic.cl | 53 ++++++++++++++++
++ 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(-)
++ utests/compiler_basic_arithmetic.cpp | 112 ++++++++++++++++++++++++++++++++++
++ utests/compiler_sub_bytes.cpp | 35 -----------
++ utests/compiler_sub_shorts.cpp | 36 -----------
++ 7 files changed, 166 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
++index 0000000..3e145d8
+--- /dev/null
++++ b/kernels/compiler_basic_arithmetic.cl
- @@ -0,0 +1,73 @@
++@@ -0,0 +1,53 @@
++#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)
+++#define DECL_KERNEL_FOR_ALL_TYPE(op) \
+++DECL_KERNEL_##op(char) \
+++DECL_KERNEL_##op(uchar) \
+++DECL_KERNEL_##op(short) \
+++DECL_KERNEL_##op(ushort) \
+++DECL_KERNEL_##op(int) \
+++DECL_KERNEL_##op(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)
+++DECL_KERNEL_FOR_ALL_TYPE(SUB)
+++DECL_KERNEL_FOR_ALL_TYPE(ADD)
+++DECL_KERNEL_FOR_ALL_TYPE(MUL)
+++DECL_KERNEL_FOR_ALL_TYPE(DIV)
+++DECL_KERNEL_FOR_ALL_TYPE(REM)
+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
++index 3326064..3740841 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
++@@ -60,8 +61,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
++index 0000000..dcdd084
+--- /dev/null
++++ b/utests/compiler_basic_arithmetic.cpp
- @@ -0,0 +1,132 @@
++@@ -0,0 +1,112 @@
++#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);
+++#define DECL_TEST_FOR_ALL_TYPE(op)\
+++DECL_TEST_##op(int8_t, char) \
+++DECL_TEST_##op(uint8_t, uchar) \
+++DECL_TEST_##op(int16_t, short) \
+++DECL_TEST_##op(uint16_t, ushort) \
+++DECL_TEST_##op(int32_t, int) \
+++DECL_TEST_##op(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);
+++DECL_TEST_FOR_ALL_TYPE(SUB)
+++DECL_TEST_FOR_ALL_TYPE(ADD)
+++DECL_TEST_FOR_ALL_TYPE(MUL)
+++DECL_TEST_FOR_ALL_TYPE(DIV)
+++DECL_TEST_FOR_ALL_TYPE(REM)
+++#undef DECL_TEST_FOR_ALL_TYPE
+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/series
index 6902ed0,0000000..90fcaf8
mode 100644,000000..100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@@ -1,12 -1,0 +1,15 @@@
+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
++0003-Add-the-builtin-function-abs-and-the-according-test-.patch
++0004-PATCH-Refine-the-get_local_id-.-builtins.patch
++0005-Add-atomic-help-functions.-Support-global-and-local-.patch
++0006-Add-all-atomic-built-in-functions-support.patch
++0007-Add-atomic-test-case.patch
++0008-support-built-in-function-rotate.patch
++0009-test-case-for-function-rotate.patch
++0010-GBE-Add-more-support-of-char-and-short-arithmetic.patch
++0011-utests-Add-basic-arithmetic-test-case.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