[Pkg-opencl-devel] [beignet] 07/47: Imported Debian patch 0.1+git20130418+0546d2e-1
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 21:45:47 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 80848e16edc3860b9224ec20619a4fd592c04501
Author: Simon Richter <sjr at debian.org>
Date: Thu Apr 18 11:51:37 2013 +0200
Imported Debian patch 0.1+git20130418+0546d2e-1
---
debian/beignet0.0.1.install | 1 +
debian/changelog | 6 +
debian/control | 3 +-
debian/intel.icd | 1 +
debian/patches/asd | 47 --
debian/patches/autogen | 974 --------------------------------------
debian/patches/clang-from-path | 20 -
debian/patches/device-info-query | 54 ---
debian/patches/fix-clean | 23 -
debian/patches/glibc-memalign | 19 -
debian/patches/ignore-missing-egl | 50 --
debian/patches/implement-gefa | 12 +-
debian/patches/llvm-3.2 | 19 -
debian/patches/missing-header | 6 +-
debian/patches/path | 21 -
debian/patches/respect-flags | 18 +-
debian/patches/series | 5 -
debian/patches/soname | 14 +-
18 files changed, 35 insertions(+), 1258 deletions(-)
diff --git a/debian/beignet0.0.1.install b/debian/beignet0.0.1.install
index 24c7e1b..3fbe284 100644
--- a/debian/beignet0.0.1.install
+++ b/debian/beignet0.0.1.install
@@ -1,2 +1,3 @@
+debian/intel.icd /etc/OpenCL/vendors
usr/lib/libcl.so.* /usr/lib
usr/lib/libgbe.so.* /usr/lib
diff --git a/debian/changelog b/debian/changelog
index ba92740..11993c3 100644
--- a/debian/changelog
+++ b/debian/changelog
@@ -1,3 +1,9 @@
+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
diff --git a/debian/control b/debian/control
index 9b6d64b..d75231b 100644
--- a/debian/control
+++ b/debian/control
@@ -1,7 +1,7 @@
Source: beignet
Priority: extra
Maintainer: Simon Richter <sjr at debian.org>
-Build-Depends: debhelper (>= 9), cmake, pkg-config, ocl-icd-opencl-dev, libdrm-dev, libgbm-dev, libxfixes-dev, libxext-dev, llvm-3.2-dev | llvm-dev (>= 1:3.2)
+Build-Depends: debhelper (>= 9), cmake, pkg-config, ocl-icd-dev, ocl-icd-opencl-dev, libdrm-dev, libgbm-dev, libxfixes-dev, libxext-dev, llvm-3.2-dev | llvm-dev (>= 1:3.2)
Build-Conflicts: libegl1-mesa-dev (<< 9)
Standards-Version: 3.9.4
Section: libs
@@ -23,6 +23,7 @@ Package: beignet0.0.1
Section: libs
Architecture: i386 amd64 kfreebsd-i386 kfreebsd-amd64
Depends: ${shlibs:Depends}, ${misc:Depends}, clang (<< 1:3.1)
+Provides: opencl-icd
Description: Intel OpenCL library
OpenCL (Open Computing Language) is a multivendor open standard for
general-purpose parallel programming of heterogeneous systems that include
diff --git a/debian/intel.icd b/debian/intel.icd
new file mode 100644
index 0000000..c2ca3f6
--- /dev/null
+++ b/debian/intel.icd
@@ -0,0 +1 @@
+/usr/lib/libcl.so.0
diff --git a/debian/patches/asd b/debian/patches/asd
deleted file mode 100644
index 6720adf..0000000
--- a/debian/patches/asd
+++ /dev/null
@@ -1,47 +0,0 @@
-Description: <short summary of the patch>
- TODO: Put a short summary on the line above and replace this paragraph
- with a longer explanation of this change. Complete the meta-information
- with other relevant fields (see below for details). To make it easier, the
- information below has been extracted from the changelog. Adjust it or drop
- it.
- .
- beignet (0.0.0+git2013.04.11+e6b503e-1) UNRELEASED; urgency=low
- .
- * New upstream release
-Author: Simon Richter <sjr at debian.org>
-
----
-The information above should follow the Patch Tagging Guidelines, please
-checkout http://dep.debian.net/deps/dep3/ to learn about the format. Here
-are templates for supplementary fields that you might want to add:
-
-Origin: <vendor|upstream|other>, <url of original patch>
-Bug: <url in upstream bugtracker>
-Bug-Debian: http://bugs.debian.org/<bugnumber>
-Bug-Ubuntu: https://launchpad.net/bugs/<bugnumber>
-Forwarded: <no|not-needed|url proving that it has been forwarded>
-Reviewed-By: <name and email of someone who approved the patch>
-Last-Update: <YYYY-MM-DD>
-
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/backend/CMakeLists.txt
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/CMakeLists.txt
-@@ -45,16 +45,16 @@ set (CMAKE_C_CXX_FLAGS "-fvisibility=hid
- if (COMPILER STREQUAL "GCC")
- set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -funroll-loops -Wstrict-aliasing=2 -fstrict-aliasing -msse2 -msse3 -mssse3 -msse4.1 -ffast-math -fPIC -Wall")
- set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} ${LLVM_CFLAGS}")
-- set (CMAKE_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -Wno-invalid-offsetof -fno-rtti -std=c++0x")
-+ set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${CMAKE_C_CXX_FLAGS} -Wno-invalid-offsetof -fno-rtti -std=c++0x")
- set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MEMORY_FLAG}")
- set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_COMPILE_UTESTS_FLAG}")
- set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wl,-E")
-- set (CMAKE_SHARED_LINKER_FLAGS "-Wl,--no-undefined ${LLVM_LFLAGS}")
-+ set (CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -Wl,--no-undefined ${LLVM_LFLAGS}")
- set (CMAKE_CXX_FLAGS_DEBUG "-g -DGBE_DEBUG=1")
- set (CMAKE_CXX_FLAGS_RELWITHDEBINFO "-O2 -g -DGBE_DEBUG=1")
- set (CMAKE_CXX_FLAGS_MINSIZEREL "-Os -DNDEBUG -DGBE_DEBUG=0")
- set (CMAKE_CXX_FLAGS_RELEASE "-O2 -DNDEBUG -DGBE_DEBUG=0")
-- set (CMAKE_C_FLAGS "${CMAKE_C_CXX_FLAGS}")
-+ set (CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${CMAKE_C_CXX_FLAGS}")
- set (CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${GBE_DEBUG_MEMORY_FLAG}")
- set (CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${GBE_COMPILE_UTESTS_FLAG}")
- set (CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Wl,-E")
diff --git a/debian/patches/autogen b/debian/patches/autogen
deleted file mode 100644
index 6698ee5..0000000
--- a/debian/patches/autogen
+++ /dev/null
@@ -1,974 +0,0 @@
-Description: <short summary of the patch>
- TODO: Put a short summary on the line above and replace this paragraph
- with a longer explanation of this change. Complete the meta-information
- with other relevant fields (see below for details). To make it easier, the
- information below has been extracted from the changelog. Adjust it or drop
- it.
- .
- beignet (0.0.0+git2013.04.11+e6b503e-1) UNRELEASED; urgency=low
- .
- * New upstream release
-Author: Simon Richter <sjr at debian.org>
-
----
-The information above should follow the Patch Tagging Guidelines, please
-checkout http://dep.debian.net/deps/dep3/ to learn about the format. Here
-are templates for supplementary fields that you might want to add:
-
-Origin: <vendor|upstream|other>, <url of original patch>
-Bug: <url in upstream bugtracker>
-Bug-Debian: http://bugs.debian.org/<bugnumber>
-Bug-Ubuntu: https://launchpad.net/bugs/<bugnumber>
-Forwarded: <no|not-needed|url proving that it has been forwarded>
-Reviewed-By: <name and email of someone who approved the patch>
-Last-Update: <YYYY-MM-DD>
-
---- /dev/null
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/src/ocl_stdlib_str.cpp
-@@ -0,0 +1,816 @@
-+#include "string"
-+namespace gbe {
-+std::string ocl_stdlib_str =
-+"/* \n"
-+"uint* Copyright © 2012 Intel Corporation\n"
-+" *\n"
-+" * This library is free software; you can redistribute it and/or\n"
-+" * modify it under the terms of the GNU Lesser General Public\n"
-+" * License as published by the Free Software Foundation; either\n"
-+" * version 2 of the License, or (at your option) any later version.\n"
-+" *\n"
-+" * This library is distributed in the hope that it will be useful,\n"
-+" * but WITHOUT ANY WARRANTY; without even the implied warranty of\n"
-+" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU\n"
-+" * Lesser General Public License for more details.\n"
-+" *\n"
-+" * You should have received a copy of the GNU Lesser General Public\n"
-+" * License along with this library. If not, see <http://www.gnu.org/licenses/>.\n"
-+" *\n"
-+" * Author: Benjamin Segovia <benjamin.segovia at intel.com>\n"
-+" */\n"
-+"\n"
-+"#ifndef __GEN_OCL_STDLIB_H__\n"
-+"#define __GEN_OCL_STDLIB_H__\n"
-+"\n"
-+"#define INLINE __attribute__((always_inline)) inline\n"
-+"#define OVERLOADABLE __attribute__((overloadable))\n"
-+"#define PURE __attribute__((pure))\n"
-+"#define CONST __attribute__((const))\n"
-+"#define INLINE_OVERLOADABLE __attribute__((overloadable,always_inline))\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL built-in scalar data types\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"typedef unsigned char uchar;\n"
-+"typedef unsigned short ushort;\n"
-+"typedef unsigned int uint;\n"
-+"typedef unsigned long ulong;\n"
-+"typedef __typeof__(sizeof(int)) size_t;\n"
-+"typedef __typeof__((int *)0-(int *)0) ptrdiff_t;\n"
-+"typedef signed int intptr_t;\n"
-+"typedef unsigned int uintptr_t;\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL address space\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define __private __attribute__((address_space(0)))\n"
-+"#define __global __attribute__((address_space(1)))\n"
-+"#define __constant __attribute__((address_space(2)))\n"
-+"#define __local __attribute__((address_space(3)))\n"
-+"#define __texture __attribute__((address_space(4)))\n"
-+"#define __sampler __attribute__((address_space(5)))\n"
-+"#define global __global\n"
-+"//#define local __local\n"
-+"#define constant __constant\n"
-+"#define private __private\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL built-in vector data types\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define DEF(type) typedef type type##2 __attribute__((ext_vector_type(2)));\\\n"
-+" typedef type type##3 __attribute__((ext_vector_type(3)));\\\n"
-+" typedef type type##4 __attribute__((ext_vector_type(4)));\\\n"
-+" typedef type type##8 __attribute__((ext_vector_type(8)));\\\n"
-+" typedef type type##16 __attribute__((ext_vector_type(16)));\n"
-+"DEF(char);\n"
-+"DEF(uchar);\n"
-+"DEF(short);\n"
-+"DEF(ushort);\n"
-+"DEF(int);\n"
-+"DEF(uint);\n"
-+"DEF(long);\n"
-+"DEF(ulong);\n"
-+"DEF(float);\n"
-+"#undef DEF\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL other built-in data types\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"struct _image2d_t;\n"
-+"typedef __texture struct _image2d_t* image2d_t;\n"
-+"struct _image3d_t;\n"
-+"typedef __texture struct _image3d_t* image3d_t;\n"
-+"typedef __sampler uint* sampler_t;\n"
-+"typedef size_t event_t;\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL conversions & type casting\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"union type_cast_4_b {\n"
-+" float f;\n"
-+" uchar4 u4;\n"
-+"};\n"
-+"uchar4 INLINE_OVERLOADABLE as_uchar4(float f) {\n"
-+" union type_cast_4_b u;\n"
-+" u.f = f;\n"
-+" return u.u4;\n"
-+"}\n"
-+"#define DEF(type, n, type2) type##n INLINE_OVERLOADABLE convert_##type##n(type2##n d) { \\\n"
-+" return (type##n)((type)(d.s0), (type)(d.s1), (type)(d.s2), (type)(d.s3)); \\\n"
-+" }\n"
-+"#define DEF2(type) DEF(type, 4, char); \\\n"
-+" DEF(type, 4, uchar); \\\n"
-+" DEF(type, 4, short); \\\n"
-+" DEF(type, 4, ushort); \\\n"
-+" DEF(type, 4, int); \\\n"
-+" DEF(type, 4, uint); \\\n"
-+" DEF(type, 4, long); \\\n"
-+" DEF(type, 4, ulong); \\\n"
-+" DEF(type, 4, float);\n"
-+"DEF2(char);\n"
-+"DEF2(uchar);\n"
-+"DEF2(short);\n"
-+"DEF2(ushort);\n"
-+"DEF2(int);\n"
-+"DEF2(uint);\n"
-+"DEF2(long);\n"
-+"DEF2(ulong);\n"
-+"DEF2(float);\n"
-+"#undef DEF2\n"
-+"#undef DEF\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL preprocessor directives & macros\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define __OPENCL_VERSION__ 110\n"
-+"#define __CL_VERSION_1_0__ 100\n"
-+"#define __CL_VERSION_1_1__ 110\n"
-+"#define __ENDIAN_LITTLE__ 1\n"
-+"#define __kernel_exec(X, TYPE) __kernel __attribute__((work_group_size_hint(X,1,1))) \\\n"
-+" __attribute__((vec_type_hint(TYPE)))\n"
-+"#define kernel_exec(X, TYPE) __kernel_exec(X, TYPE)\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL floating-point macros and pragmas\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define FLT_DIG 6\n"
-+"#define FLT_MANT_DIG 24\n"
-+"#define FLT_MAX_10_EXP +38\n"
-+"#define FLT_MAX_EXP +128\n"
-+"#define FLT_MIN_10_EXP -37\n"
-+"#define FLT_MIN_EXP -125\n"
-+"#define FLT_RADIX 2\n"
-+"#define FLT_MAX 0x1.fffffep127f\n"
-+"#define FLT_MIN 0x1.0p-126f\n"
-+"#define FLT_EPSILON 0x1.0p-23f\n"
-+"\n"
-+"#define MAXFLOAT 3.40282347e38F\n"
-+"#define HUGE_VALF (__builtin_huge_valf())\n"
-+"#define INFINITY (__builtin_inff())\n"
-+"#define NAN (__builtin_nanf(\"\"))\n"
-+"#define M_E_F 2.718281828459045F\n"
-+"#define M_LOG2E_F 1.4426950408889634F\n"
-+"#define M_LOG10E_F 0.43429448190325176F\n"
-+"#define M_LN2_F 0.6931471805599453F\n"
-+"#define M_LN10_F 2.302585092994046F\n"
-+"#define M_PI_F 3.141592653589793F\n"
-+"#define M_PI_2_F 1.5707963267948966F\n"
-+"#define M_PI_4_F 0.7853981633974483F\n"
-+"#define M_1_PI_F 0.3183098861837907F\n"
-+"#define M_2_PI_F 0.6366197723675814F\n"
-+"#define M_2_SQRTPI_F 1.1283791670955126F\n"
-+"#define M_SQRT2_F 1.4142135623730951F\n"
-+"#define M_SQRT1_2_F 0.7071067811865476F\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL integer built-in macros\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define CHAR_BIT 8\n"
-+"#define CHAR_MAX SCHAR_MAX\n"
-+"#define CHAR_MIN SCHAR_MIN\n"
-+"#define INT_MAX 2147483647\n"
-+"#define INT_MIN (-2147483647 - 1)\n"
-+"#define LONG_MAX 0x7fffffffffffffffL\n"
-+"#define LONG_MIN (-0x7fffffffffffffffL - 1)\n"
-+"#define SCHAR_MAX 127\n"
-+"#define SCHAR_MIN (-127 - 1)\n"
-+"#define SHRT_MAX 32767\n"
-+"#define SHRT_MIN (-32767 - 1)\n"
-+"#define UCHAR_MAX 255\n"
-+"#define USHRT_MAX 65535\n"
-+"#define UINT_MAX 0xffffffff\n"
-+"#define ULONG_MAX 0xffffffffffffffffUL\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// OpenCL relational built-in functions\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define DEF DECL(int, float); \\\n"
-+" DECL(int2, float2); \\\n"
-+" DECL(int3, float3); \\\n"
-+" DECL(int4, float4); \\\n"
-+" DECL(int8, float8); \\\n"
-+" DECL(int16, float16);\n"
-+"#define DECL(ret, type) ret INLINE_OVERLOADABLE isequal(type x, type y) { return x == y; }\n"
-+"DEF;\n"
-+"#undef DECL\n"
-+"#define DECL(ret, type) ret INLINE_OVERLOADABLE isnotequal(type x, type y) { return x != y; }\n"
-+"DEF;\n"
-+"#undef DECL\n"
-+"#define DECL(ret, type) ret INLINE_OVERLOADABLE isgreater(type x, type y) { return x > y; }\n"
-+"DEF;\n"
-+"#undef DECL\n"
-+"#define DECL(ret, type) ret INLINE_OVERLOADABLE isgreaterequal(type x, type y) { return x >= y; }\n"
-+"DEF;\n"
-+"#undef DECL\n"
-+"#define DECL(ret, type) ret INLINE_OVERLOADABLE isless(type x, type y) { return x < y; }\n"
-+"DEF;\n"
-+"#undef DECL\n"
-+"#define DECL(ret, type) ret INLINE_OVERLOADABLE islessequal(type x, type y) { return x <= y; }\n"
-+"DEF;\n"
-+"#undef DECL\n"
-+"#undef DEF\n"
-+"\n"
-+"#define SDEF(TYPE) \\\n"
-+"INLINE_OVERLOADABLE TYPE ocl_sadd_sat(TYPE x, TYPE y); \\\n"
-+"INLINE_OVERLOADABLE TYPE ocl_ssub_sat(TYPE x, TYPE y); \\\n"
-+"INLINE_OVERLOADABLE TYPE add_sat(TYPE x, TYPE y) { return ocl_sadd_sat(x, y); } \\\n"
-+"INLINE_OVERLOADABLE TYPE sub_sat(TYPE x, TYPE y) { return ocl_ssub_sat(x, y); }\n"
-+"SDEF(char);\n"
-+"SDEF(short);\n"
-+"SDEF(int);\n"
-+"SDEF(long);\n"
-+"#undef SDEF\n"
-+"#define UDEF(TYPE) \\\n"
-+"INLINE_OVERLOADABLE TYPE ocl_uadd_sat(TYPE x, TYPE y); \\\n"
-+"INLINE_OVERLOADABLE TYPE ocl_usub_sat(TYPE x, TYPE y); \\\n"
-+"INLINE_OVERLOADABLE TYPE add_sat(TYPE x, TYPE y) { return ocl_uadd_sat(x, y); } \\\n"
-+"INLINE_OVERLOADABLE TYPE sub_sat(TYPE x, TYPE y) { return ocl_usub_sat(x, y); }\n"
-+"UDEF(uchar);\n"
-+"UDEF(ushort);\n"
-+"UDEF(uint);\n"
-+"UDEF(ulong);\n"
-+"#undef UDEF\n"
-+"\n"
-+"\n"
-+"uchar INLINE_OVERLOADABLE convert_uchar_sat(float x) {\n"
-+" return add_sat((uchar)x, (uchar)0);\n"
-+"}\n"
-+"\n"
-+"#define DEC2(name) INLINE_OVERLOADABLE int2 name(float2 x) { return (name(x.s0), name(x.s1)); }\n"
-+"#define DEC3(name) INLINE_OVERLOADABLE int3 name(float3 x) { return (name(x.s0), name(x.s1), name(x.s2)); }\n"
-+"#define DEC4(name) INLINE_OVERLOADABLE int4 name(float4 x) { return (name(x.s0), name(x.s1), name(x.s2), name(x.s3)); }\n"
-+"#define DEC8(name) INLINE_OVERLOADABLE int8 name(float8 x) { return (name(x.s0), name(x.s1), name(x.s2), name(x.s3), name(x.s4), name(x.s5), name(x.s6), name(x.s7)); }\n"
-+"#define DEC16(name) INLINE_OVERLOADABLE int16 name(float16 x) { return (name(x.s0), name(x.s1), name(x.s2), name(x.s3), name(x.s4), name(x.s5), name(x.s6), name(x.s7), name(x.s8), name(x.s9), name(x.sA), name(x.sB), name(x.sC), name(x.sD), name(x.sE), name(x.sF)); }\n"
-+"INLINE_OVERLOADABLE int isfinite(float x) { return __builtin_isfinite(x); }\n"
-+"DEC2(isfinite);\n"
-+"DEC3(isfinite);\n"
-+"DEC4(isfinite);\n"
-+"DEC8(isfinite);\n"
-+"DEC16(isfinite);\n"
-+"INLINE_OVERLOADABLE int isinf(float x) { return __builtin_isinf(x); }\n"
-+"DEC2(isinf);\n"
-+"DEC3(isinf);\n"
-+"DEC4(isinf);\n"
-+"DEC8(isinf);\n"
-+"DEC16(isinf);\n"
-+"INLINE_OVERLOADABLE int isnan(float x) { return __builtin_isnan(x); }\n"
-+"DEC2(isnan);\n"
-+"DEC3(isnan);\n"
-+"DEC4(isnan);\n"
-+"DEC8(isnan);\n"
-+"DEC16(isnan);\n"
-+"INLINE_OVERLOADABLE int isnormal(float x) { return __builtin_isnormal(x); }\n"
-+"DEC2(isnormal);\n"
-+"DEC3(isnormal);\n"
-+"DEC4(isnormal);\n"
-+"DEC8(isnormal);\n"
-+"DEC16(isnormal);\n"
-+"INLINE_OVERLOADABLE int signbit(float x) { return __builtin_signbit(x); }\n"
-+"DEC2(signbit);\n"
-+"DEC3(signbit);\n"
-+"DEC4(signbit);\n"
-+"DEC8(signbit);\n"
-+"DEC16(signbit);\n"
-+"#undef DEC2\n"
-+"#undef DEC3\n"
-+"#undef DEC4\n"
-+"#undef DEC8\n"
-+"#undef DEC16\n"
-+"\n"
-+"#define DEC2(name) INLINE_OVERLOADABLE int2 name(float2 x, float2 y) { return (name(x.s0, y.s0), name(x.s1, y.s1)); }\n"
-+"#define DEC3(name) INLINE_OVERLOADABLE int3 name(float3 x, float3 y) { return (name(x.s0, y.s0), name(x.s1, y.s1), name(x.s2, y.s2)); }\n"
-+"#define DEC4(name) INLINE_OVERLOADABLE int4 name(float4 x, float4 y) { return (name(x.s0, y.s0), name(x.s1, y.s1), name(x.s2, y.s2), name(x.s3, y.s3)); }\n"
-+"#define DEC8(name) INLINE_OVERLOADABLE int8 name(float8 x, float8 y) { return (name(x.s0, y.s0), name(x.s1, y.s1), name(x.s2, y.s2), name(x.s3, y.s3), name(x.s4, y.s4), name(x.s5, y.s5), name(x.s6, y.s6), name(x.s7, y.s7)); }\n"
-+"#define DEC16(name) INLINE_OVERLOADABLE int16 name(float16 x, float16 y) { return (name(x.s0, y.s0), name(x.s1, y.s1), name(x.s2, y.s2), name(x.s3, y.s3), name(x.s4, y.s4), name(x.s5, y.s5), name(x.s6, y.s6), name(x.s7, y.s7), name(x.s8, y.s8), name(x.s9, y.s9), name(x.sA, y.sA), name(x.sB, y.sB), name(x.sC, y.sC), name(x.sD, y.sD), name(x.sE, y.sE), name(x.sF, y.sF)); }\n"
-+"INLINE_OVERLOADABLE int islessgreater(float x, float y) { return (x<y)||(x>y); }\n"
-+"DEC2(islessgreater);\n"
-+"DEC3(islessgreater);\n"
-+"DEC4(islessgreater);\n"
-+"DEC8(islessgreater);\n"
-+"DEC16(islessgreater);\n"
-+"INLINE_OVERLOADABLE int isordered(float x, float y) { return isequal(x,x) && isequal(y,y); }\n"
-+"DEC2(isordered);\n"
-+"DEC3(isordered);\n"
-+"DEC4(isordered);\n"
-+"DEC8(isordered);\n"
-+"DEC16(isordered);\n"
-+"INLINE_OVERLOADABLE int isunordered(float x, float y) { return isnan(x) || isnan(y); }\n"
-+"DEC2(isunordered);\n"
-+"DEC3(isunordered);\n"
-+"DEC4(isunordered);\n"
-+"DEC8(isunordered);\n"
-+"DEC16(isunordered);\n"
-+"#undef DEC2\n"
-+"#undef DEC3\n"
-+"#undef DEC4\n"
-+"#undef DEC8\n"
-+"#undef DEC16\n"
-+"#define DEC1(type) INLINE_OVERLOADABLE int any(type a) { return a<0; }\n"
-+"#define DEC2(type) INLINE_OVERLOADABLE int any(type a) { return a.s0<0 || a.s1<0; }\n"
-+"#define DEC3(type) INLINE_OVERLOADABLE int any(type a) { return a.s0<0 || a.s1<0 || a.s2<0; }\n"
-+"#define DEC4(type) INLINE_OVERLOADABLE int any(type a) { return a.s0<0 || a.s1<0 || a.s2<0 || a.s3<0; }\n"
-+"#define DEC8(type) INLINE_OVERLOADABLE int any(type a) { return a.s0<0 || a.s1<0 || a.s2<0 || a.s3<0 || a.s4<0 || a.s5<0 || a.s6<0 || a.s7<0; }\n"
-+"#define DEC16(type) INLINE_OVERLOADABLE int any(type a) { return a.s0<0 || a.s1<0 || a.s2<0 || a.s3<0 || a.s4<0 || a.s5<0 || a.s6<0 || a.s7<0 || a.s8<0 || a.s9<0 || a.sA<0 || a.sB<0 || a.sC<0 || a.sD<0 || a.sE<0 || a.sF<0; }\n"
-+"DEC1(char);\n"
-+"DEC1(short);\n"
-+"DEC1(int);\n"
-+"DEC1(long);\n"
-+"#define DEC(n) DEC##n(char##n); DEC##n(short##n); DEC##n(int##n); DEC##n(long##n);\n"
-+"DEC(2);\n"
-+"DEC(3);\n"
-+"DEC(4);\n"
-+"DEC(8);\n"
-+"DEC(16);\n"
-+"#undef DEC\n"
-+"#undef DEC1\n"
-+"#undef DEC2\n"
-+"#undef DEC3\n"
-+"#undef DEC4\n"
-+"#undef DEC8\n"
-+"#undef DEC16\n"
-+"#define DEC1(type) INLINE_OVERLOADABLE int all(type a) { return a<0; }\n"
-+"#define DEC2(type) INLINE_OVERLOADABLE int all(type a) { return a.s0<0 && a.s1<0; }\n"
-+"#define DEC3(type) INLINE_OVERLOADABLE int all(type a) { return a.s0<0 && a.s1<0 && a.s2<0; }\n"
-+"#define DEC4(type) INLINE_OVERLOADABLE int all(type a) { return a.s0<0 && a.s1<0 && a.s2<0 && a.s3<0; }\n"
-+"#define DEC8(type) INLINE_OVERLOADABLE int all(type a) { return a.s0<0 && a.s1<0 && a.s2<0 && a.s3<0 && a.s4<0 && a.s5<0 && a.s6<0 && a.s7<0; }\n"
-+"#define DEC16(type) INLINE_OVERLOADABLE int all(type a) { return a.s0<0 && a.s1<0 && a.s2<0 && a.s3<0 && a.s4<0 && a.s5<0 && a.s6<0 && a.s7<0 && a.s8<0 && a.s9<0 && a.sA<0 && a.sB<0 && a.sC<0 && a.sD<0 && a.sE<0 && a.sF<0; }\n"
-+"DEC1(char);\n"
-+"DEC1(short);\n"
-+"DEC1(int);\n"
-+"DEC1(long);\n"
-+"#define DEC(n) DEC##n(char##n); DEC##n(short##n); DEC##n(int##n); DEC##n(long##n);\n"
-+"DEC(2);\n"
-+"DEC(3);\n"
-+"DEC(4);\n"
-+"DEC(8);\n"
-+"DEC(16);\n"
-+"#undef DEC\n"
-+"#undef DEC1\n"
-+"#undef DEC2\n"
-+"#undef DEC3\n"
-+"#undef DEC4\n"
-+"#undef DEC8\n"
-+"#undef DEC16\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Work Items functions (see 6.11.1 of OCL 1.1 spec)\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// TODO get_global_offset\n"
-+"// TODO get_work_dim\n"
-+"\n"
-+"#define DECL_INTERNAL_WORK_ITEM_FN(NAME) \\\n"
-+"PURE CONST unsigned int __gen_ocl_##NAME##0(void); \\\n"
-+"PURE CONST unsigned int __gen_ocl_##NAME##1(void); \\\n"
-+"PURE CONST unsigned int __gen_ocl_##NAME##2(void);\n"
-+"DECL_INTERNAL_WORK_ITEM_FN(get_group_id)\n"
-+"DECL_INTERNAL_WORK_ITEM_FN(get_local_id)\n"
-+"DECL_INTERNAL_WORK_ITEM_FN(get_local_size)\n"
-+"DECL_INTERNAL_WORK_ITEM_FN(get_global_size)\n"
-+"DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)\n"
-+"#undef DECL_INTERNAL_WORK_ITEM_FN\n"
-+"\n"
-+"#define DECL_PUBLIC_WORK_ITEM_FN(NAME) \\\n"
-+"inline unsigned NAME(unsigned int dim) { \\\n"
-+" if (dim == 0) return __gen_ocl_##NAME##0(); \\\n"
-+" else if (dim == 1) return __gen_ocl_##NAME##1(); \\\n"
-+" else if (dim == 2) return __gen_ocl_##NAME##2(); \\\n"
-+" else return 0; \\\n"
-+"}\n"
-+"DECL_PUBLIC_WORK_ITEM_FN(get_group_id)\n"
-+"DECL_PUBLIC_WORK_ITEM_FN(get_local_id)\n"
-+"DECL_PUBLIC_WORK_ITEM_FN(get_local_size)\n"
-+"DECL_PUBLIC_WORK_ITEM_FN(get_global_size)\n"
-+"DECL_PUBLIC_WORK_ITEM_FN(get_num_groups)\n"
-+"#undef DECL_PUBLIC_WORK_ITEM_FN\n"
-+"\n"
-+"INLINE uint get_global_id(uint dim) {\n"
-+" return get_local_id(dim) + get_local_size(dim) * get_group_id(dim);\n"
-+"}\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Math Functions (see 6.11.2 of OCL 1.1 spec)\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"PURE CONST float __gen_ocl_fabs(float x);\n"
-+"PURE CONST float __gen_ocl_sin(float x);\n"
-+"PURE CONST float __gen_ocl_cos(float x);\n"
-+"PURE CONST float __gen_ocl_sqrt(float x);\n"
-+"PURE CONST float __gen_ocl_rsqrt(float x);\n"
-+"PURE CONST float __gen_ocl_log(float x);\n"
-+"PURE CONST float __gen_ocl_pow(float x, float y);\n"
-+"PURE CONST float __gen_ocl_rcp(float x);\n"
-+"PURE CONST float __gen_ocl_rndz(float x);\n"
-+"PURE CONST float __gen_ocl_rnde(float x);\n"
-+"PURE CONST float __gen_ocl_rndu(float x);\n"
-+"PURE CONST float __gen_ocl_rndd(float x);\n"
-+"INLINE OVERLOADABLE float native_cos(float x) { return __gen_ocl_cos(x); }\n"
-+"INLINE OVERLOADABLE float native_sin(float x) { return __gen_ocl_sin(x); }\n"
-+"INLINE OVERLOADABLE float native_sqrt(float x) { return __gen_ocl_sqrt(x); }\n"
-+"INLINE OVERLOADABLE float native_rsqrt(float x) { return __gen_ocl_rsqrt(x); }\n"
-+"INLINE OVERLOADABLE float native_log2(float x) { return __gen_ocl_log(x); }\n"
-+"INLINE OVERLOADABLE float native_log(float x) {\n"
-+" return native_log2(x) * 0.6931472002f;\n"
-+"}\n"
-+"INLINE OVERLOADABLE float native_log10(float x) {\n"
-+" return native_log2(x) * 0.3010299956f;\n"
-+"}\n"
-+"INLINE OVERLOADABLE float native_powr(float x, float y) { return __gen_ocl_pow(x,y); }\n"
-+"INLINE OVERLOADABLE float native_recip(float x) { return __gen_ocl_rcp(x); }\n"
-+"INLINE OVERLOADABLE float native_tan(float x) {\n"
-+" return native_sin(x) / native_cos(x);\n"
-+"}\n"
-+"#define E 2.71828182845904523536f\n"
-+"INLINE OVERLOADABLE float native_exp(float x) { return native_powr(E, x); }\n"
-+"#undef E\n"
-+"\n"
-+"// XXX work-around PTX profile\n"
-+"#define sqrt native_sqrt\n"
-+"INLINE OVERLOADABLE float rsqrt(float x) { return native_rsqrt(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_fabs(float x) { return __gen_ocl_fabs(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_trunc(float x) { return __gen_ocl_rndz(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_round(float x) { return __gen_ocl_rnde(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_floor(float x) { return __gen_ocl_rndd(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_ceil(float x) { return __gen_ocl_rndu(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_log(float x) { return native_log(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_log2(float x) { return native_log2(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_log10(float x) { return native_log10(x); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_exp(float x) { return native_exp(x); }\n"
-+"INLINE OVERLOADABLE float powr(float x, float y) { return __gen_ocl_pow(x,y); }\n"
-+"INLINE OVERLOADABLE float fmod(float x, float y) { return x-y*__gen_ocl_rndz(x/y); }\n"
-+"\n"
-+"// TODO use llvm intrinsics definitions\n"
-+"#define cos native_cos\n"
-+"#define sin native_sin\n"
-+"#define pow powr\n"
-+"\n"
-+"INLINE OVERLOADABLE float mad(float a, float b, float c) {\n"
-+" return a*b+c;\n"
-+"}\n"
-+"\n"
-+"INLINE OVERLOADABLE uint select(uint src0, uint src1, uint cond) {\n"
-+" return cond ? src1 : src0;\n"
-+"}\n"
-+"INLINE OVERLOADABLE int select(int src0, int src1, int cond) {\n"
-+" return cond ? src1 : src0;\n"
-+"}\n"
-+"INLINE OVERLOADABLE float select(float src0, float src1, int cond) {\n"
-+" return cond ? src1 : src0;\n"
-+"}\n"
-+"\n"
-+"// This will be optimized out by LLVM and will output LLVM select instructions\n"
-+"#define DECL_SELECT4(TYPE4, TYPE, COND_TYPE4, MASK) \\\n"
-+"INLINE OVERLOADABLE TYPE4 select(TYPE4 src0, TYPE4 src1, COND_TYPE4 cond) { \\\n"
-+" TYPE4 dst; \\\n"
-+" const TYPE x0 = src0.x; /* Fix performance issue with CLANG */ \\\n"
-+" const TYPE x1 = src1.x; \\\n"
-+" const TYPE y0 = src0.y; \\\n"
-+" const TYPE y1 = src1.y; \\\n"
-+" const TYPE z0 = src0.z; \\\n"
-+" const TYPE z1 = src1.z; \\\n"
-+" const TYPE w0 = src0.w; \\\n"
-+" const TYPE w1 = src1.w; \\\n"
-+" dst.x = (cond.x & MASK) ? x1 : x0; \\\n"
-+" dst.y = (cond.y & MASK) ? y1 : y0; \\\n"
-+" dst.z = (cond.z & MASK) ? z1 : z0; \\\n"
-+" dst.w = (cond.w & MASK) ? w1 : w0; \\\n"
-+" return dst; \\\n"
-+"}\n"
-+"DECL_SELECT4(int4, int, int4, 0x80000000)\n"
-+"DECL_SELECT4(float4, float, int4, 0x80000000)\n"
-+"#undef DECL_SELECT4\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Common Functions (see 6.11.4 of OCL 1.1 spec)\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define DECL_MIN_MAX(TYPE) \\\n"
-+"INLINE OVERLOADABLE TYPE max(TYPE a, TYPE b) { \\\n"
-+" return a > b ? a : b; \\\n"
-+"} \\\n"
-+"INLINE OVERLOADABLE TYPE min(TYPE a, TYPE b) { \\\n"
-+" return a < b ? a : b; \\\n"
-+"}\n"
-+"DECL_MIN_MAX(float)\n"
-+"DECL_MIN_MAX(int)\n"
-+"DECL_MIN_MAX(short)\n"
-+"DECL_MIN_MAX(char)\n"
-+"DECL_MIN_MAX(uint)\n"
-+"DECL_MIN_MAX(unsigned short)\n"
-+"DECL_MIN_MAX(unsigned char)\n"
-+"#undef DECL_MIN_MAX\n"
-+"\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_fmax(float a, float b) { return max(a,b); }\n"
-+"INLINE OVERLOADABLE float __gen_ocl_internal_fmin(float a, float b) { return min(a,b); }\n"
-+"INLINE OVERLOADABLE float mix(float x, float y, float a) { return x + (y-x)*a;}\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Geometric functions (see 6.11.5 of OCL 1.1 spec)\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"INLINE OVERLOADABLE float dot(float2 p0, float2 p1) {\n"
-+" return mad(p0.x,p1.x,p0.y*p1.y);\n"
-+"}\n"
-+"INLINE OVERLOADABLE float dot(float3 p0, float3 p1) {\n"
-+" return mad(p0.x,p1.x,mad(p0.z,p1.z,p0.y*p1.y));\n"
-+"}\n"
-+"INLINE OVERLOADABLE float dot(float4 p0, float4 p1) {\n"
-+" return mad(p0.x,p1.x,mad(p0.w,p1.w,mad(p0.z,p1.z,p0.y*p1.y)));\n"
-+"}\n"
-+"\n"
-+"INLINE OVERLOADABLE float dot(float8 p0, float8 p1) {\n"
-+" return mad(p0.x,p1.x,mad(p0.s7,p1.s7, mad(p0.s6,p1.s6,mad(p0.s5,p1.s5,\n"
-+" mad(p0.s4,p1.s4,mad(p0.w,p1.w, mad(p0.z,p1.z,p0.y*p1.y)))))));\n"
-+"}\n"
-+"INLINE OVERLOADABLE float dot(float16 p0, float16 p1) {\n"
-+" return mad(p0.sc,p1.sc,mad(p0.sd,p1.sd,mad(p0.se,p1.se,mad(p0.sf,p1.sf,\n"
-+" mad(p0.s8,p1.s8,mad(p0.s9,p1.s9,mad(p0.sa,p1.sa,mad(p0.sb,p1.sb,\n"
-+" mad(p0.x,p1.x,mad(p0.s7,p1.s7, mad(p0.s6,p1.s6,mad(p0.s5,p1.s5,\n"
-+" mad(p0.s4,p1.s4,mad(p0.w,p1.w, mad(p0.z,p1.z,p0.y*p1.y)))))))))))))));\n"
-+"}\n"
-+"\n"
-+"INLINE OVERLOADABLE float length(float x) { return __gen_ocl_fabs(x); }\n"
-+"INLINE OVERLOADABLE float length(float2 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float length(float3 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float length(float4 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float length(float8 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float length(float16 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float distance(float x, float y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float distance(float2 x, float2 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float distance(float3 x, float3 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float distance(float4 x, float4 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float distance(float8 x, float8 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float distance(float16 x, float16 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float normalize(float x) { return 1.f; }\n"
-+"INLINE OVERLOADABLE float2 normalize(float2 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float3 normalize(float3 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float4 normalize(float4 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float8 normalize(float8 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float16 normalize(float16 x) { return x * rsqrt(dot(x, x)); }\n"
-+"\n"
-+"INLINE OVERLOADABLE float fast_length(float x) { return __gen_ocl_fabs(x); }\n"
-+"INLINE OVERLOADABLE float fast_length(float2 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float fast_length(float3 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float fast_length(float4 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float fast_length(float8 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float fast_length(float16 x) { return sqrt(dot(x,x)); }\n"
-+"INLINE OVERLOADABLE float fast_distance(float x, float y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float fast_distance(float2 x, float2 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float fast_distance(float3 x, float3 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float fast_distance(float4 x, float4 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float fast_distance(float8 x, float8 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float fast_distance(float16 x, float16 y) { return length(x-y); }\n"
-+"INLINE OVERLOADABLE float fast_normalize(float x) { return 1.f; }\n"
-+"INLINE OVERLOADABLE float2 fast_normalize(float2 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float3 fast_normalize(float3 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float4 fast_normalize(float4 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float8 fast_normalize(float8 x) { return x * rsqrt(dot(x, x)); }\n"
-+"INLINE OVERLOADABLE float16 fast_normalize(float16 x) { return x * rsqrt(dot(x, x)); }\n"
-+"\n"
-+"INLINE OVERLOADABLE float3 cross(float3 v0, float3 v1) {\n"
-+" return v0.yzx*v1.zxy-v0.zxy*v1.yzx;\n"
-+"}\n"
-+"INLINE OVERLOADABLE float4 cross(float4 v0, float4 v1) {\n"
-+" return (float4)(v0.yzx*v1.zxy-v0.zxy*v1.yzx, 0.f);\n"
-+"}\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Vector loads and stores\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"\n"
-+"// These loads and stores will use untyped reads and writes, so we can just\n"
-+"// cast to vector loads / stores. Not C99 compliant BTW due to aliasing issue.\n"
-+"// Well we do not care, we do not activate TBAA in the compiler\n"
-+"#define DECL_UNTYPED_RW_SPACE_N(TYPE, DIM, SPACE) \\\n"
-+"INLINE OVERLOADABLE TYPE##DIM vload##DIM(size_t offset, const SPACE TYPE *p) { \\\n"
-+" return *(SPACE TYPE##DIM *) (p + DIM * offset); \\\n"
-+"} \\\n"
-+"INLINE OVERLOADABLE void vstore##DIM(TYPE##DIM v, size_t offset, SPACE TYPE *p) { \\\n"
-+" *(SPACE TYPE##DIM *) (p + DIM * offset) = v; \\\n"
-+"}\n"
-+"\n"
-+"#define DECL_UNTYPED_RW_ALL_SPACE(TYPE, SPACE) \\\n"
-+" DECL_UNTYPED_RW_SPACE_N(TYPE, 2, SPACE) \\\n"
-+" DECL_UNTYPED_RW_SPACE_N(TYPE, 3, SPACE) \\\n"
-+" DECL_UNTYPED_RW_SPACE_N(TYPE, 4, SPACE) \\\n"
-+" DECL_UNTYPED_RW_SPACE_N(TYPE, 8, SPACE) \\\n"
-+" DECL_UNTYPED_RW_SPACE_N(TYPE, 16, SPACE)\n"
-+"\n"
-+"#define DECL_UNTYPED_RW_ALL(TYPE) \\\n"
-+" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __global) \\\n"
-+" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __local) \\\n"
-+" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __constant) \\\n"
-+" DECL_UNTYPED_RW_ALL_SPACE(TYPE, __private)\n"
-+"\n"
-+"DECL_UNTYPED_RW_ALL(char)\n"
-+"DECL_UNTYPED_RW_ALL(uchar)\n"
-+"DECL_UNTYPED_RW_ALL(short)\n"
-+"DECL_UNTYPED_RW_ALL(ushort)\n"
-+"DECL_UNTYPED_RW_ALL(int)\n"
-+"DECL_UNTYPED_RW_ALL(uint)\n"
-+"DECL_UNTYPED_RW_ALL(long)\n"
-+"DECL_UNTYPED_RW_ALL(ulong)\n"
-+"DECL_UNTYPED_RW_ALL(float)\n"
-+"\n"
-+"#undef DECL_UNTYPED_RW_ALL\n"
-+"#undef DECL_UNTYPED_RW_ALL_SPACE\n"
-+"#undef DECL_UNTYPED_RW_SPACE_N\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Declare functions for vector types which are derived from scalar ones\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define DECL_VECTOR_1OP(NAME, TYPE) \\\n"
-+" INLINE OVERLOADABLE TYPE##2 NAME(TYPE##2 v) { \\\n"
-+" return (TYPE##2)(NAME(v.x), NAME(v.y)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##3 NAME(TYPE##3 v) { \\\n"
-+" return (TYPE##3)(NAME(v.x), NAME(v.y), NAME(v.z)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##4 NAME(TYPE##4 v) { \\\n"
-+" return (TYPE##4)(NAME(v.x), NAME(v.y), NAME(v.z), NAME(v.w)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##8 NAME(TYPE##8 v) { \\\n"
-+" TYPE##8 dst;\\\n"
-+" dst.s0123 = NAME(v.s0123);\\\n"
-+" dst.s4567 = NAME(v.s4567);\\\n"
-+" return dst;\\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##16 NAME(TYPE##16 v) { \\\n"
-+" TYPE##16 dst;\\\n"
-+" dst.s01234567 = NAME(v.s01234567);\\\n"
-+" dst.s89abcdef = NAME(v.s89abcdef);\\\n"
-+" return dst;\\\n"
-+" }\n"
-+"DECL_VECTOR_1OP(native_cos, float);\n"
-+"DECL_VECTOR_1OP(native_sin, float);\n"
-+"DECL_VECTOR_1OP(native_tan, float);\n"
-+"DECL_VECTOR_1OP(native_sqrt, float);\n"
-+"DECL_VECTOR_1OP(native_rsqrt, float);\n"
-+"DECL_VECTOR_1OP(native_log2, float);\n"
-+"DECL_VECTOR_1OP(native_recip, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_fabs, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_trunc, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_round, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_floor, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_ceil, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_log, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_log2, float);\n"
-+"DECL_VECTOR_1OP(__gen_ocl_internal_log10, float);\n"
-+"#undef DECL_VECTOR_1OP\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Arithmetic functions\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"\n"
-+"#define DECL_VECTOR_2OP(NAME, TYPE) \\\n"
-+" INLINE OVERLOADABLE TYPE##2 NAME(TYPE##2 v0, TYPE##2 v1) { \\\n"
-+" return (TYPE##2)(NAME(v0.x, v1.x), NAME(v1.y, v1.y)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##3 NAME(TYPE##3 v0, TYPE##3 v1) { \\\n"
-+" return (TYPE##3)(NAME(v0.x, v1.x), NAME(v0.y, v1.y), NAME(v0.z, v1.z)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##4 NAME(TYPE##4 v0, TYPE##4 v1) { \\\n"
-+" return (TYPE##4)(NAME(v0.x, v1.x), NAME(v0.y, v1.y), NAME(v0.z, v1.z), NAME(v0.w, v1.w)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##8 NAME(TYPE##8 v0, TYPE##8 v1) { \\\n"
-+" TYPE##8 dst;\\\n"
-+" dst.s0123 = NAME(v0.s0123, v1.s0123);\\\n"
-+" dst.s4567 = NAME(v0.s4567, v1.s4567);\\\n"
-+" return dst;\\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##16 NAME(TYPE##16 v0, TYPE##16 v1) { \\\n"
-+" TYPE##16 dst;\\\n"
-+" dst.s01234567 = NAME(v0.s01234567, v1.s01234567);\\\n"
-+" dst.s89abcdef = NAME(v0.s89abcdef, v1.s89abcdef);\\\n"
-+" return dst;\\\n"
-+" }\n"
-+"DECL_VECTOR_2OP(min, float);\n"
-+"DECL_VECTOR_2OP(max, float);\n"
-+"DECL_VECTOR_2OP(fmod, float);\n"
-+"DECL_VECTOR_2OP(powr, float);\n"
-+"#undef DECL_VECTOR_2OP\n"
-+"\n"
-+"#define DECL_VECTOR_3OP(NAME, TYPE) \\\n"
-+" INLINE OVERLOADABLE TYPE##2 NAME(TYPE##2 v0, TYPE##2 v1, TYPE##2 v2) { \\\n"
-+" return (TYPE##2)(NAME(v0.x, v1.x, v2.x), NAME(v1.y, v1.y, v2.y)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##3 NAME(TYPE##3 v0, TYPE##3 v1, TYPE##3 v2) { \\\n"
-+" return (TYPE##3)(NAME(v0.x, v1.x, v2.x), NAME(v0.y, v1.y, v2.y), NAME(v0.z, v1.z, v2.z)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##4 NAME(TYPE##4 v0, TYPE##4 v1, TYPE##4 v2) { \\\n"
-+" return (TYPE##4)(NAME(v0.x, v1.x, v2.x), NAME(v0.y, v1.y, v2.y), NAME(v0.z, v1.z, v2.z), NAME(v0.w, v1.w, v2.w)); \\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##8 NAME(TYPE##8 v0, TYPE##8 v1, TYPE##8 v2) { \\\n"
-+" TYPE##8 dst;\\\n"
-+" dst.s0123 = NAME(v0.s0123, v1.s0123, v2.s0123);\\\n"
-+" dst.s4567 = NAME(v0.s4567, v1.s4567, v2.s4567);\\\n"
-+" return dst;\\\n"
-+" }\\\n"
-+" INLINE OVERLOADABLE TYPE##16 NAME(TYPE##16 v0, TYPE##16 v1, TYPE##16 v2) { \\\n"
-+" TYPE##16 dst;\\\n"
-+" dst.s01234567 = NAME(v0.s01234567, v1.s01234567, v2.s01234567);\\\n"
-+" dst.s89abcdef = NAME(v0.s89abcdef, v1.s89abcdef, v2.s89abcdef);\\\n"
-+" return dst;\\\n"
-+" }\n"
-+"DECL_VECTOR_3OP(mad, float);\n"
-+"DECL_VECTOR_3OP(mix, float);\n"
-+"#undef DECL_VECTOR_3OP\n"
-+"\n"
-+"// mix requires more variants\n"
-+"INLINE OVERLOADABLE float2 mix(float2 x, float2 y, float a) { return mix(x,y,(float2)(a));}\n"
-+"INLINE OVERLOADABLE float3 mix(float3 x, float3 y, float a) { return mix(x,y,(float3)(a));}\n"
-+"INLINE OVERLOADABLE float4 mix(float4 x, float4 y, float a) { return mix(x,y,(float4)(a));}\n"
-+"INLINE OVERLOADABLE float8 mix(float8 x, float8 y, float a) { return mix(x,y,(float8)(a));}\n"
-+"INLINE OVERLOADABLE float16 mix(float16 x, float16 y, float a) { return mix(x,y,(float16)(a));}\n"
-+"\n"
-+"// XXX workaround ptx profile\n"
-+"#define fabs __gen_ocl_internal_fabs\n"
-+"#define trunc __gen_ocl_internal_trunc\n"
-+"#define round __gen_ocl_internal_round\n"
-+"#define floor __gen_ocl_internal_floor\n"
-+"#define ceil __gen_ocl_internal_ceil,\n"
-+"#define log __gen_ocl_internal_log\n"
-+"#define log2 __gen_ocl_internal_log2\n"
-+"#define log10 __gen_ocl_internal_log10\n"
-+"#define exp __gen_ocl_internal_exp\n"
-+"#define fmin __gen_ocl_internal_fmin\n"
-+"#define fmax __gen_ocl_internal_fmax\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Synchronization functions\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"#define CLK_LOCAL_MEM_FENCE (1 << 0)\n"
-+"#define CLK_GLOBAL_MEM_FENCE (1 << 1)\n"
-+"\n"
-+"void __gen_ocl_barrier_local(void);\n"
-+"void __gen_ocl_barrier_global(void);\n"
-+"void __gen_ocl_barrier_local_and_global(void);\n"
-+"\n"
-+"typedef uint cl_mem_fence_flags;\n"
-+"INLINE void barrier(cl_mem_fence_flags flags) {\n"
-+" if (flags == (CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE))\n"
-+" __gen_ocl_barrier_local_and_global();\n"
-+" else if (flags == CLK_LOCAL_MEM_FENCE)\n"
-+" __gen_ocl_barrier_local();\n"
-+" else if (flags == CLK_GLOBAL_MEM_FENCE)\n"
-+" __gen_ocl_barrier_global();\n"
-+"}\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Force the compilation to SIMD8 or SIMD16\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"\n"
-+"int __gen_ocl_force_simd8(void);\n"
-+"int __gen_ocl_force_simd16(void);\n"
-+"\n"
-+"#define NULL ((void*)0)\n"
-+"\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"// Image access functions\n"
-+"/////////////////////////////////////////////////////////////////////////////\n"
-+"\n"
-+"OVERLOADABLE int4 __gen_ocl_read_imagei(uint surface_id, uint sampler, int u, int v);\n"
-+"OVERLOADABLE int4 __gen_ocl_read_imagei(uint surface_id, uint sampler, float u, float v);\n"
-+"OVERLOADABLE uint4 __gen_ocl_read_imageui(uint surface_id, uint sampler, int u, int v);\n"
-+"OVERLOADABLE uint4 __gen_ocl_read_imageui(uint surface_id, uint sampler, float u, float v);\n"
-+"OVERLOADABLE float4 __gen_ocl_read_imagef(uint surface_id, uint sampler, int u, int v);\n"
-+"OVERLOADABLE float4 __gen_ocl_read_imagef(uint surface_id, uint sampler, float u, float v);\n"
-+"OVERLOADABLE void __gen_ocl_write_imagei(uint surface_id, int u, int v, int4 color);\n"
-+"OVERLOADABLE void __gen_ocl_write_imagei(uint surface_id, float u, float v, int4 color);\n"
-+"OVERLOADABLE void __gen_ocl_write_imageui(uint surface_id, int u, int v, uint4 color);\n"
-+"OVERLOADABLE void __gen_ocl_write_imageui(uint surface_id, float u, float v, uint4 color);\n"
-+"OVERLOADABLE void __gen_ocl_write_imagef(uint surface_id, int u, int v, float4 color);\n"
-+"OVERLOADABLE void __gen_ocl_write_imagef(uint surface_id, float u, float v, float4 color);\n"
-+"\n"
-+"#define GET_IMAGE(cl_image, surface_id) \\\n"
-+" uint surface_id = (uint)cl_image\n"
-+"\n"
-+"#define DECL_READ_IMAGE(type, suffix, coord_type) \\\n"
-+" INLINE_OVERLOADABLE type read_image ##suffix(image2d_t cl_image, sampler_t sampler, coord_type coord) \\\n"
-+" {\\\n"
-+" GET_IMAGE(cl_image, surface_id);\\\n"
-+" return __gen_ocl_read_image ##suffix(surface_id, (uint)sampler, coord.s0, coord.s1);\\\n"
-+" }\n"
-+"\n"
-+"#define DECL_WRITE_IMAGE(type, suffix, coord_type) \\\n"
-+" INLINE_OVERLOADABLE void write_image ##suffix(image2d_t cl_image, coord_type coord, type color)\\\n"
-+" {\\\n"
-+" GET_IMAGE(cl_image, surface_id);\\\n"
-+" __gen_ocl_write_image ##suffix(surface_id, coord.s0, coord.s1, color);\\\n"
-+" }\n"
-+"\n"
-+"#define DECL_IMAGE(type, suffix) \\\n"
-+" DECL_READ_IMAGE(type, suffix, int2) \\\n"
-+" DECL_READ_IMAGE(type, suffix, float2) \\\n"
-+" DECL_WRITE_IMAGE(type, suffix, int2) \\\n"
-+" DECL_WRITE_IMAGE(type, suffix, float2)\n"
-+"\n"
-+"DECL_IMAGE(int4, i)\n"
-+"DECL_IMAGE(uint4, ui)\n"
-+"DECL_IMAGE(float4, f)\n"
-+"\n"
-+"#undef GET_IMAGE\n"
-+"#undef DECL_IMAGE\n"
-+"#undef DECL_READ_IMAGE\n"
-+"#undef DECL_WRITE_IMAGE\n"
-+"#undef INLINE_OVERLOADABLE\n"
-+"\n"
-+"#undef PURE\n"
-+"#undef CONST\n"
-+"#undef OVERLOADABLE\n"
-+"#undef INLINE\n"
-+"#endif /* __GEN_OCL_STDLIB_H__ */\n"
-+"\n"
-+;
-+}
-+
---- /dev/null
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/src/ocl_common_defines_str.cpp
-@@ -0,0 +1,127 @@
-+#include "string"
-+namespace gbe {
-+std::string ocl_common_defines_str =
-+"// This file includes defines that are common to both kernel code and\n"
-+"// the NVPTX back-end.\n"
-+"\n"
-+"//\n"
-+"// Common defines for Image intrinsics\n"
-+"// Channel order\n"
-+"enum {\n"
-+" CLK_R = 0x10B0,\n"
-+" CLK_A = 0x10B1,\n"
-+" CLK_RG = 0x10B2,\n"
-+" CLK_RA = 0x10B3,\n"
-+" CLK_RGB = 0x10B4,\n"
-+" CLK_RGBA = 0x10B5,\n"
-+" CLK_BGRA = 0x10B6,\n"
-+" CLK_ARGB = 0x10B7,\n"
-+"\n"
-+"#if (__NV_CL_C_VERSION == __NV_CL_C_VERSION_1_0)\n"
-+" CLK_xRGB = 0x10B7,\n"
-+"#endif\n"
-+"\n"
-+" CLK_INTENSITY = 0x10B8,\n"
-+" CLK_LUMINANCE = 0x10B9\n"
-+"\n"
-+"#if (__NV_CL_C_VERSION >= __NV_CL_C_VERSION_1_1)\n"
-+" ,\n"
-+" CLK_Rx = 0x10BA,\n"
-+" CLK_RGx = 0x10BB,\n"
-+" CLK_RGBx = 0x10BC\n"
-+"#endif\n"
-+"};\n"
-+"\n"
-+"\n"
-+"typedef enum clk_channel_type {\n"
-+" // valid formats for float return types\n"
-+" CLK_SNORM_INT8 = 0x10D0, // four channel RGBA unorm8\n"
-+" CLK_SNORM_INT16 = 0x10D1, // four channel RGBA unorm16\n"
-+" CLK_UNORM_INT8 = 0x10D2, // four channel RGBA unorm8\n"
-+" CLK_UNORM_INT16 = 0x10D3, // four channel RGBA unorm16\n"
-+" CLK_HALF_FLOAT = 0x10DD, // four channel RGBA half\n"
-+" CLK_FLOAT = 0x10DE, // four channel RGBA float\n"
-+"\n"
-+"#if (__NV_CL_C_VERSION >= __NV_CL_C_VERSION_1_1)\n"
-+" CLK_UNORM_SHORT_565 = 0x10D4,\n"
-+" CLK_UNORM_SHORT_555 = 0x10D5,\n"
-+" CLK_UNORM_INT_101010 = 0x10D6,\n"
-+"#endif\n"
-+"\n"
-+" // valid only for integer return types\n"
-+" CLK_SIGNED_INT8 = 0x10D7,\n"
-+" CLK_SIGNED_INT16 = 0x10D8,\n"
-+" CLK_SIGNED_INT32 = 0x10D9,\n"
-+" CLK_UNSIGNED_INT8 = 0x10DA,\n"
-+" CLK_UNSIGNED_INT16 = 0x10DB,\n"
-+" CLK_UNSIGNED_INT32 = 0x10DC,\n"
-+"\n"
-+" // CI SPI for CPU\n"
-+" __CLK_UNORM_INT8888 , // four channel ARGB unorm8\n"
-+" __CLK_UNORM_INT8888R, // four channel BGRA unorm8\n"
-+"\n"
-+" __CLK_VALID_IMAGE_TYPE_COUNT,\n"
-+" __CLK_INVALID_IMAGE_TYPE = __CLK_VALID_IMAGE_TYPE_COUNT,\n"
-+" __CLK_VALID_IMAGE_TYPE_MASK_BITS = 4, // number of bits required to\n"
-+" // represent any image type\n"
-+" __CLK_VALID_IMAGE_TYPE_MASK = ( 1 << __CLK_VALID_IMAGE_TYPE_MASK_BITS ) - 1\n"
-+"}clk_channel_type;\n"
-+"\n"
-+"typedef enum clk_sampler_type {\n"
-+" __CLK_ADDRESS_BASE = 0,\n"
-+" CLK_ADDRESS_NONE = 0 << __CLK_ADDRESS_BASE,\n"
-+" CLK_ADDRESS_CLAMP = 1 << __CLK_ADDRESS_BASE,\n"
-+" CLK_ADDRESS_CLAMP_TO_EDGE = 2 << __CLK_ADDRESS_BASE,\n"
-+" CLK_ADDRESS_REPEAT = 3 << __CLK_ADDRESS_BASE,\n"
-+" CLK_ADDRESS_MIRROR = 4 << __CLK_ADDRESS_BASE,\n"
-+"\n"
-+"#if (__NV_CL_C_VERSION >= __NV_CL_C_VERSION_1_1)\n"
-+" CLK_ADDRESS_MIRRORED_REPEAT = CLK_ADDRESS_MIRROR,\n"
-+"#endif\n"
-+" __CLK_ADDRESS_MASK = CLK_ADDRESS_NONE | CLK_ADDRESS_CLAMP |\n"
-+" CLK_ADDRESS_CLAMP_TO_EDGE |\n"
-+" CLK_ADDRESS_REPEAT | CLK_ADDRESS_MIRROR,\n"
-+" __CLK_ADDRESS_BITS = 3, // number of bits required to\n"
-+" // represent address info\n"
-+"\n"
-+" __CLK_NORMALIZED_BASE = __CLK_ADDRESS_BITS,\n"
-+" CLK_NORMALIZED_COORDS_FALSE = 0,\n"
-+" CLK_NORMALIZED_COORDS_TRUE = 1 << __CLK_NORMALIZED_BASE,\n"
-+" __CLK_NORMALIZED_MASK = CLK_NORMALIZED_COORDS_FALSE |\n"
-+" CLK_NORMALIZED_COORDS_TRUE,\n"
-+" __CLK_NORMALIZED_BITS = 1, // number of bits required to\n"
-+" // represent normalization\n"
-+"\n"
-+" __CLK_FILTER_BASE = __CLK_NORMALIZED_BASE +\n"
-+" __CLK_NORMALIZED_BITS,\n"
-+" CLK_FILTER_NEAREST = 0 << __CLK_FILTER_BASE,\n"
-+" CLK_FILTER_LINEAR = 1 << __CLK_FILTER_BASE,\n"
-+" CLK_FILTER_ANISOTROPIC = 2 << __CLK_FILTER_BASE,\n"
-+" __CLK_FILTER_MASK = CLK_FILTER_NEAREST | CLK_FILTER_LINEAR |\n"
-+" CLK_FILTER_ANISOTROPIC,\n"
-+" __CLK_FILTER_BITS = 2, // number of bits required to\n"
-+" // represent address info\n"
-+"\n"
-+" __CLK_MIP_BASE = __CLK_FILTER_BASE + __CLK_FILTER_BITS,\n"
-+" CLK_MIP_NEAREST = 0 << __CLK_MIP_BASE,\n"
-+" CLK_MIP_LINEAR = 1 << __CLK_MIP_BASE,\n"
-+" CLK_MIP_ANISOTROPIC = 2 << __CLK_MIP_BASE,\n"
-+" __CLK_MIP_MASK = CLK_MIP_NEAREST | CLK_MIP_LINEAR |\n"
-+" CLK_MIP_ANISOTROPIC,\n"
-+" __CLK_MIP_BITS = 2,\n"
-+"\n"
-+" __CLK_SAMPLER_BITS = __CLK_MIP_BASE + __CLK_MIP_BITS,\n"
-+" __CLK_SAMPLER_MASK = __CLK_MIP_MASK | __CLK_FILTER_MASK |\n"
-+" __CLK_NORMALIZED_MASK | __CLK_ADDRESS_MASK,\n"
-+"\n"
-+" __CLK_ANISOTROPIC_RATIO_BITS = 5,\n"
-+" __CLK_ANISOTROPIC_RATIO_MASK = (int) 0x80000000 >>\n"
-+" (__CLK_ANISOTROPIC_RATIO_BITS-1)\n"
-+"} clk_sampler_type;\n"
-+"\n"
-+"// Memory synchronization\n"
-+"#define CLK_LOCAL_MEM_FENCE (1 << 0)\n"
-+"#define CLK_GLOBAL_MEM_FENCE (1 << 1)\n"
-+;
-+}
-+
diff --git a/debian/patches/clang-from-path b/debian/patches/clang-from-path
deleted file mode 100644
index 24bc834..0000000
--- a/debian/patches/clang-from-path
+++ /dev/null
@@ -1,20 +0,0 @@
-Description: Use clang from PATH
-Author: Simon Richter <sjr at debian.org>
-Last-Update: 2013-04-02
-
-Index: beignet-0.0.0+git2013.04.11+e6b503e/backend/src/backend/program.cpp
-===================================================================
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/backend/src/backend/program.cpp 2013-04-12 08:13:48.000000000 +0200
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/src/backend/program.cpp 2013-04-15 18:33:25.392352309 +0200
-@@ -119,9 +119,9 @@
-
- // Now compile the code to llvm using clang
- #if LLVM_VERSION_MINOR <= 1
-- std::string compileCmd = LLVM_PREFIX "/bin/clang -x cl -fno-color-diagnostics -emit-llvm -O3 -ccc-host-triple ptx32 -c ";
-+ std::string compileCmd = "clang -x cl -fno-color-diagnostics -emit-llvm -O3 -ccc-host-triple ptx32 -c ";
- #else
-- std::string compileCmd = LLVM_PREFIX "/bin/clang -ffp-contract=off -emit-llvm -O3 -target nvptx -x cl -c ";
-+ std::string compileCmd = "clang -ffp-contract=off -emit-llvm -O3 -target nvptx -x cl -c ";
- #endif /* LLVM_VERSION_MINOR <= 1 */
- compileCmd += clName;
- compileCmd += " ";
diff --git a/debian/patches/device-info-query b/debian/patches/device-info-query
deleted file mode 100644
index 38a13b2..0000000
--- a/debian/patches/device-info-query
+++ /dev/null
@@ -1,54 +0,0 @@
-commit ab7af52f7e9cd4d6b5a7f1fce4552c3e97b49099
-Author: Simon Richter <Simon.Richter at hogyros.de>
-Date: Thu Apr 11 13:59:30 2013 +0200
-
- Allow size queries in device_id accessors
-
- Query functions with variable output length can be called with an output
- buffer length and address of zero to find out the number of elements
- available.
-
-Index: beignet-0.0.0+git2013.04.11+e6b503e/src/cl_device_id.c
-===================================================================
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/src/cl_device_id.c 2013-04-12 08:13:48.000000000 +0200
-+++ beignet-0.0.0+git2013.04.11+e6b503e/src/cl_device_id.c 2013-04-15 18:34:02.220354446 +0200
-@@ -131,21 +131,25 @@
-
- #define DECL_FIELD(CASE,FIELD) \
- case JOIN(CL_DEVICE_,CASE): \
-- if (param_value_size < sizeof(((cl_device_id)NULL)->FIELD)) \
-- return CL_INVALID_VALUE; \
- if (param_value_size_ret != NULL) \
- *param_value_size_ret = sizeof(((cl_device_id)NULL)->FIELD);\
-+ if (param_value == NULL) \
-+ return CL_SUCCESS; \
-+ if (param_value_size < sizeof(((cl_device_id)NULL)->FIELD)) \
-+ return CL_INVALID_VALUE; \
- memcpy(param_value, \
- &device->FIELD, \
- sizeof(((cl_device_id)NULL)->FIELD)); \
-- return CL_SUCCESS;
-+ return CL_SUCCESS;
-
- #define DECL_STRING_FIELD(CASE,FIELD) \
- case JOIN(CL_DEVICE_,CASE): \
-- if (param_value_size < device->JOIN(FIELD,_sz)) \
-- return CL_INVALID_VALUE; \
- if (param_value_size_ret != NULL) \
- *param_value_size_ret = device->JOIN(FIELD,_sz); \
-+ if (param_value == NULL) \
-+ return CL_SUCCESS; \
-+ if (param_value_size < device->JOIN(FIELD,_sz)) \
-+ return CL_INVALID_VALUE; \
- memcpy(param_value, device->FIELD, device->JOIN(FIELD,_sz)); \
- return CL_SUCCESS;
-
-@@ -282,8 +286,6 @@
- if (UNLIKELY(device != &intel_ivb_gt1_device &&
- device != &intel_ivb_gt2_device))
- return CL_INVALID_DEVICE;
-- if (UNLIKELY(param_value == NULL))
-- return CL_INVALID_VALUE;
-
- switch (param_name) {
- DECL_FIELD(WORK_GROUP_SIZE, wg_sz)
diff --git a/debian/patches/fix-clean b/debian/patches/fix-clean
deleted file mode 100644
index 4f9b581..0000000
--- a/debian/patches/fix-clean
+++ /dev/null
@@ -1,23 +0,0 @@
-Description: Fix cleaning
-Author: Simon Richter <sjr at debian.org>
-Last-Update: 2013-04-01
-
-Index: beignet-0.0.0+git2013.04.01+d1b234c/backend/Makefile
-===================================================================
---- beignet-0.0.0+git2013.04.01+d1b234c.orig/backend/Makefile 2013-04-01 04:00:50.847138659 +0200
-+++ beignet-0.0.0+git2013.04.01+d1b234c/backend/Makefile 2013-04-01 04:00:55.667138444 +0200
-@@ -1,4 +1,4 @@
- TOP=.
--SUBDIRS=src/backend src/backend/gen src/backend/sim src/ir src/llvm src/sys
-+SUBDIRS=src src/backend src/backend/gen src/backend/sim src/ir src/llvm src/sys
-
- include $(TOP)/Makefile.shared
---- beignet-0.0.0+git2013.04.01+d1b234c.orig/utests/Makefile
-+++ beignet-0.0.0+git2013.04.01+d1b234c/utests/Makefile
-@@ -74,5 +74,6 @@ runtime_flat_address_space: $(OBJ) runti
-
- clean:
- rm -f $(OBJ)
-+ rm -f *.bmp
- rm -f utest_run utest_run.o
- rm -f runtime_flat_address_space runtime_flat_address_space.o
diff --git a/debian/patches/glibc-memalign b/debian/patches/glibc-memalign
deleted file mode 100644
index 653799c..0000000
--- a/debian/patches/glibc-memalign
+++ /dev/null
@@ -1,19 +0,0 @@
-commit 3be62a8aacd8cd18484b381ab9536160c7de7119
-Author: Simon Richter <Simon.Richter at hogyros.de>
-Date: Fri Apr 12 11:21:19 2013 +0200
-
- Accept glibc's implementation of memalign()
-
-Index: beignet-0.0.0+git2013.04.11+e6b503e/backend/src/sys/alloc.cpp
-===================================================================
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/backend/src/sys/alloc.cpp 2013-04-12 08:13:48.000000000 +0200
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/src/sys/alloc.cpp 2013-04-15 18:34:31.608356151 +0200
-@@ -271,7 +271,7 @@
- /// Linux Platform
- ////////////////////////////////////////////////////////////////////////////////
-
--#if defined(__LINUX__)
-+#if defined(__LINUX__) || defined(__GLIBC__)
-
- #include <unistd.h>
- #include <sys/mman.h>
diff --git a/debian/patches/ignore-missing-egl b/debian/patches/ignore-missing-egl
deleted file mode 100644
index 19f0721..0000000
--- a/debian/patches/ignore-missing-egl
+++ /dev/null
@@ -1,50 +0,0 @@
-Description: Ignore missing EGL support (version in Debian is too old)
-Author: Simon Richter <sjr at debian.org>
-Last-Update: 2013-04-15
-
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/backend/src/backend/program.cpp
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/src/backend/program.cpp
-@@ -118,7 +118,7 @@ namespace gbe {
- fclose(clFile);
-
- // Now compile the code to llvm using clang
--#if LLVM_VERSION_MINOR <= 1
-+#if 1 //LLVM_VERSION_MINOR <= 1
- std::string compileCmd = "clang -x cl -fno-color-diagnostics -emit-llvm -O3 -ccc-host-triple ptx32 -c ";
- #else
- std::string compileCmd = "clang -ffp-contract=off -emit-llvm -O3 -target nvptx -x cl -c ";
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/src/CMakeLists.txt
-+++ beignet-0.0.0+git2013.04.11+e6b503e/src/CMakeLists.txt
-@@ -32,6 +32,9 @@ if (EGL_FOUND)
- set (OPENCL_SRC ${OPENCL_SRC} cl_mem_gl.c cl_gl_api.c x11/gbm_dri2_x11_platform.c)
- SET(CMAKE_CXX_FLAGS "-DHAS_EGL ${CMAKE_CXX_FLAGS}")
- SET(CMAKE_C_FLAGS "-DHAS_EGL ${CMAKE_C_FLAGS}")
-+SET(OPTIONAL_EGL_LIBRARY "${EGL_LIBRARY}")
-+else(EGL_FOUND)
-+SET(OPTIONAL_EGL_LIBRARY "")
- endif (EGL_FOUND)
-
- link_directories (${LLVM_LIBRARY_DIR})
-@@ -45,6 +48,6 @@ target_link_libraries(
- ${DRM_INTEL_LIBRARY}
- ${DRM_LIBRARY}
- ${OPENGL_LIBRARIES}
-- ${EGL_LIBRARY}
-+ ${OPTIONAL_EGL_LIBRARY}
- ${GBM_LIBRARY})
- install (TARGETS cl LIBRARY DESTINATION lib)
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/utests/utest_helper.hpp
-+++ beignet-0.0.0+git2013.04.11+e6b503e/utests/utest_helper.hpp
-@@ -128,9 +128,9 @@ extern cl_mem buf[MAX_BUFFER_N];
- extern void* buf_data[MAX_BUFFER_N];
- extern size_t globals[3];
- extern size_t locals[3];
--extern Display *xDisplay;
--extern EGLDisplay eglDisplay;
--extern EGLSurface eglSurface;
-+//extern Display *xDisplay;
-+//extern EGLDisplay eglDisplay;
-+//extern EGLSurface eglSurface;
-
-
- enum {
diff --git a/debian/patches/implement-gefa b/debian/patches/implement-gefa
index 39c9764..09c6362 100644
--- a/debian/patches/implement-gefa
+++ b/debian/patches/implement-gefa
@@ -6,13 +6,13 @@ Date: Tue Apr 2 15:11:01 2013 +0200
This function can legally return NULL in the current implementation.
-Index: beignet-0.0.0+git2013.04.11+e6b503e/src/cl_api.c
+Index: beignet-0.1+git20130418+0546d2e/src/cl_api.c
===================================================================
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/src/cl_api.c 2013-04-12 08:13:48.000000000 +0200
-+++ beignet-0.0.0+git2013.04.11+e6b503e/src/cl_api.c 2013-04-15 18:33:32.788352738 +0200
-@@ -1169,7 +1169,7 @@
- void*
- clGetExtensionFunctionAddress(const char *func_name)
+--- beignet-0.1+git20130418+0546d2e.orig/src/cl_api.c 2013-04-18 05:21:35.000000000 +0200
++++ beignet-0.1+git20130418+0546d2e/src/cl_api.c 2013-04-18 11:52:29.225428977 +0200
+@@ -1006,7 +1006,7 @@
+ cl_event * event,
+ cl_int * errcode_ret)
{
- NOT_IMPLEMENTED;
+ /* No extensions supported at present */
diff --git a/debian/patches/llvm-3.2 b/debian/patches/llvm-3.2
deleted file mode 100644
index 22d5b2f..0000000
--- a/debian/patches/llvm-3.2
+++ /dev/null
@@ -1,19 +0,0 @@
-Description: Look harder for LLVM 3.2
-Author: Simon Richter <sjr at debian.org>
-Last-Update: 2013-04-15
-
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/CMake/FindLLVM.cmake
-+++ beignet-0.0.0+git2013.04.11+e6b503e/CMake/FindLLVM.cmake
-@@ -7,9 +7,9 @@
- # LLVM_MODULE_LIBS - list of llvm libs for working with modules.
- # LLVM_FOUND - True if llvm found.
- if (LLVM_INSTALL_DIR)
-- find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config llvm-config-32 llvm-config-3.0 DOC "llvm-config executable" PATHS ${LLVM_INSTALL_DIR} NO_DEFAULT_PATH)
--else (LLVM_INSTALL_DIR)
-- find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config llvm-config-32 llvm-config-3.0 DOC "llvm-config executable")
-+ find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config-32 llvm-config-3.2 llvm-config DOC "llvm-config executable" PATHS ${LLVM_INSTALL_DIR} NO_DEFAULT_PATH)
-+else (LLVM_INSTALL_DIR)
-+ find_program(LLVM_CONFIG_EXECUTABLE NAMES llvm-config-32 llvm-config-3.2 llvm-config DOC "llvm-config executable")
- endif (LLVM_INSTALL_DIR)
-
- if (LLVM_CONFIG_EXECUTABLE)
diff --git a/debian/patches/missing-header b/debian/patches/missing-header
index adfcafe..dab28b7 100644
--- a/debian/patches/missing-header
+++ b/debian/patches/missing-header
@@ -2,10 +2,10 @@ Description: Fix missing include
Author: Simon Richter <sjr at debian.org>
Last-Update: 2013-04-03
-Index: beignet-0.0.0+git2013.04.11+e6b503e/backend/src/sys/alloc.hpp
+Index: beignet-0.1+git20130418+0546d2e/backend/src/sys/alloc.hpp
===================================================================
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/backend/src/sys/alloc.hpp 2013-04-12 08:13:48.000000000 +0200
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/src/sys/alloc.hpp 2013-04-15 18:33:38.340353060 +0200
+--- beignet-0.1+git20130418+0546d2e.orig/backend/src/sys/alloc.hpp 2013-04-18 05:21:35.000000000 +0200
++++ beignet-0.1+git20130418+0546d2e/backend/src/sys/alloc.hpp 2013-04-18 11:53:00.793430809 +0200
@@ -27,6 +27,7 @@
#include "sys/platform.hpp"
#include "sys/assert.hpp"
diff --git a/debian/patches/path b/debian/patches/path
deleted file mode 100644
index dfdf558..0000000
--- a/debian/patches/path
+++ /dev/null
@@ -1,21 +0,0 @@
-Description: Fix linker path
-Author: Simon Richter <sjr at debian.org>
-Last-Update: 2013-04-01
-
-Index: beignet-0.0.0+git2013.04.01+d1b234c/utests/Makefile
-===================================================================
---- beignet-0.0.0+git2013.04.01+d1b234c.orig/utests/Makefile 2013-04-04 13:12:58.000000000 +0200
-+++ beignet-0.0.0+git2013.04.01+d1b234c/utests/Makefile 2013-04-04 13:13:31.575880044 +0200
-@@ -67,10 +67,10 @@
- all: utest_run runtime_flat_address_space
-
- utest_run: $(OBJ) utest_run.o $(TOP)/$(LIBBASE)
-- $(CXX) -o $@ $(OBJ) utest_run.o $(TOP)/$(LIBBASE)
-+ $(CXX) -o $@ $(OBJ) utest_run.o -L$(TOP)/backend -Wl,-rpath,$(TOP)/backend -Wl,-rpath,$(TOP) $(TOP)/$(LIBBASE)
-
- runtime_flat_address_space: $(OBJ) runtime_flat_address_space.o
-- $(CXX) -o $@ $(OBJ) runtime_flat_address_space.o $(TOP)/$(LIBBASE)
-+ $(CXX) -o $@ $(OBJ) runtime_flat_address_space.o -Wl,-rpath,$(TOP)/backend -Wl,-rpath,$(TOP) $(TOP)/$(LIBBASE)
-
- clean:
- rm -f $(OBJ)
diff --git a/debian/patches/respect-flags b/debian/patches/respect-flags
index 726c46b..cc57dcb 100644
--- a/debian/patches/respect-flags
+++ b/debian/patches/respect-flags
@@ -2,10 +2,10 @@ Description: Respect CFLAGS/CXXFLAGS from Debian
Author: Simon Richter <sjr at debian.org>
Last-Update: 2013-04-16
-Index: beignet-0.0.0+git2013.04.11+e6b503e/CMakeLists.txt
+Index: beignet-0.1+git20130418+0546d2e/CMakeLists.txt
===================================================================
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/CMakeLists.txt 2013-04-16 14:49:50.450369637 +0200
-+++ beignet-0.0.0+git2013.04.11+e6b503e/CMakeLists.txt 2013-04-16 14:50:20.298368306 +0200
+--- beignet-0.1+git20130418+0546d2e.orig/CMakeLists.txt 2013-04-18 05:21:35.000000000 +0200
++++ beignet-0.1+git20130418+0546d2e/CMakeLists.txt 2013-04-18 11:54:52.065437266 +0200
@@ -18,7 +18,6 @@
INCLUDE_DIRECTORIES(${CMAKE_CURRENT_BINARY_DIR} ${CMAKE_CURRENT_SOURCE_DIR})
@@ -18,20 +18,20 @@ Index: beignet-0.0.0+git2013.04.11+e6b503e/CMakeLists.txt
ADD_DEFINITIONS(-DUSE_FULSIM=0)
ENDIF (USE_FULSIM)
--SET(CMAKE_CXX_FLAGS "-Wall -Wno-invalid-offsetof -mfpmath=sse --no-rtti -Wcast-align -std=c++0x -msse2 -msse3 -mssse3 -msse4.1 ")
+-SET(CMAKE_CXX_FLAGS "-Wall -Wno-invalid-offsetof -mfpmath=sse -fno-rtti -Wcast-align -std=c++0x -msse2 -msse3 -mssse3 -msse4.1 ")
-SET(CMAKE_C_FLAGS "-Wall -mfpmath=sse -msse2 -Wcast-align -msse2 -msse3 -mssse3 -msse4.1")
-+SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wno-invalid-offsetof -mfpmath=sse --no-rtti -Wcast-align -std=c++0x -msse2 -msse3 -mssse3 -msse4.1 ")
++SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wno-invalid-offsetof -mfpmath=sse -fno-rtti -Wcast-align -std=c++0x -msse2 -msse3 -mssse3 -msse4.1 ")
+SET(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Wall -mfpmath=sse -msse2 -Wcast-align -msse2 -msse3 -mssse3 -msse4.1")
# Front end stuff we need
#INCLUDE(CMake/FindLLVM.cmake)
-Index: beignet-0.0.0+git2013.04.11+e6b503e/backend/CMakeLists.txt
+Index: beignet-0.1+git20130418+0546d2e/backend/CMakeLists.txt
===================================================================
---- beignet-0.0.0+git2013.04.11+e6b503e.orig/backend/CMakeLists.txt 2013-04-16 14:49:48.290369733 +0200
-+++ beignet-0.0.0+git2013.04.11+e6b503e/backend/CMakeLists.txt 2013-04-16 14:50:20.298368306 +0200
+--- beignet-0.1+git20130418+0546d2e.orig/backend/CMakeLists.txt 2013-04-18 05:21:35.000000000 +0200
++++ beignet-0.1+git20130418+0546d2e/backend/CMakeLists.txt 2013-04-18 11:54:19.629435384 +0200
@@ -45,16 +45,16 @@
if (COMPILER STREQUAL "GCC")
- set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -funroll-loops -Wstrict-aliasing=2 -fstrict-aliasing -msse2 -msse3 -mssse3 -msse4.1 -ffast-math -fPIC -Wall")
+ set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -funroll-loops -Wstrict-aliasing=2 -fstrict-aliasing -msse2 -msse3 -mssse3 -msse4.1 -fPIC -Wall")
set (CMAKE_C_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} ${LLVM_CFLAGS}")
- set (CMAKE_CXX_FLAGS "${CMAKE_C_CXX_FLAGS} -Wno-invalid-offsetof -fno-rtti -std=c++0x")
+ set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${CMAKE_C_CXX_FLAGS} -Wno-invalid-offsetof -fno-rtti -std=c++0x")
diff --git a/debian/patches/series b/debian/patches/series
index 6c20ce6..110055d 100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@ -1,11 +1,6 @@
khronos
verbose
-clang-from-path
implement-gefa
missing-header
-glibc-memalign
-llvm-3.2
-autogen
-ignore-missing-egl
soname
respect-flags
diff --git a/debian/patches/soname b/debian/patches/soname
index 2e270fd..32597fa 100644
--- a/debian/patches/soname
+++ b/debian/patches/soname
@@ -2,10 +2,10 @@ Description: Use proper SONAME
Author: Simon Richter <sjr at debian.org>
Last-Update: 2013-04-15
-Index: beignet-0.1/backend/src/CMakeLists.txt
+Index: beignet-0.1+git20130418+0546d2e/backend/src/CMakeLists.txt
===================================================================
---- beignet-0.1.orig/backend/src/CMakeLists.txt 2013-04-12 08:13:48.000000000 +0200
-+++ beignet-0.1/backend/src/CMakeLists.txt 2013-04-16 17:20:10.497967484 +0200
+--- beignet-0.1+git20130418+0546d2e.orig/backend/src/CMakeLists.txt 2013-04-18 05:21:35.000000000 +0200
++++ beignet-0.1+git20130418+0546d2e/backend/src/CMakeLists.txt 2013-04-18 11:53:59.985434244 +0200
@@ -114,7 +114,10 @@
${LLVM_MODULE_LIBS}
${CMAKE_THREAD_LIBS_INIT}
@@ -18,11 +18,11 @@ Index: beignet-0.1/backend/src/CMakeLists.txt
install (TARGETS gbe LIBRARY DESTINATION lib)
install (FILES backend/program.h DESTINATION include/gen)
-Index: beignet-0.1/src/CMakeLists.txt
+Index: beignet-0.1+git20130418+0546d2e/src/CMakeLists.txt
===================================================================
---- beignet-0.1.orig/src/CMakeLists.txt 2013-04-16 17:18:36.000000000 +0200
-+++ beignet-0.1/src/CMakeLists.txt 2013-04-16 17:19:57.913968045 +0200
-@@ -50,4 +50,8 @@
+--- beignet-0.1+git20130418+0546d2e.orig/src/CMakeLists.txt 2013-04-18 05:21:35.000000000 +0200
++++ beignet-0.1+git20130418+0546d2e/src/CMakeLists.txt 2013-04-18 11:53:59.985434244 +0200
+@@ -58,4 +58,8 @@
${OPENGL_LIBRARIES}
${OPTIONAL_EGL_LIBRARY}
${GBM_LIBRARY})
--
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