[Pkg-opencl-devel] [beignet] 17/47: Imported Debian patch 0.1+git20130614+89b5e40-1
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 21:45:48 UTC 2014
This is an automated email from the git hooks/post-receive script.
anbe pushed a commit to branch master
in repository beignet.
commit b897c2c1fe6174561d11b3013224dba2fe92a812
Author: Simon Richter <sjr at debian.org>
Date: Fri Jun 14 15:22:18 2013 +0200
Imported Debian patch 0.1+git20130614+89b5e40-1
---
debian/changelog | 6 +
...0001-Generate-all-supported-as_-functions.patch | 106 +++++----
.../0002-Define-all-convert_-functions.patch | 115 +++++----
...g-and-ulong-types-to-generated-functions.patch} | 145 ++++++------
.../0004-Add-vector-argument-test-case.patch | 8 +-
.../0005-Add-more-get-image-info-functions.patch | 203 ----------------
...nt-get_image_size-cases-to-other-informat.patch | 180 --------------
...ge-clang-system-call-to-libclang-api-call.patch | 261 ---------------------
debian/patches/series | 5 +-
9 files changed, 221 insertions(+), 808 deletions(-)
diff --git a/debian/changelog b/debian/changelog
index 4ea12cc..c66d98c 100644
--- a/debian/changelog
+++ b/debian/changelog
@@ -1,3 +1,9 @@
+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
diff --git a/debian/patches/0001-Generate-all-supported-as_-functions.patch b/debian/patches/0001-Generate-all-supported-as_-functions.patch
index 5520e63..45276c2 100644
--- a/debian/patches/0001-Generate-all-supported-as_-functions.patch
+++ b/debian/patches/0001-Generate-all-supported-as_-functions.patch
@@ -1,36 +1,35 @@
-From 67668ad9be1186f021a7ee2dd56d00297593b6d0 Mon Sep 17 00:00:00 2001
+From a1926ba22c15aee973d651d700fdc7b94cd8bf4d Mon Sep 17 00:00:00 2001
From: Simon Richter <Simon.Richter at hogyros.de>
Date: Mon, 13 May 2013 22:43:34 +0200
-Subject: [PATCH 1/7] Generate all supported as_* functions
+Subject: [PATCH 1/4] Generate all supported as_* functions
To: beignet at lists.freedesktop.org
-This adds support for all type conversions currently possible.
+This adds support for all reinterpreting type conversions currently
+possible.
The conversion functions can be updated by invoking the
-update_conversions.sh script.
+update_as.sh script.
---
- backend/src/gen_conversions.sh | 89 +++
- backend/src/ocl_stdlib.h | 1138 ++++++++++++++++++++++++++++++++++++-
- backend/src/update_conversions.sh | 11 +
- 3 files changed, 1231 insertions(+), 7 deletions(-)
- create mode 100755 backend/src/gen_conversions.sh
- create mode 100755 backend/src/update_conversions.sh
+ backend/src/gen_as.sh | 83 ++++
+ backend/src/genconfig.sh | 11 +
+ backend/src/ocl_stdlib.h | 1137 +++++++++++++++++++++++++++++++++++++++++++++-
+ backend/src/update.sh | 2 +
+ backend/src/update_as.sh | 11 +
+ 5 files changed, 1237 insertions(+), 7 deletions(-)
+ create mode 100755 backend/src/gen_as.sh
+ create mode 100644 backend/src/genconfig.sh
+ create mode 100755 backend/src/update.sh
+ create mode 100755 backend/src/update_as.sh
-diff --git a/backend/src/gen_conversions.sh b/backend/src/gen_conversions.sh
+diff --git a/backend/src/gen_as.sh b/backend/src/gen_as.sh
new file mode 100755
-index 0000000..50a3668
+index 0000000..76fedf8
--- /dev/null
-+++ b/backend/src/gen_conversions.sh
-@@ -0,0 +1,89 @@
++++ b/backend/src/gen_as.sh
+@@ -0,0 +1,83 @@
+#! /bin/sh -e
+
-+# Supported base types and their lengths
-+TYPES="int:4 uint:4 short:2 ushort:2 char:1 uchar:1 float:4"
-+
-+# Supported vector lengths
-+VECTOR_LENGTHS="1 2 3 4 8 16"
-+
-+## No user serviceable parts below here
++. ./genconfig.sh
+
+# Generate list of union sizes
+for type in $TYPES; do
@@ -111,15 +110,28 @@ index 0000000..50a3668
+ done
+
+done
+diff --git a/backend/src/genconfig.sh b/backend/src/genconfig.sh
+new file mode 100644
+index 0000000..60edafd
+--- /dev/null
++++ b/backend/src/genconfig.sh
+@@ -0,0 +1,11 @@
++#! /bin/false
++# This is to be sourced by the generation scripts
++
++# Supported base types and their lengths
++TYPES="int:4 uint:4 short:2 ushort:2 char:1 uchar:1 float:4"
++
++# Supported vector lengths
++VECTOR_LENGTHS="1 2 3 4 8 16"
++
++## No user serviceable parts below here
++
diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
-index 92f9ba9..a051803 100644
+index 46b81e1..d00de44 100644
--- a/backend/src/ocl_stdlib.h
+++ b/backend/src/ocl_stdlib.h
-@@ -78,18 +78,1142 @@ struct _image3d_t;
- typedef __texture struct _image3d_t* image3d_t;
- typedef uint sampler_t;
- typedef size_t event_t;
-+
+@@ -91,15 +91,1138 @@ typedef size_t __event_t;
/////////////////////////////////////////////////////////////////////////////
// OpenCL conversions & type casting
/////////////////////////////////////////////////////////////////////////////
@@ -127,7 +139,7 @@ index 92f9ba9..a051803 100644
- float f;
- uchar4 u4;
+
-+// ##BEGIN_CONVERSIONS##
++// ##BEGIN_AS##
+union _type_cast_1_b {
+ char _char;
+ uchar _uchar;
@@ -150,11 +162,7 @@ index 92f9ba9..a051803 100644
+ ushort _ushort;
+ char2 _char2;
+ uchar2 _uchar2;
- };
--uchar4 INLINE_OVERLOADABLE as_uchar4(float f) {
-- union type_cast_4_b u;
-- u.f = f;
-- return u.u4;
++};
+
+INLINE OVERLOADABLE ushort as_ushort(short v) {
+ union _type_cast_2_b u;
@@ -700,7 +708,7 @@ index 92f9ba9..a051803 100644
+ union _type_cast_8_b u;
+ u._char8 = v;
+ return u._ushort4;
- }
++}
+
+INLINE OVERLOADABLE uchar8 as_uchar8(char8 v) {
+ union _type_cast_8_b u;
@@ -1096,7 +1104,11 @@ index 92f9ba9..a051803 100644
+ short16 _short16;
+ ushort16 _ushort16;
+ float8 _float8;
-+};
+ };
+-uchar4 INLINE_OVERLOADABLE as_uchar4(float f) {
+- union type_cast_4_b u;
+- u.f = f;
+- return u.u4;
+
+INLINE OVERLOADABLE uint8 as_uint8(int8 v) {
+ union _type_cast_32_b u;
@@ -1126,7 +1138,7 @@ index 92f9ba9..a051803 100644
+ union _type_cast_32_b u;
+ u._uint8 = v;
+ return u._int8;
-+}
+ }
+
+INLINE OVERLOADABLE short16 as_short16(uint8 v) {
+ union _type_cast_32_b u;
@@ -1260,25 +1272,33 @@ index 92f9ba9..a051803 100644
+ return u._uint16;
+}
+
-+// ##END_CONVERSIONS##
++// ##END_AS##
+
#define DEF(type, n, type2) type##n INLINE_OVERLOADABLE convert_##type##n(type2##n d) { \
return (type##n)((type)(d.s0), (type)(d.s1), (type)(d.s2), (type)(d.s3)); \
}
-diff --git a/backend/src/update_conversions.sh b/backend/src/update_conversions.sh
+diff --git a/backend/src/update.sh b/backend/src/update.sh
+new file mode 100755
+index 0000000..4f9af8c
+--- /dev/null
++++ b/backend/src/update.sh
+@@ -0,0 +1,2 @@
++#! /bin/sh -e
++./update_as.sh
+diff --git a/backend/src/update_as.sh b/backend/src/update_as.sh
new file mode 100755
-index 0000000..bd9099f
+index 0000000..54b4191
--- /dev/null
-+++ b/backend/src/update_conversions.sh
++++ b/backend/src/update_as.sh
@@ -0,0 +1,11 @@
+#! /bin/sh -e
+
+STDLIB_HEADER=ocl_stdlib.h
+
+exec >$STDLIB_HEADER.tmp
-+sed -n -e '1,/##BEGIN_CONVERSIONS##/p' $STDLIB_HEADER
-+./gen_conversions.sh
-+sed -n -e '/##END_CONVERSIONS##/,$p' $STDLIB_HEADER
++sed -n -e '1,/##BEGIN_AS##/p' $STDLIB_HEADER
++./gen_as.sh
++sed -n -e '/##END_AS##/,$p' $STDLIB_HEADER
+exec >&2
+
+mv $STDLIB_HEADER.tmp $STDLIB_HEADER
diff --git a/debian/patches/0002-Define-all-convert_-functions.patch b/debian/patches/0002-Define-all-convert_-functions.patch
index f47e282..a0e5c20 100644
--- a/debian/patches/0002-Define-all-convert_-functions.patch
+++ b/debian/patches/0002-Define-all-convert_-functions.patch
@@ -1,24 +1,29 @@
-From dd0b29265577ad8cf2cc4373bf73a6f549a8263d Mon Sep 17 00:00:00 2001
+From 1900bf07f138edbf956e01618c304a7b10c59a9b Mon Sep 17 00:00:00 2001
From: Simon Richter <Simon.Richter at hogyros.de>
Date: Tue, 14 May 2013 17:04:56 +0200
-Subject: [PATCH 2/7] Define all convert_* functions.
+Subject: [PATCH 2/4] Define all convert_* functions.
To: beignet at lists.freedesktop.org
These functions convert between vectors of the same length by casting each
member in turn.
---
- backend/src/gen_conversions.sh | 49 +++
- backend/src/ocl_stdlib.h | 863 ++++++++++++++++++++++++++++++++++++++--
- 2 files changed, 889 insertions(+), 23 deletions(-)
+ backend/src/gen_convert.sh | 52 +++
+ backend/src/ocl_stdlib.h | 866 +++++++++++++++++++++++++++++++++++++++--
+ backend/src/update.sh | 1 +
+ backend/src/update_convert.sh | 11 +
+ 4 files changed, 907 insertions(+), 23 deletions(-)
+ create mode 100755 backend/src/gen_convert.sh
+ create mode 100755 backend/src/update_convert.sh
-diff --git a/backend/src/gen_conversions.sh b/backend/src/gen_conversions.sh
-index 50a3668..d0b1b58 100755
---- a/backend/src/gen_conversions.sh
-+++ b/backend/src/gen_conversions.sh
-@@ -87,3 +87,52 @@ for union_size in $union_sizes; do
- done
-
- done
+diff --git a/backend/src/gen_convert.sh b/backend/src/gen_convert.sh
+new file mode 100755
+index 0000000..74fc73c
+--- /dev/null
++++ b/backend/src/gen_convert.sh
+@@ -0,0 +1,52 @@
++#! /bin/sh -e
++
++. ./genconfig.sh
+
+# For all vector lengths and types, generate conversion functions
+for vector_length in $VECTOR_LENGTHS; do
@@ -69,13 +74,37 @@ index 50a3668..d0b1b58 100755
+ done
+done
diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
-index a051803..4e0459a 100644
+index d00de44..913917f 100644
--- a/backend/src/ocl_stdlib.h
+++ b/backend/src/ocl_stdlib.h
-@@ -1212,31 +1212,848 @@ INLINE OVERLOADABLE uint16 as_uint16(float16 v) {
- return u._uint16;
- }
+@@ -1223,29 +1223,849 @@ INLINE OVERLOADABLE uint16 as_uint16(float16 v) {
+ // ##END_AS##
+
+-#define DEF(type, n, type2) type##n INLINE_OVERLOADABLE convert_##type##n(type2##n d) { \
+- return (type##n)((type)(d.s0), (type)(d.s1), (type)(d.s2), (type)(d.s3)); \
+- }
+-#define DEF2(type) DEF(type, 4, char); \
+- DEF(type, 4, uchar); \
+- DEF(type, 4, short); \
+- DEF(type, 4, ushort); \
+- DEF(type, 4, int); \
+- DEF(type, 4, uint); \
+- DEF(type, 4, long); \
+- DEF(type, 4, ulong); \
+- DEF(type, 4, float);
+-DEF2(char);
+-DEF2(uchar);
+-DEF2(short);
+-DEF2(ushort);
+-DEF2(int);
+-DEF2(uint);
+-DEF2(long);
+-DEF2(ulong);
+-DEF2(float);
+-#undef DEF2
+-#undef DEF
++// ##BEGIN_CONVERT##
+INLINE OVERLOADABLE uint2 convert_uint2(int2 v) {
+ return (uint2)((uint)(v.s0), (uint)(v.s1));
+}
@@ -916,34 +945,36 @@ index a051803..4e0459a 100644
+ return (uchar16)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3), (uchar)(v.s4), (uchar)(v.s5), (uchar)(v.s6), (uchar)(v.s7), (uchar)(v.s8), (uchar)(v.s9), (uchar)(v.sA), (uchar)(v.sB), (uchar)(v.sC), (uchar)(v.sD), (uchar)(v.sE), (uchar)(v.sF));
+}
+
- // ##END_CONVERSIONS##
-
--#define DEF(type, n, type2) type##n INLINE_OVERLOADABLE convert_##type##n(type2##n d) { \
-- return (type##n)((type)(d.s0), (type)(d.s1), (type)(d.s2), (type)(d.s3)); \
-- }
--#define DEF2(type) DEF(type, 4, char); \
-- DEF(type, 4, uchar); \
-- DEF(type, 4, short); \
-- DEF(type, 4, ushort); \
-- DEF(type, 4, int); \
-- DEF(type, 4, uint); \
-- DEF(type, 4, long); \
-- DEF(type, 4, ulong); \
-- DEF(type, 4, float);
--DEF2(char);
--DEF2(uchar);
--DEF2(short);
--DEF2(ushort);
--DEF2(int);
--DEF2(uint);
--DEF2(long);
--DEF2(ulong);
--DEF2(float);
--#undef DEF2
--#undef DEF
++// ##END_CONVERT##
++
/////////////////////////////////////////////////////////////////////////////
// OpenCL preprocessor directives & macros
/////////////////////////////////////////////////////////////////////////////
+diff --git a/backend/src/update.sh b/backend/src/update.sh
+index 4f9af8c..0e5f8c0 100755
+--- a/backend/src/update.sh
++++ b/backend/src/update.sh
+@@ -1,2 +1,3 @@
+ #! /bin/sh -e
+ ./update_as.sh
++./update_convert.sh
+diff --git a/backend/src/update_convert.sh b/backend/src/update_convert.sh
+new file mode 100755
+index 0000000..f1fcd36
+--- /dev/null
++++ b/backend/src/update_convert.sh
+@@ -0,0 +1,11 @@
++#! /bin/sh -e
++
++STDLIB_HEADER=ocl_stdlib.h
++
++exec >$STDLIB_HEADER.tmp
++sed -n -e '1,/##BEGIN_CONVERT##/p' $STDLIB_HEADER
++./gen_convert.sh
++sed -n -e '/##END_CONVERT##/,$p' $STDLIB_HEADER
++exec >&2
++
++mv $STDLIB_HEADER.tmp $STDLIB_HEADER
--
1.7.10.4
diff --git a/debian/patches/0003-Add-long-and-ulong-types-to-conversions.patch b/debian/patches/0003-Add-long-and-ulong-types-to-generated-functions.patch
similarity index 92%
rename from debian/patches/0003-Add-long-and-ulong-types-to-conversions.patch
rename to debian/patches/0003-Add-long-and-ulong-types-to-generated-functions.patch
index c2253ee..7e7f81d 100644
--- a/debian/patches/0003-Add-long-and-ulong-types-to-conversions.patch
+++ b/debian/patches/0003-Add-long-and-ulong-types-to-generated-functions.patch
@@ -1,21 +1,21 @@
-From 809e0d28d8b628885d1024020b5f479154e2e915 Mon Sep 17 00:00:00 2001
+From 44161ff1568479390464c0b0a282f5aeeb86915d Mon Sep 17 00:00:00 2001
From: Simon Richter <Simon.Richter at hogyros.de>
Date: Tue, 14 May 2013 17:04:57 +0200
-Subject: [PATCH 3/7] Add long and ulong types to conversions.
+Subject: [PATCH 3/4] Add long and ulong types to generated functions.
To: beignet at lists.freedesktop.org
-This enables as_* and convert_* for the long and ulong data types.
+This enables all generated functions for 64 bit integers.
---
- backend/src/gen_conversions.sh | 2 +-
- backend/src/ocl_stdlib.h | 1248 +++++++++++++++++++++++++++++++++++++++-
+ backend/src/genconfig.sh | 2 +-
+ backend/src/ocl_stdlib.h | 1248 +++++++++++++++++++++++++++++++++++++++++++++-
2 files changed, 1234 insertions(+), 16 deletions(-)
-diff --git a/backend/src/gen_conversions.sh b/backend/src/gen_conversions.sh
-index d0b1b58..083fc38 100755
---- a/backend/src/gen_conversions.sh
-+++ b/backend/src/gen_conversions.sh
-@@ -1,7 +1,7 @@
- #! /bin/sh -e
+diff --git a/backend/src/genconfig.sh b/backend/src/genconfig.sh
+index 60edafd..a3ba3f9 100644
+--- a/backend/src/genconfig.sh
++++ b/backend/src/genconfig.sh
+@@ -2,7 +2,7 @@
+ # This is to be sourced by the generation scripts
# Supported base types and their lengths
-TYPES="int:4 uint:4 short:2 ushort:2 char:1 uchar:1 float:4"
@@ -24,10 +24,10 @@ index d0b1b58..083fc38 100755
# Supported vector lengths
VECTOR_LENGTHS="1 2 3 4 8 16"
diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
-index 4e0459a..0d09aff 100644
+index 913917f..2f55184 100644
--- a/backend/src/ocl_stdlib.h
+++ b/backend/src/ocl_stdlib.h
-@@ -477,6 +477,8 @@ INLINE OVERLOADABLE short3 as_short3(ushort3 v) {
+@@ -486,6 +486,8 @@ INLINE OVERLOADABLE short3 as_short3(ushort3 v) {
}
union _type_cast_8_b {
@@ -36,7 +36,7 @@ index 4e0459a..0d09aff 100644
int2 _int2;
uint2 _uint2;
short4 _short4;
-@@ -486,6 +488,114 @@ union _type_cast_8_b {
+@@ -495,6 +497,114 @@ union _type_cast_8_b {
float2 _float2;
};
@@ -151,7 +151,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE uint2 as_uint2(int2 v) {
union _type_cast_8_b u;
u._int2 = v;
-@@ -522,6 +632,18 @@ INLINE OVERLOADABLE float2 as_float2(int2 v) {
+@@ -531,6 +641,18 @@ INLINE OVERLOADABLE float2 as_float2(int2 v) {
return u._float2;
}
@@ -170,7 +170,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 as_int2(uint2 v) {
union _type_cast_8_b u;
u._uint2 = v;
-@@ -558,6 +680,18 @@ INLINE OVERLOADABLE float2 as_float2(uint2 v) {
+@@ -567,6 +689,18 @@ INLINE OVERLOADABLE float2 as_float2(uint2 v) {
return u._float2;
}
@@ -189,7 +189,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 as_int2(short4 v) {
union _type_cast_8_b u;
u._short4 = v;
-@@ -594,6 +728,18 @@ INLINE OVERLOADABLE float2 as_float2(short4 v) {
+@@ -603,6 +737,18 @@ INLINE OVERLOADABLE float2 as_float2(short4 v) {
return u._float2;
}
@@ -208,7 +208,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 as_int2(ushort4 v) {
union _type_cast_8_b u;
u._ushort4 = v;
-@@ -630,6 +776,18 @@ INLINE OVERLOADABLE float2 as_float2(ushort4 v) {
+@@ -639,6 +785,18 @@ INLINE OVERLOADABLE float2 as_float2(ushort4 v) {
return u._float2;
}
@@ -227,7 +227,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 as_int2(char8 v) {
union _type_cast_8_b u;
u._char8 = v;
-@@ -666,6 +824,18 @@ INLINE OVERLOADABLE float2 as_float2(char8 v) {
+@@ -675,6 +833,18 @@ INLINE OVERLOADABLE float2 as_float2(char8 v) {
return u._float2;
}
@@ -246,7 +246,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 as_int2(uchar8 v) {
union _type_cast_8_b u;
u._uchar8 = v;
-@@ -702,6 +872,18 @@ INLINE OVERLOADABLE float2 as_float2(uchar8 v) {
+@@ -711,6 +881,18 @@ INLINE OVERLOADABLE float2 as_float2(uchar8 v) {
return u._float2;
}
@@ -265,7 +265,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 as_int2(float2 v) {
union _type_cast_8_b u;
u._float2 = v;
-@@ -781,6 +963,8 @@ INLINE OVERLOADABLE uint3 as_uint3(float3 v) {
+@@ -790,6 +972,8 @@ INLINE OVERLOADABLE uint3 as_uint3(float3 v) {
}
union _type_cast_16_b {
@@ -274,7 +274,7 @@ index 4e0459a..0d09aff 100644
int4 _int4;
uint4 _uint4;
short8 _short8;
-@@ -790,6 +974,114 @@ union _type_cast_16_b {
+@@ -799,6 +983,114 @@ union _type_cast_16_b {
float4 _float4;
};
@@ -389,7 +389,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE uint4 as_uint4(int4 v) {
union _type_cast_16_b u;
u._int4 = v;
-@@ -826,6 +1118,18 @@ INLINE OVERLOADABLE float4 as_float4(int4 v) {
+@@ -835,6 +1127,18 @@ INLINE OVERLOADABLE float4 as_float4(int4 v) {
return u._float4;
}
@@ -408,7 +408,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 as_int4(uint4 v) {
union _type_cast_16_b u;
u._uint4 = v;
-@@ -862,6 +1166,18 @@ INLINE OVERLOADABLE float4 as_float4(uint4 v) {
+@@ -871,6 +1175,18 @@ INLINE OVERLOADABLE float4 as_float4(uint4 v) {
return u._float4;
}
@@ -427,7 +427,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 as_int4(short8 v) {
union _type_cast_16_b u;
u._short8 = v;
-@@ -898,6 +1214,18 @@ INLINE OVERLOADABLE float4 as_float4(short8 v) {
+@@ -907,6 +1223,18 @@ INLINE OVERLOADABLE float4 as_float4(short8 v) {
return u._float4;
}
@@ -446,7 +446,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 as_int4(ushort8 v) {
union _type_cast_16_b u;
u._ushort8 = v;
-@@ -934,6 +1262,18 @@ INLINE OVERLOADABLE float4 as_float4(ushort8 v) {
+@@ -943,6 +1271,18 @@ INLINE OVERLOADABLE float4 as_float4(ushort8 v) {
return u._float4;
}
@@ -465,7 +465,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 as_int4(char16 v) {
union _type_cast_16_b u;
u._char16 = v;
-@@ -970,6 +1310,18 @@ INLINE OVERLOADABLE float4 as_float4(char16 v) {
+@@ -979,6 +1319,18 @@ INLINE OVERLOADABLE float4 as_float4(char16 v) {
return u._float4;
}
@@ -484,7 +484,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 as_int4(uchar16 v) {
union _type_cast_16_b u;
u._uchar16 = v;
-@@ -1006,6 +1358,18 @@ INLINE OVERLOADABLE float4 as_float4(uchar16 v) {
+@@ -1015,6 +1367,18 @@ INLINE OVERLOADABLE float4 as_float4(uchar16 v) {
return u._float4;
}
@@ -503,7 +503,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 as_int4(float4 v) {
union _type_cast_16_b u;
u._float4 = v;
-@@ -1042,7 +1406,26 @@ INLINE OVERLOADABLE uchar16 as_uchar16(float4 v) {
+@@ -1051,7 +1415,26 @@ INLINE OVERLOADABLE uchar16 as_uchar16(float4 v) {
return u._uchar16;
}
@@ -530,7 +530,7 @@ index 4e0459a..0d09aff 100644
int8 _int8;
uint8 _uint8;
short16 _short16;
-@@ -1050,30 +1433,126 @@ union _type_cast_32_b {
+@@ -1059,30 +1442,126 @@ union _type_cast_32_b {
float8 _float8;
};
@@ -664,7 +664,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 as_int8(uint8 v) {
union _type_cast_32_b u;
u._uint8 = v;
-@@ -1098,6 +1577,18 @@ INLINE OVERLOADABLE float8 as_float8(uint8 v) {
+@@ -1107,6 +1586,18 @@ INLINE OVERLOADABLE float8 as_float8(uint8 v) {
return u._float8;
}
@@ -683,7 +683,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 as_int8(short16 v) {
union _type_cast_32_b u;
u._short16 = v;
-@@ -1122,6 +1613,18 @@ INLINE OVERLOADABLE float8 as_float8(short16 v) {
+@@ -1131,6 +1622,18 @@ INLINE OVERLOADABLE float8 as_float8(short16 v) {
return u._float8;
}
@@ -702,7 +702,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 as_int8(ushort16 v) {
union _type_cast_32_b u;
u._ushort16 = v;
-@@ -1146,6 +1649,18 @@ INLINE OVERLOADABLE float8 as_float8(ushort16 v) {
+@@ -1155,6 +1658,18 @@ INLINE OVERLOADABLE float8 as_float8(ushort16 v) {
return u._float8;
}
@@ -721,7 +721,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 as_int8(float8 v) {
union _type_cast_32_b u;
u._float8 = v;
-@@ -1171,11 +1686,73 @@ INLINE OVERLOADABLE ushort16 as_ushort16(float8 v) {
+@@ -1180,11 +1695,73 @@ INLINE OVERLOADABLE ushort16 as_ushort16(float8 v) {
}
union _type_cast_64_b {
@@ -795,7 +795,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE uint16 as_uint16(int16 v) {
union _type_cast_64_b u;
u._int16 = v;
-@@ -1188,6 +1765,18 @@ INLINE OVERLOADABLE float16 as_float16(int16 v) {
+@@ -1197,6 +1774,18 @@ INLINE OVERLOADABLE float16 as_float16(int16 v) {
return u._float16;
}
@@ -814,7 +814,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int16 as_int16(uint16 v) {
union _type_cast_64_b u;
u._uint16 = v;
-@@ -1200,6 +1789,18 @@ INLINE OVERLOADABLE float16 as_float16(uint16 v) {
+@@ -1209,6 +1798,18 @@ INLINE OVERLOADABLE float16 as_float16(uint16 v) {
return u._float16;
}
@@ -833,7 +833,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int16 as_int16(float16 v) {
union _type_cast_64_b u;
u._float16 = v;
-@@ -1212,6 +1813,95 @@ INLINE OVERLOADABLE uint16 as_uint16(float16 v) {
+@@ -1221,9 +1822,98 @@ INLINE OVERLOADABLE uint16 as_uint16(float16 v) {
return u._uint16;
}
@@ -854,6 +854,9 @@ index 4e0459a..0d09aff 100644
+ return u._long16;
+}
+
+ // ##END_AS##
+
+ // ##BEGIN_CONVERT##
+INLINE OVERLOADABLE ulong2 convert_ulong2(long2 v) {
+ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
+}
@@ -929,7 +932,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE uint2 convert_uint2(int2 v) {
return (uint2)((uint)(v.s0), (uint)(v.s1));
}
-@@ -1236,6 +1926,14 @@ INLINE OVERLOADABLE float2 convert_float2(int2 v) {
+@@ -1248,6 +1938,14 @@ INLINE OVERLOADABLE float2 convert_float2(int2 v) {
return (float2)((float)(v.s0), (float)(v.s1));
}
@@ -944,7 +947,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 convert_int2(uint2 v) {
return (int2)((int)(v.s0), (int)(v.s1));
}
-@@ -1260,6 +1958,14 @@ INLINE OVERLOADABLE float2 convert_float2(uint2 v) {
+@@ -1272,6 +1970,14 @@ INLINE OVERLOADABLE float2 convert_float2(uint2 v) {
return (float2)((float)(v.s0), (float)(v.s1));
}
@@ -959,7 +962,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 convert_int2(short2 v) {
return (int2)((int)(v.s0), (int)(v.s1));
}
-@@ -1284,6 +1990,14 @@ INLINE OVERLOADABLE float2 convert_float2(short2 v) {
+@@ -1296,6 +2002,14 @@ INLINE OVERLOADABLE float2 convert_float2(short2 v) {
return (float2)((float)(v.s0), (float)(v.s1));
}
@@ -974,7 +977,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 convert_int2(ushort2 v) {
return (int2)((int)(v.s0), (int)(v.s1));
}
-@@ -1308,6 +2022,14 @@ INLINE OVERLOADABLE float2 convert_float2(ushort2 v) {
+@@ -1320,6 +2034,14 @@ INLINE OVERLOADABLE float2 convert_float2(ushort2 v) {
return (float2)((float)(v.s0), (float)(v.s1));
}
@@ -989,7 +992,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 convert_int2(char2 v) {
return (int2)((int)(v.s0), (int)(v.s1));
}
-@@ -1332,6 +2054,14 @@ INLINE OVERLOADABLE float2 convert_float2(char2 v) {
+@@ -1344,6 +2066,14 @@ INLINE OVERLOADABLE float2 convert_float2(char2 v) {
return (float2)((float)(v.s0), (float)(v.s1));
}
@@ -1004,7 +1007,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 convert_int2(uchar2 v) {
return (int2)((int)(v.s0), (int)(v.s1));
}
-@@ -1356,6 +2086,14 @@ INLINE OVERLOADABLE float2 convert_float2(uchar2 v) {
+@@ -1368,6 +2098,14 @@ INLINE OVERLOADABLE float2 convert_float2(uchar2 v) {
return (float2)((float)(v.s0), (float)(v.s1));
}
@@ -1019,7 +1022,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int2 convert_int2(float2 v) {
return (int2)((int)(v.s0), (int)(v.s1));
}
-@@ -1364,20 +2102,92 @@ INLINE OVERLOADABLE uint2 convert_uint2(float2 v) {
+@@ -1376,20 +2114,92 @@ INLINE OVERLOADABLE uint2 convert_uint2(float2 v) {
return (uint2)((uint)(v.s0), (uint)(v.s1));
}
@@ -1120,7 +1123,7 @@ index 4e0459a..0d09aff 100644
}
INLINE OVERLOADABLE uint3 convert_uint3(int3 v) {
-@@ -1404,6 +2214,14 @@ INLINE OVERLOADABLE float3 convert_float3(int3 v) {
+@@ -1416,6 +2226,14 @@ INLINE OVERLOADABLE float3 convert_float3(int3 v) {
return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
}
@@ -1135,7 +1138,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int3 convert_int3(uint3 v) {
return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
}
-@@ -1428,6 +2246,14 @@ INLINE OVERLOADABLE float3 convert_float3(uint3 v) {
+@@ -1440,6 +2258,14 @@ INLINE OVERLOADABLE float3 convert_float3(uint3 v) {
return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
}
@@ -1150,7 +1153,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int3 convert_int3(short3 v) {
return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
}
-@@ -1452,6 +2278,14 @@ INLINE OVERLOADABLE float3 convert_float3(short3 v) {
+@@ -1464,6 +2290,14 @@ INLINE OVERLOADABLE float3 convert_float3(short3 v) {
return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
}
@@ -1165,7 +1168,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int3 convert_int3(ushort3 v) {
return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
}
-@@ -1476,6 +2310,14 @@ INLINE OVERLOADABLE float3 convert_float3(ushort3 v) {
+@@ -1488,6 +2322,14 @@ INLINE OVERLOADABLE float3 convert_float3(ushort3 v) {
return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
}
@@ -1180,7 +1183,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int3 convert_int3(char3 v) {
return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
}
-@@ -1500,6 +2342,14 @@ INLINE OVERLOADABLE float3 convert_float3(char3 v) {
+@@ -1512,6 +2354,14 @@ INLINE OVERLOADABLE float3 convert_float3(char3 v) {
return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
}
@@ -1195,7 +1198,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int3 convert_int3(uchar3 v) {
return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
}
-@@ -1524,6 +2374,14 @@ INLINE OVERLOADABLE float3 convert_float3(uchar3 v) {
+@@ -1536,6 +2386,14 @@ INLINE OVERLOADABLE float3 convert_float3(uchar3 v) {
return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
}
@@ -1210,7 +1213,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int3 convert_int3(float3 v) {
return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
}
-@@ -1548,6 +2406,78 @@ INLINE OVERLOADABLE uchar3 convert_uchar3(float3 v) {
+@@ -1560,6 +2418,78 @@ INLINE OVERLOADABLE uchar3 convert_uchar3(float3 v) {
return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
}
@@ -1289,7 +1292,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE uint4 convert_uint4(int4 v) {
return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
}
-@@ -1572,6 +2502,14 @@ INLINE OVERLOADABLE float4 convert_float4(int4 v) {
+@@ -1584,6 +2514,14 @@ INLINE OVERLOADABLE float4 convert_float4(int4 v) {
return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
}
@@ -1304,7 +1307,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 convert_int4(uint4 v) {
return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
}
-@@ -1596,6 +2534,14 @@ INLINE OVERLOADABLE float4 convert_float4(uint4 v) {
+@@ -1608,6 +2546,14 @@ INLINE OVERLOADABLE float4 convert_float4(uint4 v) {
return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
}
@@ -1319,7 +1322,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 convert_int4(short4 v) {
return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
}
-@@ -1620,6 +2566,14 @@ INLINE OVERLOADABLE float4 convert_float4(short4 v) {
+@@ -1632,6 +2578,14 @@ INLINE OVERLOADABLE float4 convert_float4(short4 v) {
return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
}
@@ -1334,7 +1337,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 convert_int4(ushort4 v) {
return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
}
-@@ -1644,6 +2598,14 @@ INLINE OVERLOADABLE float4 convert_float4(ushort4 v) {
+@@ -1656,6 +2610,14 @@ INLINE OVERLOADABLE float4 convert_float4(ushort4 v) {
return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
}
@@ -1349,7 +1352,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 convert_int4(char4 v) {
return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
}
-@@ -1668,6 +2630,14 @@ INLINE OVERLOADABLE float4 convert_float4(char4 v) {
+@@ -1680,6 +2642,14 @@ INLINE OVERLOADABLE float4 convert_float4(char4 v) {
return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
}
@@ -1364,7 +1367,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 convert_int4(uchar4 v) {
return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
}
-@@ -1692,6 +2662,14 @@ INLINE OVERLOADABLE float4 convert_float4(uchar4 v) {
+@@ -1704,6 +2674,14 @@ INLINE OVERLOADABLE float4 convert_float4(uchar4 v) {
return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
}
@@ -1379,7 +1382,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int4 convert_int4(float4 v) {
return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
}
-@@ -1716,6 +2694,78 @@ INLINE OVERLOADABLE uchar4 convert_uchar4(float4 v) {
+@@ -1728,6 +2706,78 @@ INLINE OVERLOADABLE uchar4 convert_uchar4(float4 v) {
return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
}
@@ -1458,7 +1461,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE uint8 convert_uint8(int8 v) {
return (uint8)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3), (uint)(v.s4), (uint)(v.s5), (uint)(v.s6), (uint)(v.s7));
}
-@@ -1740,6 +2790,14 @@ INLINE OVERLOADABLE float8 convert_float8(int8 v) {
+@@ -1752,6 +2802,14 @@ INLINE OVERLOADABLE float8 convert_float8(int8 v) {
return (float8)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7));
}
@@ -1473,7 +1476,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 convert_int8(uint8 v) {
return (int8)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7));
}
-@@ -1764,6 +2822,14 @@ INLINE OVERLOADABLE float8 convert_float8(uint8 v) {
+@@ -1776,6 +2834,14 @@ INLINE OVERLOADABLE float8 convert_float8(uint8 v) {
return (float8)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7));
}
@@ -1488,7 +1491,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 convert_int8(short8 v) {
return (int8)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7));
}
-@@ -1788,6 +2854,14 @@ INLINE OVERLOADABLE float8 convert_float8(short8 v) {
+@@ -1800,6 +2866,14 @@ INLINE OVERLOADABLE float8 convert_float8(short8 v) {
return (float8)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7));
}
@@ -1503,7 +1506,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 convert_int8(ushort8 v) {
return (int8)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7));
}
-@@ -1812,6 +2886,14 @@ INLINE OVERLOADABLE float8 convert_float8(ushort8 v) {
+@@ -1824,6 +2898,14 @@ INLINE OVERLOADABLE float8 convert_float8(ushort8 v) {
return (float8)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7));
}
@@ -1518,7 +1521,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 convert_int8(char8 v) {
return (int8)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7));
}
-@@ -1836,6 +2918,14 @@ INLINE OVERLOADABLE float8 convert_float8(char8 v) {
+@@ -1848,6 +2930,14 @@ INLINE OVERLOADABLE float8 convert_float8(char8 v) {
return (float8)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7));
}
@@ -1533,7 +1536,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 convert_int8(uchar8 v) {
return (int8)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7));
}
-@@ -1860,6 +2950,14 @@ INLINE OVERLOADABLE float8 convert_float8(uchar8 v) {
+@@ -1872,6 +2962,14 @@ INLINE OVERLOADABLE float8 convert_float8(uchar8 v) {
return (float8)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7));
}
@@ -1548,7 +1551,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int8 convert_int8(float8 v) {
return (int8)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7));
}
-@@ -1884,6 +2982,78 @@ INLINE OVERLOADABLE uchar8 convert_uchar8(float8 v) {
+@@ -1896,6 +2994,78 @@ INLINE OVERLOADABLE uchar8 convert_uchar8(float8 v) {
return (uchar8)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3), (uchar)(v.s4), (uchar)(v.s5), (uchar)(v.s6), (uchar)(v.s7));
}
@@ -1627,7 +1630,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE uint16 convert_uint16(int16 v) {
return (uint16)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3), (uint)(v.s4), (uint)(v.s5), (uint)(v.s6), (uint)(v.s7), (uint)(v.s8), (uint)(v.s9), (uint)(v.sA), (uint)(v.sB), (uint)(v.sC), (uint)(v.sD), (uint)(v.sE), (uint)(v.sF));
}
-@@ -1908,6 +3078,14 @@ INLINE OVERLOADABLE float16 convert_float16(int16 v) {
+@@ -1920,6 +3090,14 @@ INLINE OVERLOADABLE float16 convert_float16(int16 v) {
return (float16)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7), (float)(v.s8), (float)(v.s9), (float)(v.sA), (float)(v.sB), (float)(v.sC), (float)(v.sD), (float)(v.sE), (float)(v.sF));
}
@@ -1642,7 +1645,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int16 convert_int16(uint16 v) {
return (int16)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7), (int)(v.s8), (int)(v.s9), (int)(v.sA), (int)(v.sB), (int)(v.sC), (int)(v.sD), (int)(v.sE), (int)(v.sF));
}
-@@ -1932,6 +3110,14 @@ INLINE OVERLOADABLE float16 convert_float16(uint16 v) {
+@@ -1944,6 +3122,14 @@ INLINE OVERLOADABLE float16 convert_float16(uint16 v) {
return (float16)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7), (float)(v.s8), (float)(v.s9), (float)(v.sA), (float)(v.sB), (float)(v.sC), (float)(v.sD), (float)(v.sE), (float)(v.sF));
}
@@ -1657,7 +1660,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int16 convert_int16(short16 v) {
return (int16)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7), (int)(v.s8), (int)(v.s9), (int)(v.sA), (int)(v.sB), (int)(v.sC), (int)(v.sD), (int)(v.sE), (int)(v.sF));
}
-@@ -1956,6 +3142,14 @@ INLINE OVERLOADABLE float16 convert_float16(short16 v) {
+@@ -1968,6 +3154,14 @@ INLINE OVERLOADABLE float16 convert_float16(short16 v) {
return (float16)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7), (float)(v.s8), (float)(v.s9), (float)(v.sA), (float)(v.sB), (float)(v.sC), (float)(v.sD), (float)(v.sE), (float)(v.sF));
}
@@ -1672,7 +1675,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int16 convert_int16(ushort16 v) {
return (int16)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7), (int)(v.s8), (int)(v.s9), (int)(v.sA), (int)(v.sB), (int)(v.sC), (int)(v.sD), (int)(v.sE), (int)(v.sF));
}
-@@ -1980,6 +3174,14 @@ INLINE OVERLOADABLE float16 convert_float16(ushort16 v) {
+@@ -1992,6 +3186,14 @@ INLINE OVERLOADABLE float16 convert_float16(ushort16 v) {
return (float16)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7), (float)(v.s8), (float)(v.s9), (float)(v.sA), (float)(v.sB), (float)(v.sC), (float)(v.sD), (float)(v.sE), (float)(v.sF));
}
@@ -1687,7 +1690,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int16 convert_int16(char16 v) {
return (int16)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7), (int)(v.s8), (int)(v.s9), (int)(v.sA), (int)(v.sB), (int)(v.sC), (int)(v.sD), (int)(v.sE), (int)(v.sF));
}
-@@ -2004,6 +3206,14 @@ INLINE OVERLOADABLE float16 convert_float16(char16 v) {
+@@ -2016,6 +3218,14 @@ INLINE OVERLOADABLE float16 convert_float16(char16 v) {
return (float16)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7), (float)(v.s8), (float)(v.s9), (float)(v.sA), (float)(v.sB), (float)(v.sC), (float)(v.sD), (float)(v.sE), (float)(v.sF));
}
@@ -1702,7 +1705,7 @@ index 4e0459a..0d09aff 100644
INLINE OVERLOADABLE int16 convert_int16(uchar16 v) {
return (int16)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3), (int)(v.s4), (int)(v.s5), (int)(v.s6), (int)(v.s7), (int)(v.s8), (int)(v.s9), (int)(v.sA), (int)(v.sB), (int)(v.sC), (int)(v.sD), (int)(v.sE), (int)(v.sF));
}
-@@ -2028,6 +3238,14 @@ INLINE OVERLOADABLE float16 convert_float16(uchar16 v) {
+@@ -2040,6 +3250,14 @@ INLINE OVERLOADABLE float16 convert_float16(uchar16 v) {
return (float16)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3), (float)(v.s4), (float)(v.s5), (float)(v.s6), (float)(v.s7), (float)(v.s8), (float)(v.s9), (float)(v.sA), (float)(v.sB), (float)(v.sC), (float)(v.sD), (float)(v.sE), (float)(v.sF));
}
diff --git a/debian/patches/0004-Add-vector-argument-test-case.patch b/debian/patches/0004-Add-vector-argument-test-case.patch
index 4a1fbbe..cf43d98 100644
--- a/debian/patches/0004-Add-vector-argument-test-case.patch
+++ b/debian/patches/0004-Add-vector-argument-test-case.patch
@@ -1,7 +1,7 @@
-From 123085e021ecb265228500c9deea407139762ba8 Mon Sep 17 00:00:00 2001
+From e9f476243902f2f3989c880030b267c384d7c040 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 4/7] Add vector argument test case.
+Subject: [PATCH 4/4] Add vector argument test case.
To: beignet at lists.freedesktop.org
Signed-off-by: Yang Rong <rong.r.yang at intel.com>
@@ -26,10 +26,10 @@ index 0000000..0985dbd
+ dst[id] = value.w;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
-index 2ba01c4..717383b 100644
+index e5c03ee..f0bbe46 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
-@@ -32,6 +32,7 @@ set (utests_sources
+@@ -34,6 +34,7 @@ set (utests_sources
compiler_fill_image_3d_2.cpp
compiler_function_argument0.cpp
compiler_function_argument1.cpp
diff --git a/debian/patches/0005-Add-more-get-image-info-functions.patch b/debian/patches/0005-Add-more-get-image-info-functions.patch
deleted file mode 100644
index 4682e51..0000000
--- a/debian/patches/0005-Add-more-get-image-info-functions.patch
+++ /dev/null
@@ -1,203 +0,0 @@
-From 503bee58b3080f4b863e61be22065c2fd2923827 Mon Sep 17 00:00:00 2001
-From: Zhigang Gong <zhigang.gong at linux.intel.com>
-Date: Mon, 20 May 2013 17:00:37 +0800
-Subject: [PATCH 5/7] Add more get image info functions.
-To: beignet at lists.freedesktop.org
-
-Add get image depth/channel data type/channel order/dim support.
-Now, only those functions for the unsupported image type have not
-been implemented. The unsupported image types are as below:
-image1d_t,image1d_buffer_t,image1d_array_t,image2d_array_t.
-
-Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
----
- backend/src/ir/image.cpp | 14 ++++++++---
- backend/src/ir/instruction.hpp | 10 ++++++--
- backend/src/llvm/llvm_gen_backend.cpp | 6 +++++
- backend/src/llvm/llvm_gen_ocl_function.hxx | 3 +++
- backend/src/ocl_stdlib.h | 37 ++++++++++++++++++----------
- src/cl_command_queue.c | 6 +++++
- 6 files changed, 57 insertions(+), 19 deletions(-)
-
-diff --git a/backend/src/ir/image.cpp b/backend/src/ir/image.cpp
-index 9398e22..486fde1 100644
---- a/backend/src/ir/image.cpp
-+++ b/backend/src/ir/image.cpp
-@@ -31,8 +31,11 @@ namespace ir {
- static uint32_t getInfoOffset4Type(struct ImageInfo *imageInfo, int type)
- {
- switch (type) {
-- case GetImageInfoInstruction::WIDTH: return imageInfo->wSlot;
-- case GetImageInfoInstruction::HEIGHT: return imageInfo->hSlot;
-+ case GetImageInfoInstruction::WIDTH: return imageInfo->wSlot;
-+ case GetImageInfoInstruction::HEIGHT: return imageInfo->hSlot;
-+ case GetImageInfoInstruction::DEPTH: return imageInfo->depthSlot;
-+ case GetImageInfoInstruction::CHANNEL_DATA_TYPE: return imageInfo->dataTypeSlot;
-+ case GetImageInfoInstruction::CHANNEL_ORDER: return imageInfo->channelOrderSlot;
- default:
- NOT_IMPLEMENTED;
- }
-@@ -42,8 +45,11 @@ namespace ir {
- static uint32_t setInfoOffset4Type(struct ImageInfo *imageInfo, int type, uint32_t offset)
- {
- switch (type) {
-- case GetImageInfoInstruction::WIDTH: imageInfo->wSlot = offset; break;
-- case GetImageInfoInstruction::HEIGHT: imageInfo->hSlot = offset; break;
-+ case GetImageInfoInstruction::WIDTH: imageInfo->wSlot = offset; break;
-+ case GetImageInfoInstruction::HEIGHT: imageInfo->hSlot = offset; break;
-+ case GetImageInfoInstruction::DEPTH: imageInfo->depthSlot = offset; break;
-+ case GetImageInfoInstruction::CHANNEL_DATA_TYPE: imageInfo->dataTypeSlot = offset; break;
-+ case GetImageInfoInstruction::CHANNEL_ORDER: imageInfo->channelOrderSlot = offset; break;
- default:
- NOT_IMPLEMENTED;
- }
-diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp
-index c948d2c..7662b6a 100644
---- a/backend/src/ir/instruction.hpp
-+++ b/backend/src/ir/instruction.hpp
-@@ -335,12 +335,18 @@ namespace ir {
- enum {
- WIDTH = 0,
- HEIGHT = 1,
-+ DEPTH = 2,
-+ CHANNEL_DATA_TYPE = 3,
-+ CHANNEL_ORDER = 4,
- };
-
- static INLINE uint32_t getDstNum4Type(int infoType) {
- switch (infoType) {
-- case GetImageInfoInstruction::WIDTH:
-- case GetImageInfoInstruction::HEIGHT:
-+ case WIDTH:
-+ case HEIGHT:
-+ case DEPTH:
-+ case CHANNEL_DATA_TYPE:
-+ case CHANNEL_ORDER:
- return 1;
- break;
- default:
-diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
-index deda687..a0946ef 100644
---- a/backend/src/llvm/llvm_gen_backend.cpp
-+++ b/backend/src/llvm/llvm_gen_backend.cpp
-@@ -1662,6 +1662,9 @@ namespace gbe
- case GEN_OCL_RNDD:
- case GEN_OCL_GET_IMAGE_WIDTH:
- case GEN_OCL_GET_IMAGE_HEIGHT:
-+ case GEN_OCL_GET_IMAGE_CHANNEL_DATA_TYPE:
-+ case GEN_OCL_GET_IMAGE_CHANNEL_ORDER:
-+ case GEN_OCL_GET_IMAGE_DEPTH:
- // No structure can be returned
- this->newRegister(&I);
- break;
-@@ -1827,6 +1830,9 @@ namespace gbe
- case GEN_OCL_LGBARRIER: ctx.SYNC(ir::syncLocalBarrier | ir::syncGlobalBarrier); break;
- case GEN_OCL_GET_IMAGE_WIDTH:
- case GEN_OCL_GET_IMAGE_HEIGHT:
-+ case GEN_OCL_GET_IMAGE_DEPTH:
-+ case GEN_OCL_GET_IMAGE_CHANNEL_DATA_TYPE:
-+ case GEN_OCL_GET_IMAGE_CHANNEL_ORDER:
- {
- GBE_ASSERT(AI != AE); const ir::Register surface_id = this->getRegister(*AI); ++AI;
- uint32_t elemNum;
-diff --git a/backend/src/llvm/llvm_gen_ocl_function.hxx b/backend/src/llvm/llvm_gen_ocl_function.hxx
-index 2fb33c0..0524744 100644
---- a/backend/src/llvm/llvm_gen_ocl_function.hxx
-+++ b/backend/src/llvm/llvm_gen_ocl_function.hxx
-@@ -73,6 +73,9 @@ DECL_LLVM_GEN_FUNCTION(WRITE_IMAGE15, _Z22__gen_ocl_write_imagefjfffDv4_f)
- // To get image info function
- DECL_LLVM_GEN_FUNCTION(GET_IMAGE_WIDTH, __gen_ocl_get_image_width)
- DECL_LLVM_GEN_FUNCTION(GET_IMAGE_HEIGHT, __gen_ocl_get_image_height)
-+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)
-
- // saturation related functions.
- DECL_LLVM_GEN_FUNCTION(SADD_SAT_CHAR, _Z12ocl_sadd_satcc)
-diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
-index 0d09aff..78a8da3 100644
---- a/backend/src/ocl_stdlib.h
-+++ b/backend/src/ocl_stdlib.h
-@@ -4257,7 +4257,9 @@ OVERLOADABLE void __gen_ocl_write_imagef(uint surface_id, int u, int v, int w, f
- OVERLOADABLE void __gen_ocl_write_imagef(uint surface_id, float u, float v, float w, float4 color);
- int __gen_ocl_get_image_width(uint surface_id);
- int __gen_ocl_get_image_height(uint surface_id);
--//OVERLOADABLE int __gen_ocl_get_image_depth(image3d_t image);
-+int __gen_ocl_get_image_channel_data_type(uint surface_id);
-+int __gen_ocl_get_image_channel_order(uint surface_id);
-+int __gen_ocl_get_image_depth(uint surface_id);
-
- #define GET_IMAGE(cl_image, surface_id) \
- uint surface_id = (uint)cl_image
-@@ -4309,17 +4311,32 @@ DECL_IMAGE(float4, f)
- { \
- GET_IMAGE(image, surface_id);\
- return __gen_ocl_get_image_height(surface_id); \
-- }
--#if 0
-+ } \
- INLINE_OVERLOADABLE int get_image_channel_data_type(image_type image)\
-- { NOT_IMPLEMENTED; }\
-+ { \
-+ GET_IMAGE(image, surface_id);\
-+ return __gen_ocl_get_image_channel_data_type(surface_id); \
-+ }\
- INLINE_OVERLOADABLE int get_image_channel_order(image_type image)\
-- { NOT_IMPLEMENTED; }
--#endif
--
-+ { \
-+ GET_IMAGE(image, surface_id);\
-+ return __gen_ocl_get_image_channel_order(surface_id); \
-+ }
-
- DECL_IMAGE_INFO(image2d_t)
- DECL_IMAGE_INFO(image3d_t)
-+
-+INLINE_OVERLOADABLE int get_image_depth(image3d_t image)
-+ {
-+ GET_IMAGE(image, surface_id);
-+ return __gen_ocl_get_image_depth(surface_id);
-+ }
-+
-+INLINE_OVERLOADABLE int2 get_image_dim(image2d_t image)
-+ { return (int2){get_image_width(image), get_image_height(image)}; }
-+
-+INLINE_OVERLOADABLE int4 get_image_dim(image3d_t image)
-+ { return (int4){get_image_width(image), get_image_height(image), get_image_depth(image), 0}; }
- #if 0
- /* The following functions are not implemented yet. */
- DECL_IMAGE_INFO(image1d_t)
-@@ -4327,12 +4344,6 @@ DECL_IMAGE_INFO(image1d_buffer_t)
- DECL_IMAGE_INFO(image1d_array_t)
- DECL_IMAGE_INFO(image2d_array_t)
-
--INLINE_OVERLOADABLE int get_image_depth(image3d_t image)
-- { return __gen_ocl_get_image_depth(image); }
--
--INLINE_OVERLOADABLE int2 get_image_dim(image2d_t image)
-- { return __gen_ocl_get_image_dim(image); }
--
- INLINE_OVERLOADABLE int2 get_image_dim(image2d_array_t image)
- { return __gen_ocl_get_image_dim(image); }
-
-diff --git a/src/cl_command_queue.c b/src/cl_command_queue.c
-index 6ce9016..1a37c78 100644
---- a/src/cl_command_queue.c
-+++ b/src/cl_command_queue.c
-@@ -105,6 +105,12 @@ set_image_info(char *curbe, struct ImageInfo * image_info, cl_mem image)
- *(uint32_t*)(curbe + image_info->wSlot) = image->w;
- if (image_info->hSlot >= 0)
- *(uint32_t*)(curbe + image_info->hSlot) = image->h;
-+ if (image_info->depthSlot >= 0)
-+ *(uint32_t*)(curbe + image_info->depthSlot) = image->depth;
-+ if (image_info->channelOrderSlot >= 0)
-+ *(uint32_t*)(curbe + image_info->channelOrderSlot) = image->fmt.image_channel_order;
-+ if (image_info->dataTypeSlot >= 0)
-+ *(uint32_t*)(curbe + image_info->dataTypeSlot) = image->fmt.image_channel_data_type;
- }
-
- LOCAL cl_int
---
-1.7.10.4
-
diff --git a/debian/patches/0006-utests-extent-get_image_size-cases-to-other-informat.patch b/debian/patches/0006-utests-extent-get_image_size-cases-to-other-informat.patch
deleted file mode 100644
index fe06e44..0000000
--- a/debian/patches/0006-utests-extent-get_image_size-cases-to-other-informat.patch
+++ /dev/null
@@ -1,180 +0,0 @@
-From 94021c43047de2593393777c39189b45cb8043ad Mon Sep 17 00:00:00 2001
-From: Zhigang Gong <zhigang.gong at linux.intel.com>
-Date: Tue, 21 May 2013 16:14:52 +0800
-Subject: [PATCH 6/7] utests: extent get_image_size cases to other
- informations..
-To: beignet at lists.freedesktop.org
-
-Please ignore the previous patch, I sent the wrong patch. Please try this one.
-Sorry for that.
-
-From 6107ba1fee7091dee4ead72fbd780f94803be42f Mon Sep 17 00:00:00 2001
-From: Zhigang Gong <zhigang.gong at linux.intel.com>
-Date: Mon, 20 May 2013 16:41:28 +0800
-Subject: [PATCH v2 2/2] utests: extent get_image_size cases to other
- informations..
-
-Extent it to test all the supported image informations.
-
-Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
----
- kernels/test_get_image_info.cl | 13 ++++++++++
- kernels/test_get_image_size.cl | 9 -------
- utests/CMakeLists.txt | 2 +-
- utests/compiler_get_image_info.cpp | 50 ++++++++++++++++++++++++++++++++++++
- utests/compiler_get_image_size.cpp | 37 --------------------------
- 5 files changed, 64 insertions(+), 47 deletions(-)
- create mode 100644 kernels/test_get_image_info.cl
- delete mode 100644 kernels/test_get_image_size.cl
- create mode 100644 utests/compiler_get_image_info.cpp
- delete mode 100644 utests/compiler_get_image_size.cpp
-
-diff --git a/kernels/test_get_image_info.cl b/kernels/test_get_image_info.cl
-new file mode 100644
-index 0000000..8f69b75
---- /dev/null
-+++ b/kernels/test_get_image_info.cl
-@@ -0,0 +1,13 @@
-+__kernel void
-+test_get_image_info(__write_only image3d_t src, __global int *size, __global int *fmt)
-+{
-+ int id = (int)get_global_id(0);
-+ int w, h, depth;
-+ w = get_image_width(src);
-+ h = get_image_height(src);
-+ depth = get_image_depth(src);
-+ int channel_data_type = get_image_channel_data_type(src);
-+ int channel_order = get_image_channel_order(src);
-+ size[id] = (w << 20 | h << 8 | depth);
-+ fmt[id] = (channel_data_type << 16 | channel_order);
-+}
-diff --git a/kernels/test_get_image_size.cl b/kernels/test_get_image_size.cl
-deleted file mode 100644
-index aeb7d66..0000000
---- a/kernels/test_get_image_size.cl
-+++ /dev/null
-@@ -1,9 +0,0 @@
--__kernel void
--test_get_image_size(__write_only image2d_t src, __global int *info)
--{
-- int id = (int)get_global_id(0);
-- int w, h;
-- w = get_image_width(src);
-- h = get_image_height(src);
-- info[id] = (w << 16 | h);
--}
-diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
-index 717383b..8498d73 100644
---- a/utests/CMakeLists.txt
-+++ b/utests/CMakeLists.txt
-@@ -76,7 +76,7 @@ set (utests_sources
- compiler_movforphi_undef.cpp
- compiler_volatile.cpp
- compiler_copy_image1.cpp
-- compiler_get_image_size.cpp
-+ compiler_get_image_info.cpp
- runtime_createcontext.cpp
- utest_assert.cpp
- utest.cpp
-diff --git a/utests/compiler_get_image_info.cpp b/utests/compiler_get_image_info.cpp
-new file mode 100644
-index 0000000..3b9d132
---- /dev/null
-+++ b/utests/compiler_get_image_info.cpp
-@@ -0,0 +1,50 @@
-+#include "utest_helper.hpp"
-+
-+static void compiler_get_image_info(void)
-+{
-+ const size_t w = 256;
-+ const size_t h = 512;
-+ const size_t depth = 3;
-+ cl_image_format format;
-+ cl_image_desc desc;
-+
-+ format.image_channel_order = CL_RGBA;
-+ format.image_channel_data_type = CL_UNSIGNED_INT8;
-+ desc.image_type = CL_MEM_OBJECT_IMAGE3D;
-+ desc.image_width = w;
-+ desc.image_height = h;
-+ desc.image_depth = depth;
-+ desc.image_row_pitch = 0;
-+ desc.image_slice_pitch = 0;
-+ desc.num_mip_levels = 0;
-+ desc.num_samples = 0;
-+ desc.buffer = NULL;
-+
-+ // Setup kernel and images
-+ OCL_CREATE_KERNEL("test_get_image_info");
-+
-+ OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
-+ OCL_CREATE_BUFFER(buf[1], 0, 32 * sizeof(int), NULL);
-+ OCL_CREATE_BUFFER(buf[2], 0, 32 * sizeof(int), 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] = 32;
-+ locals[0] = 16;
-+ OCL_NDRANGE(1);
-+
-+ // Check result
-+ OCL_MAP_BUFFER(1);
-+ OCL_MAP_BUFFER(2);
-+ for (uint32_t i = 0; i < 32; i++)
-+ {
-+ OCL_ASSERT(((uint32_t*)buf_data[1])[i] == ((w << 20) | (h << 8) | depth));
-+ OCL_ASSERT(((uint32_t*)buf_data[2])[i] == ((CL_UNSIGNED_INT8 << 16) | CL_RGBA));
-+ }
-+ OCL_UNMAP_BUFFER(1);
-+ OCL_UNMAP_BUFFER(2);
-+}
-+
-+MAKE_UTEST_FROM_FUNCTION(compiler_get_image_info);
-diff --git a/utests/compiler_get_image_size.cpp b/utests/compiler_get_image_size.cpp
-deleted file mode 100644
-index 49c08ad..0000000
---- a/utests/compiler_get_image_size.cpp
-+++ /dev/null
-@@ -1,37 +0,0 @@
--#include "utest_helper.hpp"
--
--static void compiler_get_image_size(void)
--{
-- const size_t w = 256;
-- const size_t h = 512;
-- cl_image_format format;
-- cl_image_desc desc;
--
-- format.image_channel_order = CL_RGBA;
-- format.image_channel_data_type = CL_UNSIGNED_INT8;
-- desc.image_type = CL_MEM_OBJECT_IMAGE2D;
-- desc.image_width = w;
-- desc.image_height = h;
-- desc.image_row_pitch = 0;
--
-- // Setup kernel and images
-- OCL_CREATE_KERNEL("test_get_image_size");
--
-- OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
-- OCL_CREATE_BUFFER(buf[1], 0, 32 * sizeof(int), NULL);
--
-- // Run the kernel
-- OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
-- OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
-- globals[0] = 32;
-- locals[0] = 16;
-- OCL_NDRANGE(1);
--
-- // Check result
-- OCL_MAP_BUFFER(1);
-- for (uint32_t i = 0; i < 32; i++)
-- OCL_ASSERT(((uint32_t*)buf_data[1])[i] == ((w << 16) | (h)));
-- OCL_UNMAP_BUFFER(0);
--}
--
--MAKE_UTEST_FROM_FUNCTION(compiler_get_image_size);
---
-1.7.10.4
-
diff --git a/debian/patches/0007-Change-clang-system-call-to-libclang-api-call.patch b/debian/patches/0007-Change-clang-system-call-to-libclang-api-call.patch
deleted file mode 100644
index bdb289f..0000000
--- a/debian/patches/0007-Change-clang-system-call-to-libclang-api-call.patch
+++ /dev/null
@@ -1,261 +0,0 @@
-From 2d98eff3e09fe2b63812401949af8aea81308c8c Mon Sep 17 00:00:00 2001
-From: Yang Rong <rong.r.yang at intel.com>
-Date: Tue, 21 May 2013 12:45:56 +0800
-Subject: [PATCH 7/7] Change clang system call to libclang api call.
-To: beignet at lists.freedesktop.org
-
-The original call clang command directly as frontend. The implement is not very flexible.
-I change to call libclang apis, now support both clang 3.1 and clang 3.2.
-Now still write the intermediate to the file, for code simply.
-Also fix llvm 3.1 build errors for my vector scalarize commit.
-
-Signed-off-by: Yang Rong <rong.r.yang at intel.com>
----
- CMake/FindLLVM.cmake | 27 ++++++++-
- backend/src/CMakeLists.txt | 1 +
- backend/src/backend/program.cpp | 103 ++++++++++++++++++++++++++-------
- backend/src/llvm/llvm_gen_backend.cpp | 7 +--
- backend/src/llvm/llvm_scalarize.cpp | 10 ++--
- 5 files changed, 116 insertions(+), 32 deletions(-)
-
-diff --git a/CMake/FindLLVM.cmake b/CMake/FindLLVM.cmake
-index c06b8a4..b320639 100644
---- a/CMake/FindLLVM.cmake
-+++ b/CMake/FindLLVM.cmake
-@@ -8,7 +8,7 @@
- # LLVM_FOUND - True if llvm found.
- if (LLVM_INSTALL_DIR)
- 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)
-+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)
-
-@@ -66,3 +66,28 @@ execute_process(
- OUTPUT_VARIABLE LLVM_MODULE_LIBS
- OUTPUT_STRIP_TRAILING_WHITESPACE
- )
-+
-+macro(add_one_lib name)
-+ FIND_LIBRARY(CLANG_LIB
-+ NAMES ${name}
-+ PATHS ${LLVM_LIBRARY_DIR} )
-+ set(CLANG_LIBRARIES ${CLANG_LIBRARIES} ${CLANG_LIB})
-+ unset(CLANG_LIB CACHE)
-+endmacro()
-+
-+#Assume clang lib path same as llvm lib path
-+add_one_lib("clangFrontend")
-+add_one_lib("clangSerialization")
-+add_one_lib("clangDriver")
-+add_one_lib("clangCodeGen")
-+add_one_lib("clangSema")
-+add_one_lib("clangStaticAnalyzerFrontend")
-+add_one_lib("clangStaticAnalyzerCheckers")
-+add_one_lib("clangStaticAnalyzerCore")
-+add_one_lib("clangAnalysis")
-+add_one_lib("clangEdit")
-+add_one_lib("clangAST")
-+add_one_lib("clangParse")
-+add_one_lib("clangSema")
-+add_one_lib("clangLex")
-+add_one_lib("clangBasic")
-diff --git a/backend/src/CMakeLists.txt b/backend/src/CMakeLists.txt
-index 183517a..a0fe198 100644
---- a/backend/src/CMakeLists.txt
-+++ b/backend/src/CMakeLists.txt
-@@ -116,6 +116,7 @@ target_link_libraries(
- ${DRM_INTEL_LIBRARY}
- ${DRM_LIBRARY}
- ${OPENGL_LIBRARIES}
-+ ${CLANG_LIBRARIES}
- ${LLVM_MODULE_LIBS}
- ${CMAKE_THREAD_LIBS_INIT}
- ${CMAKE_DL_LIBS})
-diff --git a/backend/src/backend/program.cpp b/backend/src/backend/program.cpp
-index c46c681..6816a13 100644
---- a/backend/src/backend/program.cpp
-+++ b/backend/src/backend/program.cpp
-@@ -46,6 +46,23 @@
- #define LLVM_VERSION_MINOR 0
- #endif /* !defined(LLVM_VERSION_MINOR) */
-
-+#include <clang/CodeGen/CodeGenAction.h>
-+#include <clang/Frontend/CompilerInstance.h>
-+#include <clang/Frontend/CompilerInvocation.h>
-+#if LLVM_VERSION_MINOR <= 1
-+#include <clang/Frontend/DiagnosticOptions.h>
-+#else
-+#include <clang/Basic/DiagnosticOptions.h>
-+#endif /* LLVM_VERSION_MINOR <= 1 */
-+#include <clang/Frontend/TextDiagnosticPrinter.h>
-+#include <clang/Basic/TargetInfo.h>
-+#include <clang/Basic/TargetOptions.h>
-+#include <llvm/ADT/IntrusiveRefCntPtr.h>
-+#include <llvm/ADT/OwningPtr.h>
-+#include <llvm/Module.h>
-+#include <llvm/Bitcode/ReaderWriter.h>
-+#include <llvm/Support/raw_ostream.h>
-+
- namespace gbe {
-
- Kernel::Kernel(const std::string &name) :
-@@ -104,6 +121,71 @@ namespace gbe {
- GBE_SAFE_DELETE(program);
- }
-
-+ static void buildModuleFromSource(const char* input, const char* output) {
-+ // Arguments to pass to the clang frontend
-+ vector<const char *> args;
-+ args.push_back("-emit-llvm");
-+ args.push_back("-O3");
-+ args.push_back("-triple");
-+ args.push_back("nvptx");
-+ args.push_back(input);
-+
-+ // The compiler invocation needs a DiagnosticsEngine so it can report problems
-+#if LLVM_VERSION_MINOR <= 1
-+ args.push_back("-triple");
-+ args.push_back("ptx32");
-+
-+ clang::TextDiagnosticPrinter *DiagClient =
-+ new clang::TextDiagnosticPrinter(llvm::errs(), clang::DiagnosticOptions());
-+ llvm::IntrusiveRefCntPtr<clang::DiagnosticIDs> DiagID(new clang::DiagnosticIDs());
-+ clang::DiagnosticsEngine Diags(DiagID, DiagClient);
-+#else
-+ args.push_back("-ffp-contract=off");
-+ args.push_back("-triple");
-+ args.push_back("nvptx");
-+
-+ llvm::IntrusiveRefCntPtr<clang::DiagnosticOptions> DiagOpts = new clang::DiagnosticOptions();
-+ clang::TextDiagnosticPrinter *DiagClient =
-+ new clang::TextDiagnosticPrinter(llvm::errs(), &*DiagOpts);
-+ llvm::IntrusiveRefCntPtr<clang::DiagnosticIDs> DiagID(new clang::DiagnosticIDs());
-+ clang::DiagnosticsEngine Diags(DiagID, &*DiagOpts, DiagClient);
-+#endif /* LLVM_VERSION_MINOR <= 1 */
-+
-+ // Create the compiler invocation
-+ llvm::OwningPtr<clang::CompilerInvocation> CI(new clang::CompilerInvocation);
-+ clang::CompilerInvocation::CreateFromArgs(*CI,
-+ &args[0],
-+ &args[0] + args.size(),
-+ Diags);
-+
-+ // Create the compiler instance
-+ clang::CompilerInstance Clang;
-+ Clang.setInvocation(CI.take());
-+ // Get ready to report problems
-+ Clang.createDiagnostics(args.size(), &args[0]);
-+ if (!Clang.hasDiagnostics())
-+ return;
-+
-+ // Set Language
-+ clang::LangOptions & lang_opts = Clang.getLangOpts();
-+ lang_opts.OpenCL = 1;
-+
-+ // Create an action and make the compiler instance carry it out
-+ llvm::OwningPtr<clang::CodeGenAction> Act(new clang::EmitLLVMOnlyAction());
-+ if (!Clang.ExecuteAction(*Act))
-+ return;
-+
-+ llvm::Module *module = Act->takeModule();
-+
-+ std::string ErrorInfo;
-+ llvm::raw_fd_ostream OS(output, ErrorInfo,llvm::raw_fd_ostream::F_Binary);
-+ //still write to temp file for code simply, otherwise need add another function.
-+ //because gbe_program_new_from_llvm also be used by cl_program_create_from_llvm, can't be removed
-+ //TODO: Pass module to llvmToGen, if use module, should return Act and use OwningPtr out of this funciton
-+ llvm::WriteBitcodeToFile(module, OS);
-+ OS.close();
-+ }
-+
- extern std::string ocl_stdlib_str;
- extern std::string ocl_common_defines_str;
- static gbe_program programNewFromSource(const char *source,
-@@ -124,26 +206,7 @@ namespace gbe {
- fwrite(source, strlen(source), 1, clFile);
- fclose(clFile);
-
-- // Now compile the code to llvm using clang
--#if 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 ";
--#endif /* LLVM_VERSION_MINOR <= 1 */
-- compileCmd += clName;
-- compileCmd += " ";
-- if(options)
-- compileCmd += options;
-- compileCmd += " -o ";
-- compileCmd += llName;
--
-- // Open a pipe and compile from here. Using Clang API instead is better
-- FILE *pipe = popen(compileCmd.c_str(), "r");
-- FATAL_IF (pipe == NULL, "Unable to run extern compilation command");
-- char msg[256];
-- while (fgets(msg, sizeof(msg), pipe))
-- std::cout << msg;
-- pclose(pipe);
-+ buildModuleFromSource(clName.c_str(), llName.c_str());
- remove(clName.c_str());
-
- // Now build the program from llvm
-diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
-index a0946ef..3fe0cbf 100644
---- a/backend/src/llvm/llvm_gen_backend.cpp
-+++ b/backend/src/llvm/llvm_gen_backend.cpp
-@@ -853,11 +853,8 @@ namespace gbe
- // Insert a new register for each function argument
- #if LLVM_VERSION_MINOR <= 1
- const AttrListPtr &PAL = F.getAttributes();
-- uint32_t argID = 1; // Start at one actually
-- for (; I != E; ++I, ++argID) {
--#else
-- for (; I != E; ++I, ++argID) {
- #endif /* LLVM_VERSION_MINOR <= 1 */
-+ for (; I != E; ++I, ++argID) {
- const std::string &argName = I->getName().str();
- Type *type = I->getType();
-
-@@ -892,7 +889,7 @@ namespace gbe
- PointerType *pointerType = dyn_cast<PointerType>(type);
- // By value structure
- #if LLVM_VERSION_MINOR <= 1
-- if (PAL.paramHasAttr(argID, Attribute::ByVal)) {
-+ if (PAL.paramHasAttr(argID+1, Attribute::ByVal)) {
- #else
- if (I->hasByValAttr()) {
- #endif /* LLVM_VERSION_MINOR <= 1 */
-diff --git a/backend/src/llvm/llvm_scalarize.cpp b/backend/src/llvm/llvm_scalarize.cpp
-index f71401f..ef431f2 100644
---- a/backend/src/llvm/llvm_scalarize.cpp
-+++ b/backend/src/llvm/llvm_scalarize.cpp
-@@ -71,7 +71,11 @@
- #include "llvm/IntrinsicInst.h"
- #include "llvm/Module.h"
- #include "llvm/Pass.h"
-+#if LLVM_VERSION_MINOR <= 1
-+#include "llvm/Support/IRBuilder.h"
-+#else
- #include "llvm/IRBuilder.h"
-+#endif /* LLVM_VERSION_MINOR <= 1 */
- #include "llvm/Support/CallSite.h"
- #include "llvm/Support/CFG.h"
- #include "llvm/Support/raw_ostream.h"
-@@ -730,13 +734,7 @@ namespace gbe {
-
- Function::arg_iterator I = F.arg_begin(), E = F.arg_end();
-
--#if LLVM_VERSION_MINOR <= 1
-- const AttrListPtr &PAL = F.getAttributes();
-- uint32_t argID = 1; // Start at one actually
-- for (; I != E; ++I, ++argID) {
--#else
- for (; I != E; ++I) {
--#endif /* LLVM_VERSION_MINOR <= 1 */
- Type *type = I->getType();
-
- if(type->isVectorTy())
---
-1.7.10.4
-
diff --git a/debian/patches/series b/debian/patches/series
index f33ca92..9971719 100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@ -4,8 +4,5 @@ khronos
private
0001-Generate-all-supported-as_-functions.patch
0002-Define-all-convert_-functions.patch
-0003-Add-long-and-ulong-types-to-conversions.patch
+0003-Add-long-and-ulong-types-to-generated-functions.patch
0004-Add-vector-argument-test-case.patch
-0005-Add-more-get-image-info-functions.patch
-0006-utests-extent-get_image_size-cases-to-other-informat.patch
-0007-Change-clang-system-call-to-libclang-api-call.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