[Pkg-opencl-devel] [beignet] 27/66: Imported Debian patch 0.1+git20130614+89b5e40-1
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 07:27:04 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 d0a602b01fe7b9fa24ed076d348b1215e36a00a1
Merge: 0dad311 dd82e26
Author: Simon Richter <sjr at debian.org>
Date: Fri Jun 14 15:22:18 2013 +0200
Imported Debian patch 0.1+git20130614+89b5e40-1
CMake/FindLLVM.cmake | 31 +-
CMakeLists.txt | 2 +-
README.html | 31 +-
README.md | 33 +-
backend/doc/TODO.html | 7 +-
backend/doc/TODO.md | 12 +-
backend/src/CMakeLists.txt | 1 +
backend/src/backend/context.cpp | 8 +-
backend/src/backend/gen/gen_mesa_disasm.c | 44 ++-
backend/src/backend/gen_context.cpp | 4 +
backend/src/backend/gen_encoder.cpp | 39 +-
backend/src/backend/gen_insn_selection.cpp | 35 +-
backend/src/backend/gen_insn_selection.hpp | 8 +-
backend/src/backend/gen_reg_allocation.cpp | 7 +-
backend/src/backend/program.cpp | 137 ++++++-
backend/src/backend/program.h | 1 +
backend/src/ir/image.cpp | 14 +-
backend/src/ir/instruction.hpp | 10 +-
backend/src/ir/profile.cpp | 1 +
backend/src/ir/profile.hpp | 3 +-
backend/src/ir/unit.cpp | 15 +-
backend/src/ir/unit.hpp | 7 +
backend/src/llvm/llvm_gen_backend.cpp | 106 +++--
backend/src/llvm/llvm_gen_ocl_function.hxx | 4 +
backend/src/llvm/llvm_passes.cpp | 20 +-
backend/src/llvm/llvm_scalarize.cpp | 40 +-
backend/src/llvm/llvm_to_gen.cpp | 13 +-
backend/src/ocl_stdlib.h | 433 ++++++++++++---------
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 +-
kernels/buildin_work_dim.cl | 3 +
kernels/compiler_box_blur.cl | 51 +--
kernels/compiler_box_blur_image.cl | 18 +
kernels/compiler_box_blur_ref.bmp | Bin 49206 -> 49206 bytes
kernels/compiler_clod.cl | 4 +-
kernels/compiler_displacement_map_element.cl | 11 +
kernels/compiler_global_memory_barrier.cl | 7 +
kernels/compiler_group_size.cl | 12 +
kernels/compiler_julia.cl | 4 +-
kernels/compiler_julia_no_break.cl | 4 +-
kernels/compiler_local_memory.cl | 5 -
kernels/compiler_local_memory_two_ptr.cl | 1 +
kernels/compiler_mandelbrot.cl | 6 +-
kernels/compiler_mandelbrot_alternate.cl | 6 +-
kernels/compiler_menger_sponge_no_shadow.cl | 14 +-
kernels/compiler_ribbon.cl | 177 +++++----
kernels/compiler_vector_load_store.cl | 52 ++-
kernels/compiler_write_only_bytes.cl | 13 +-
kernels/compiler_write_only_shorts.cl | 13 +-
kernels/null_kernel_arg.cl | 9 +
kernels/test_cl_finish.cl | 12 +
kernels/test_copy_buffer.cl | 13 +-
kernels/test_copy_buffer_row.cl | 17 +-
kernels/test_get_image_info.cl | 13 +
kernels/test_get_image_size.cl | 9 -
kernels/test_write_only.cl | 13 +-
src/cl_api.c | 244 +++++++++---
src/cl_command_queue.c | 22 +-
src/cl_command_queue.h | 2 +-
src/cl_command_queue_gen7.c | 15 +-
src/cl_context.c | 27 +-
src/cl_device_id.c | 34 +-
src/cl_driver.h | 4 +
src/cl_driver_defs.c | 1 +
src/cl_kernel.c | 63 ++-
src/cl_khr_icd.c | 2 +-
src/cl_mem.c | 129 ++----
src/cl_platform_id.c | 19 +-
src/cl_program.c | 7 -
src/intel/intel_batchbuffer.c | 3 +
src/intel/intel_batchbuffer.h | 2 +
src/intel/intel_gpgpu.c | 31 +-
utests/CMakeLists.txt | 11 +-
utests/buildin_work_dim.cpp | 37 ++
utests/compiler_box_blur_image.cpp | 52 +++
utests/compiler_cl_finish.cpp | 50 +++
utests/compiler_displacement_map_element.cpp | 64 +++
..._image_size.cpp => compiler_get_image_info.cpp} | 25 +-
utests/compiler_global_memory_barrier.cpp | 28 ++
utests/compiler_group_size.cpp | 86 ++++
utests/compiler_local_memory.cpp | 47 ---
utests/compiler_vector_load_store.cpp | 57 ++-
utests/runtime_null_kernel_arg.cpp | 27 ++
utests/utest_helper.hpp | 11 +
91 files changed, 1993 insertions(+), 1689 deletions(-)
diff --cc debian/changelog
index 4ea12cc,0000000..c66d98c
mode 100644,000000..100644
--- a/debian/changelog
+++ b/debian/changelog
@@@ -1,100 -1,0 +1,106 @@@
++beignet (0.1+git20130614+89b5e40-1) experimental; urgency=low
++
++ * New upstream release
++
++ -- Simon Richter <sjr at debian.org> Fri, 14 Jun 2013 15:22:18 +0200
++
+beignet (0.1+git20130521+a7ea35c-1) experimental; urgency=low
+
+ * Rename binary package
+
+ -- Simon Richter <sjr at debian.org> Tue, 21 May 2013 10:48:39 +0200
+
+beignet (0.1+git20130521+a7ea35c-1~prerename) experimental; urgency=low
+
+ * New upstream release
+ * Move libraries to /usr/lib/beignet (should not be used directly)
+
+ -- Simon Richter <sjr at debian.org> Tue, 21 May 2013 09:17:45 +0200
+
+beignet (0.1+git20130514+19e9c58-1) experimental; urgency=low
+
+ * New upstream release
+ * Added a number of tentative patches
+
+ -- Simon Richter <sjr at debian.org> Tue, 14 May 2013 20:04:29 +0200
+
+beignet (0.1+git20130502+63e60ed-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Mon, 06 May 2013 06:30:32 +0200
+
+beignet (0.1+git20130426+0c8f6fe-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Fri, 26 Apr 2013 14:42:21 +0200
+
+beignet (0.1+git20130422+003fac5-2) experimental; urgency=low
+
+ * Add patch for select()
+ * Add patch for fmin() / fmax()
+
+ -- Simon Richter <sjr at debian.org> Mon, 22 Apr 2013 18:26:01 +0200
+
+beignet (0.1+git20130422+003fac5-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Mon, 22 Apr 2013 15:10:54 +0200
+
+beignet (0.1+git20130419+9c11c18-1) experimental; urgency=low
+
+ * Add more functionality patches
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Fri, 19 Apr 2013 14:14:39 +0200
+
+beignet (0.1+git20130418+0546d2e-2) experimental; urgency=low
+
+ * Add functionality patches
+ * Use clang 3.0 command line syntax
+
+ -- Simon Richter <sjr at debian.org> Fri, 19 Apr 2013 09:53:23 +0200
+
+beignet (0.1+git20130418+0546d2e-1) experimental; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Thu, 18 Apr 2013 11:51:37 +0200
+
+beignet (0.1-1) unstable; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Tue, 16 Apr 2013 17:16:18 +0200
+
+beignet (0.0.0+git2013.04.11+e6b503e-1) unstable; urgency=low
+
+ * New upstream release
+
+ -- Simon Richter <sjr at debian.org> Mon, 15 Apr 2013 18:22:45 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-4) unstable; urgency=low
+
+ * Build fix for kfreebsd-*
+
+ -- Simon Richter <sjr at debian.org> Fri, 12 Apr 2013 11:22:36 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-3) unstable; urgency=low
+
+ * Adjust Build-Depends, Architecture list
+
+ -- Simon Richter <sjr at debian.org> Fri, 12 Apr 2013 10:32:36 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-2) unstable; urgency=low
+
+ * Add patch to support size queries in device info
+
+ -- Simon Richter <sjr at debian.org> Thu, 11 Apr 2013 14:00:59 +0200
+
+beignet (0.0.0+git2013.04.01+d1b234c-1) unstable; urgency=low
+
+ * Initial release.
+
+ -- Simon Richter <sjr at debian.org> Tue, 09 Apr 2013 17:14:00 +0200
diff --cc debian/patches/0001-Generate-all-supported-as_-functions.patch
index 5520e63,0000000..45276c2
mode 100644,000000..100644
--- a/debian/patches/0001-Generate-all-supported-as_-functions.patch
+++ b/debian/patches/0001-Generate-all-supported-as_-functions.patch
@@@ -1,1287 -1,0 +1,1307 @@@
- 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
++ size=`IFS=:; set -- dummy $type; echo $3`
++ for vector_length in $VECTOR_LENGTHS; do
++ union_sizes="$union_sizes `expr $vector_length \* $size`"
++ done
++done
++union_sizes="`echo $union_sizes | tr ' ' '\n' | sort -n | uniq`"
++
++# For each union size
++for union_size in $union_sizes; do
++
++ # Define an union that contains all vector types that have the same size as the union
++ unionname="union _type_cast_${union_size}_b"
++ echo "$unionname {"
++ for type in $TYPES; do
++ basetype=`IFS=:; set -- dummy $type; echo $2`
++ basesize=`IFS=:; set -- dummy $type; echo $3`
++ for vector_length in $VECTOR_LENGTHS; do
++ vector_size_in_union="`expr $vector_length \* $basesize`"
++ if test $union_size -ne $vector_size_in_union; then
++ continue
++ fi
++ if test $vector_length -eq 1; then
++ vectortype=$basetype
++ else
++ vectortype=$basetype$vector_length
++ fi
++ echo " $vectortype _$vectortype;"
++ done
++
++ done
++ echo "};"
++ echo
++
++ # For each tuple of vector types that has the same size as the current union size,
++ # define an as_* function that converts types without changing binary representation.
++ for ftype in $TYPES; do
++ fbasetype=`IFS=:; set -- dummy $ftype; echo $2`
++ fbasesize=`IFS=:; set -- dummy $ftype; echo $3`
++ for fvector_length in $VECTOR_LENGTHS; do
++ fvector_size_in_union="`expr $fvector_length \* $fbasesize`"
++ if test $union_size -ne $fvector_size_in_union; then
++ continue
++ fi
++ if test $fvector_length -eq 1; then
++ fvectortype=$fbasetype
++ else
++ fvectortype=$fbasetype$fvector_length
++ fi
++ for ttype in $TYPES; do
++ tbasetype=`IFS=:; set -- dummy $ttype; echo $2`
++ tbasesize=`IFS=:; set -- dummy $ttype; echo $3`
++ if test $fbasetype = $tbasetype; then
++ continue
++ fi
++ for tvector_length in $VECTOR_LENGTHS; do
++ tvector_size_in_union="`expr $tvector_length \* $tbasesize`"
++ if test $union_size -ne $tvector_size_in_union; then
++ continue
++ fi
++ if test $tvector_length -eq 1; then
++ tvectortype=$tbasetype
++ else
++ tvectortype=$tbasetype$tvector_length
++ fi
++ echo "INLINE OVERLOADABLE $tvectortype as_$tvectortype($fvectortype v) {"
++ echo " $unionname u;"
++ echo " u._$fvectortype = v;"
++ echo " return u._$tvectortype;"
++ echo "}"
++ echo
++ done
++ done
++ done
++
++ 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
+ /////////////////////////////////////////////////////////////////////////////
+-union type_cast_4_b {
+- float f;
+- uchar4 u4;
++
- +// ##BEGIN_CONVERSIONS##
+++// ##BEGIN_AS##
++union _type_cast_1_b {
++ char _char;
++ uchar _uchar;
++};
++
++INLINE OVERLOADABLE uchar as_uchar(char v) {
++ union _type_cast_1_b u;
++ u._char = v;
++ return u._uchar;
++}
++
++INLINE OVERLOADABLE char as_char(uchar v) {
++ union _type_cast_1_b u;
++ u._uchar = v;
++ return u._char;
++}
++
++union _type_cast_2_b {
++ short _short;
++ 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;
++ u._short = v;
++ return u._ushort;
++}
++
++INLINE OVERLOADABLE char2 as_char2(short v) {
++ union _type_cast_2_b u;
++ u._short = v;
++ return u._char2;
++}
++
++INLINE OVERLOADABLE uchar2 as_uchar2(short v) {
++ union _type_cast_2_b u;
++ u._short = v;
++ return u._uchar2;
++}
++
++INLINE OVERLOADABLE short as_short(ushort v) {
++ union _type_cast_2_b u;
++ u._ushort = v;
++ return u._short;
++}
++
++INLINE OVERLOADABLE char2 as_char2(ushort v) {
++ union _type_cast_2_b u;
++ u._ushort = v;
++ return u._char2;
++}
++
++INLINE OVERLOADABLE uchar2 as_uchar2(ushort v) {
++ union _type_cast_2_b u;
++ u._ushort = v;
++ return u._uchar2;
++}
++
++INLINE OVERLOADABLE short as_short(char2 v) {
++ union _type_cast_2_b u;
++ u._char2 = v;
++ return u._short;
++}
++
++INLINE OVERLOADABLE ushort as_ushort(char2 v) {
++ union _type_cast_2_b u;
++ u._char2 = v;
++ return u._ushort;
++}
++
++INLINE OVERLOADABLE uchar2 as_uchar2(char2 v) {
++ union _type_cast_2_b u;
++ u._char2 = v;
++ return u._uchar2;
++}
++
++INLINE OVERLOADABLE short as_short(uchar2 v) {
++ union _type_cast_2_b u;
++ u._uchar2 = v;
++ return u._short;
++}
++
++INLINE OVERLOADABLE ushort as_ushort(uchar2 v) {
++ union _type_cast_2_b u;
++ u._uchar2 = v;
++ return u._ushort;
++}
++
++INLINE OVERLOADABLE char2 as_char2(uchar2 v) {
++ union _type_cast_2_b u;
++ u._uchar2 = v;
++ return u._char2;
++}
++
++union _type_cast_3_b {
++ char3 _char3;
++ uchar3 _uchar3;
++};
++
++INLINE OVERLOADABLE uchar3 as_uchar3(char3 v) {
++ union _type_cast_3_b u;
++ u._char3 = v;
++ return u._uchar3;
++}
++
++INLINE OVERLOADABLE char3 as_char3(uchar3 v) {
++ union _type_cast_3_b u;
++ u._uchar3 = v;
++ return u._char3;
++}
++
++union _type_cast_4_b {
++ int _int;
++ uint _uint;
++ short2 _short2;
++ ushort2 _ushort2;
++ char4 _char4;
++ uchar4 _uchar4;
++ float _float;
++};
++
++INLINE OVERLOADABLE uint as_uint(int v) {
++ union _type_cast_4_b u;
++ u._int = v;
++ return u._uint;
++}
++
++INLINE OVERLOADABLE short2 as_short2(int v) {
++ union _type_cast_4_b u;
++ u._int = v;
++ return u._short2;
++}
++
++INLINE OVERLOADABLE ushort2 as_ushort2(int v) {
++ union _type_cast_4_b u;
++ u._int = v;
++ return u._ushort2;
++}
++
++INLINE OVERLOADABLE char4 as_char4(int v) {
++ union _type_cast_4_b u;
++ u._int = v;
++ return u._char4;
++}
++
++INLINE OVERLOADABLE uchar4 as_uchar4(int v) {
++ union _type_cast_4_b u;
++ u._int = v;
++ return u._uchar4;
++}
++
++INLINE OVERLOADABLE float as_float(int v) {
++ union _type_cast_4_b u;
++ u._int = v;
++ return u._float;
++}
++
++INLINE OVERLOADABLE int as_int(uint v) {
++ union _type_cast_4_b u;
++ u._uint = v;
++ return u._int;
++}
++
++INLINE OVERLOADABLE short2 as_short2(uint v) {
++ union _type_cast_4_b u;
++ u._uint = v;
++ return u._short2;
++}
++
++INLINE OVERLOADABLE ushort2 as_ushort2(uint v) {
++ union _type_cast_4_b u;
++ u._uint = v;
++ return u._ushort2;
++}
++
++INLINE OVERLOADABLE char4 as_char4(uint v) {
++ union _type_cast_4_b u;
++ u._uint = v;
++ return u._char4;
++}
++
++INLINE OVERLOADABLE uchar4 as_uchar4(uint v) {
++ union _type_cast_4_b u;
++ u._uint = v;
++ return u._uchar4;
++}
++
++INLINE OVERLOADABLE float as_float(uint v) {
++ union _type_cast_4_b u;
++ u._uint = v;
++ return u._float;
++}
++
++INLINE OVERLOADABLE int as_int(short2 v) {
++ union _type_cast_4_b u;
++ u._short2 = v;
++ return u._int;
++}
++
++INLINE OVERLOADABLE uint as_uint(short2 v) {
++ union _type_cast_4_b u;
++ u._short2 = v;
++ return u._uint;
++}
++
++INLINE OVERLOADABLE ushort2 as_ushort2(short2 v) {
++ union _type_cast_4_b u;
++ u._short2 = v;
++ return u._ushort2;
++}
++
++INLINE OVERLOADABLE char4 as_char4(short2 v) {
++ union _type_cast_4_b u;
++ u._short2 = v;
++ return u._char4;
++}
++
++INLINE OVERLOADABLE uchar4 as_uchar4(short2 v) {
++ union _type_cast_4_b u;
++ u._short2 = v;
++ return u._uchar4;
++}
++
++INLINE OVERLOADABLE float as_float(short2 v) {
++ union _type_cast_4_b u;
++ u._short2 = v;
++ return u._float;
++}
++
++INLINE OVERLOADABLE int as_int(ushort2 v) {
++ union _type_cast_4_b u;
++ u._ushort2 = v;
++ return u._int;
++}
++
++INLINE OVERLOADABLE uint as_uint(ushort2 v) {
++ union _type_cast_4_b u;
++ u._ushort2 = v;
++ return u._uint;
++}
++
++INLINE OVERLOADABLE short2 as_short2(ushort2 v) {
++ union _type_cast_4_b u;
++ u._ushort2 = v;
++ return u._short2;
++}
++
++INLINE OVERLOADABLE char4 as_char4(ushort2 v) {
++ union _type_cast_4_b u;
++ u._ushort2 = v;
++ return u._char4;
++}
++
++INLINE OVERLOADABLE uchar4 as_uchar4(ushort2 v) {
++ union _type_cast_4_b u;
++ u._ushort2 = v;
++ return u._uchar4;
++}
++
++INLINE OVERLOADABLE float as_float(ushort2 v) {
++ union _type_cast_4_b u;
++ u._ushort2 = v;
++ return u._float;
++}
++
++INLINE OVERLOADABLE int as_int(char4 v) {
++ union _type_cast_4_b u;
++ u._char4 = v;
++ return u._int;
++}
++
++INLINE OVERLOADABLE uint as_uint(char4 v) {
++ union _type_cast_4_b u;
++ u._char4 = v;
++ return u._uint;
++}
++
++INLINE OVERLOADABLE short2 as_short2(char4 v) {
++ union _type_cast_4_b u;
++ u._char4 = v;
++ return u._short2;
++}
++
++INLINE OVERLOADABLE ushort2 as_ushort2(char4 v) {
++ union _type_cast_4_b u;
++ u._char4 = v;
++ return u._ushort2;
++}
++
++INLINE OVERLOADABLE uchar4 as_uchar4(char4 v) {
++ union _type_cast_4_b u;
++ u._char4 = v;
++ return u._uchar4;
++}
++
++INLINE OVERLOADABLE float as_float(char4 v) {
++ union _type_cast_4_b u;
++ u._char4 = v;
++ return u._float;
++}
++
++INLINE OVERLOADABLE int as_int(uchar4 v) {
++ union _type_cast_4_b u;
++ u._uchar4 = v;
++ return u._int;
++}
++
++INLINE OVERLOADABLE uint as_uint(uchar4 v) {
++ union _type_cast_4_b u;
++ u._uchar4 = v;
++ return u._uint;
++}
++
++INLINE OVERLOADABLE short2 as_short2(uchar4 v) {
++ union _type_cast_4_b u;
++ u._uchar4 = v;
++ return u._short2;
++}
++
++INLINE OVERLOADABLE ushort2 as_ushort2(uchar4 v) {
++ union _type_cast_4_b u;
++ u._uchar4 = v;
++ return u._ushort2;
++}
++
++INLINE OVERLOADABLE char4 as_char4(uchar4 v) {
++ union _type_cast_4_b u;
++ u._uchar4 = v;
++ return u._char4;
++}
++
++INLINE OVERLOADABLE float as_float(uchar4 v) {
++ union _type_cast_4_b u;
++ u._uchar4 = v;
++ return u._float;
++}
++
++INLINE OVERLOADABLE int as_int(float v) {
++ union _type_cast_4_b u;
++ u._float = v;
++ return u._int;
++}
++
++INLINE OVERLOADABLE uint as_uint(float v) {
++ union _type_cast_4_b u;
++ u._float = v;
++ return u._uint;
++}
++
++INLINE OVERLOADABLE short2 as_short2(float v) {
++ union _type_cast_4_b u;
++ u._float = v;
++ return u._short2;
++}
++
++INLINE OVERLOADABLE ushort2 as_ushort2(float v) {
++ union _type_cast_4_b u;
++ u._float = v;
++ return u._ushort2;
++}
++
++INLINE OVERLOADABLE char4 as_char4(float v) {
++ union _type_cast_4_b u;
++ u._float = v;
++ return u._char4;
++}
++
++INLINE OVERLOADABLE uchar4 as_uchar4(float v) {
++ union _type_cast_4_b u;
++ u._float = v;
++ return u._uchar4;
++}
++
++union _type_cast_6_b {
++ short3 _short3;
++ ushort3 _ushort3;
++};
++
++INLINE OVERLOADABLE ushort3 as_ushort3(short3 v) {
++ union _type_cast_6_b u;
++ u._short3 = v;
++ return u._ushort3;
++}
++
++INLINE OVERLOADABLE short3 as_short3(ushort3 v) {
++ union _type_cast_6_b u;
++ u._ushort3 = v;
++ return u._short3;
++}
++
++union _type_cast_8_b {
++ int2 _int2;
++ uint2 _uint2;
++ short4 _short4;
++ ushort4 _ushort4;
++ char8 _char8;
++ uchar8 _uchar8;
++ float2 _float2;
++};
++
++INLINE OVERLOADABLE uint2 as_uint2(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._ushort4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE uchar8 as_uchar8(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._uchar8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE int2 as_int2(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._ushort4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE uchar8 as_uchar8(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._uchar8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE int2 as_int2(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE uint2 as_uint2(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._ushort4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE uchar8 as_uchar8(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._uchar8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE int2 as_int2(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE uint2 as_uint2(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE uchar8 as_uchar8(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._uchar8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE int2 as_int2(char8 v) {
++ union _type_cast_8_b u;
++ u._char8 = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE uint2 as_uint2(char8 v) {
++ union _type_cast_8_b u;
++ u._char8 = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(char8 v) {
++ union _type_cast_8_b u;
++ u._char8 = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(char8 v) {
++ 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;
++ u._char8 = v;
++ return u._uchar8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(char8 v) {
++ union _type_cast_8_b u;
++ u._char8 = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE int2 as_int2(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE uint2 as_uint2(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._ushort4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE int2 as_int2(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE uint2 as_uint2(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._ushort4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE uchar8 as_uchar8(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._uchar8;
++}
++
++union _type_cast_12_b {
++ int3 _int3;
++ uint3 _uint3;
++ float3 _float3;
++};
++
++INLINE OVERLOADABLE uint3 as_uint3(int3 v) {
++ union _type_cast_12_b u;
++ u._int3 = v;
++ return u._uint3;
++}
++
++INLINE OVERLOADABLE float3 as_float3(int3 v) {
++ union _type_cast_12_b u;
++ u._int3 = v;
++ return u._float3;
++}
++
++INLINE OVERLOADABLE int3 as_int3(uint3 v) {
++ union _type_cast_12_b u;
++ u._uint3 = v;
++ return u._int3;
++}
++
++INLINE OVERLOADABLE float3 as_float3(uint3 v) {
++ union _type_cast_12_b u;
++ u._uint3 = v;
++ return u._float3;
++}
++
++INLINE OVERLOADABLE int3 as_int3(float3 v) {
++ union _type_cast_12_b u;
++ u._float3 = v;
++ return u._int3;
++}
++
++INLINE OVERLOADABLE uint3 as_uint3(float3 v) {
++ union _type_cast_12_b u;
++ u._float3 = v;
++ return u._uint3;
++}
++
++union _type_cast_16_b {
++ int4 _int4;
++ uint4 _uint4;
++ short8 _short8;
++ ushort8 _ushort8;
++ char16 _char16;
++ uchar16 _uchar16;
++ float4 _float4;
++};
++
++INLINE OVERLOADABLE uint4 as_uint4(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._uchar16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE int4 as_int4(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._uchar16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE int4 as_int4(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE uint4 as_uint4(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._uchar16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE int4 as_int4(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE uint4 as_uint4(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._uchar16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE int4 as_int4(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE uint4 as_uint4(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._uchar16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE int4 as_int4(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE uint4 as_uint4(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE int4 as_int4(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE uint4 as_uint4(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._uchar16;
++}
++
++union _type_cast_32_b {
++ int8 _int8;
++ uint8 _uint8;
++ 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;
++ u._int8 = v;
++ return u._uint8;
++}
++
++INLINE OVERLOADABLE short16 as_short16(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._short16;
++}
++
++INLINE OVERLOADABLE ushort16 as_ushort16(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._ushort16;
++}
++
++INLINE OVERLOADABLE float8 as_float8(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._float8;
++}
++
++INLINE OVERLOADABLE int8 as_int8(uint8 v) {
++ 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;
++ u._uint8 = v;
++ return u._short16;
++}
++
++INLINE OVERLOADABLE ushort16 as_ushort16(uint8 v) {
++ union _type_cast_32_b u;
++ u._uint8 = v;
++ return u._ushort16;
++}
++
++INLINE OVERLOADABLE float8 as_float8(uint8 v) {
++ union _type_cast_32_b u;
++ u._uint8 = v;
++ return u._float8;
++}
++
++INLINE OVERLOADABLE int8 as_int8(short16 v) {
++ union _type_cast_32_b u;
++ u._short16 = v;
++ return u._int8;
++}
++
++INLINE OVERLOADABLE uint8 as_uint8(short16 v) {
++ union _type_cast_32_b u;
++ u._short16 = v;
++ return u._uint8;
++}
++
++INLINE OVERLOADABLE ushort16 as_ushort16(short16 v) {
++ union _type_cast_32_b u;
++ u._short16 = v;
++ return u._ushort16;
++}
++
++INLINE OVERLOADABLE float8 as_float8(short16 v) {
++ union _type_cast_32_b u;
++ u._short16 = v;
++ return u._float8;
++}
++
++INLINE OVERLOADABLE int8 as_int8(ushort16 v) {
++ union _type_cast_32_b u;
++ u._ushort16 = v;
++ return u._int8;
++}
++
++INLINE OVERLOADABLE uint8 as_uint8(ushort16 v) {
++ union _type_cast_32_b u;
++ u._ushort16 = v;
++ return u._uint8;
++}
++
++INLINE OVERLOADABLE short16 as_short16(ushort16 v) {
++ union _type_cast_32_b u;
++ u._ushort16 = v;
++ return u._short16;
++}
++
++INLINE OVERLOADABLE float8 as_float8(ushort16 v) {
++ union _type_cast_32_b u;
++ u._ushort16 = v;
++ return u._float8;
++}
++
++INLINE OVERLOADABLE int8 as_int8(float8 v) {
++ union _type_cast_32_b u;
++ u._float8 = v;
++ return u._int8;
++}
++
++INLINE OVERLOADABLE uint8 as_uint8(float8 v) {
++ union _type_cast_32_b u;
++ u._float8 = v;
++ return u._uint8;
++}
++
++INLINE OVERLOADABLE short16 as_short16(float8 v) {
++ union _type_cast_32_b u;
++ u._float8 = v;
++ return u._short16;
++}
++
++INLINE OVERLOADABLE ushort16 as_ushort16(float8 v) {
++ union _type_cast_32_b u;
++ u._float8 = v;
++ return u._ushort16;
++}
++
++union _type_cast_64_b {
++ int16 _int16;
++ uint16 _uint16;
++ float16 _float16;
++};
++
++INLINE OVERLOADABLE uint16 as_uint16(int16 v) {
++ union _type_cast_64_b u;
++ u._int16 = v;
++ return u._uint16;
++}
++
++INLINE OVERLOADABLE float16 as_float16(int16 v) {
++ union _type_cast_64_b u;
++ u._int16 = v;
++ return u._float16;
++}
++
++INLINE OVERLOADABLE int16 as_int16(uint16 v) {
++ union _type_cast_64_b u;
++ u._uint16 = v;
++ return u._int16;
++}
++
++INLINE OVERLOADABLE float16 as_float16(uint16 v) {
++ union _type_cast_64_b u;
++ u._uint16 = v;
++ return u._float16;
++}
++
++INLINE OVERLOADABLE int16 as_int16(float16 v) {
++ union _type_cast_64_b u;
++ u._float16 = v;
++ return u._int16;
++}
++
++INLINE OVERLOADABLE uint16 as_uint16(float16 v) {
++ union _type_cast_64_b u;
++ u._float16 = v;
++ 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
+--
+1.7.10.4
+
diff --cc debian/patches/0002-Define-all-convert_-functions.patch
index f47e282,0000000..a0e5c20
mode 100644,000000..100644
--- a/debian/patches/0002-Define-all-convert_-functions.patch
+++ b/debian/patches/0002-Define-all-convert_-functions.patch
@@@ -1,949 -1,0 +1,980 @@@
- 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
++ if test $vector_length -eq 1; then
++ continue;
++ fi
++ for ftype in $TYPES; do
++ fbasetype=`IFS=:; set -- dummy $ftype; echo $2`
++ for ttype in $TYPES; do
++ tbasetype=`IFS=:; set -- dummy $ttype; echo $2`
++ if test $fbasetype = $tbasetype; then
++ continue
++ fi
++ fvectortype=$fbasetype$vector_length
++ tvectortype=$tbasetype$vector_length
++ construct="($tbasetype)(v.s0)"
++ if test $vector_length -gt 1; then
++ construct="$construct, ($tbasetype)(v.s1)"
++ fi
++ if test $vector_length -gt 2; then
++ construct="$construct, ($tbasetype)(v.s2)"
++ fi
++ if test $vector_length -gt 3; then
++ construct="$construct, ($tbasetype)(v.s3)"
++ fi
++ if test $vector_length -gt 4; then
++ construct="$construct, ($tbasetype)(v.s4)"
++ construct="$construct, ($tbasetype)(v.s5)"
++ construct="$construct, ($tbasetype)(v.s6)"
++ construct="$construct, ($tbasetype)(v.s7)"
++ fi
++ if test $vector_length -gt 8; then
++ construct="$construct, ($tbasetype)(v.s8)"
++ construct="$construct, ($tbasetype)(v.s9)"
++ construct="$construct, ($tbasetype)(v.sA)"
++ construct="$construct, ($tbasetype)(v.sB)"
++ construct="$construct, ($tbasetype)(v.sC)"
++ construct="$construct, ($tbasetype)(v.sD)"
++ construct="$construct, ($tbasetype)(v.sE)"
++ construct="$construct, ($tbasetype)(v.sF)"
++ fi
++
++ echo "INLINE OVERLOADABLE $tvectortype convert_$tvectortype($fvectortype v) {"
++ echo " return ($tvectortype)($construct);"
++ echo "}"
++ echo
++ done
++ 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));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(int2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(int2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(int2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(int2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(int2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(uint2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(uint2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(uint2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(uint2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(uint2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(uint2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(short2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE uint2 convert_uint2(short2 v) {
++ return (uint2)((uint)(v.s0), (uint)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(short2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(short2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(short2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(short2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(ushort2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE uint2 convert_uint2(ushort2 v) {
++ return (uint2)((uint)(v.s0), (uint)(v.s1));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(ushort2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(ushort2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(ushort2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(ushort2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(char2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE uint2 convert_uint2(char2 v) {
++ return (uint2)((uint)(v.s0), (uint)(v.s1));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(char2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(char2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(char2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(char2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(uchar2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE uint2 convert_uint2(uchar2 v) {
++ return (uint2)((uint)(v.s0), (uint)(v.s1));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(uchar2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(uchar2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(uchar2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(uchar2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(float2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE uint2 convert_uint2(float2 v) {
++ return (uint2)((uint)(v.s0), (uint)(v.s1));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(float2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(float2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(float2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(float2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(int3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(int3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(int3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(int3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(int3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
++}
++
++INLINE OVERLOADABLE float3 convert_float3(int3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(uint3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(uint3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(uint3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(uint3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(uint3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
++}
++
++INLINE OVERLOADABLE float3 convert_float3(uint3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(short3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(short3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(short3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(short3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(short3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
++}
++
++INLINE OVERLOADABLE float3 convert_float3(short3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(ushort3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(ushort3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(ushort3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(ushort3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(ushort3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
++}
++
++INLINE OVERLOADABLE float3 convert_float3(ushort3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(char3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(char3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(char3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(char3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(char3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
++}
++
++INLINE OVERLOADABLE float3 convert_float3(char3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(uchar3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(uchar3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(uchar3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(uchar3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(uchar3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE float3 convert_float3(uchar3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(float3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(float3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(float3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(float3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(float3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(float3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(int4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(int4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(int4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(int4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(int4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(int4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(uint4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(uint4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(uint4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(uint4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(uint4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(uint4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(short4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(short4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(short4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(short4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(short4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(short4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(ushort4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(ushort4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(ushort4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(ushort4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(ushort4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(ushort4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(char4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(char4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(char4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(char4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(char4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(char4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(uchar4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(uchar4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(uchar4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(uchar4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(uchar4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(uchar4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(float4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(float4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(float4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(float4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(float4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(float4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(int8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(int8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(int8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++INLINE OVERLOADABLE uchar8 convert_uchar8(int8 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(uint8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(uint8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(uint8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++INLINE OVERLOADABLE uchar8 convert_uchar8(uint8 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint8 convert_uint8(short8 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));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(short8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(short8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++INLINE OVERLOADABLE uchar8 convert_uchar8(short8 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint8 convert_uint8(ushort8 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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(ushort8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(ushort8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++INLINE OVERLOADABLE uchar8 convert_uchar8(ushort8 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint8 convert_uint8(char8 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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(char8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(char8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE uchar8 convert_uchar8(char8 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint8 convert_uint8(uchar8 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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(uchar8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(uchar8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(uchar8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint8 convert_uint8(float8 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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(float8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(float8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(float8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(int16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(int16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(int16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(int16 v) {
++ 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(uint16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(uint16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(uint16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(uint16 v) {
++ 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint16 convert_uint16(short16 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));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(short16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(short16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(short16 v) {
++ 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint16 convert_uint16(ushort16 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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(ushort16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(ushort16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(ushort16 v) {
++ 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint16 convert_uint16(char16 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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(char16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(char16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(char16 v) {
++ 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));
++}
++
++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));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE uint16 convert_uint16(uchar16 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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(uchar16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(uchar16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(uchar16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++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));
++}
++
++INLINE OVERLOADABLE int16 convert_int16(float16 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));
++}
++
++INLINE OVERLOADABLE uint16 convert_uint16(float16 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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(float16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(float16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(float16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(float16 v) {
++ 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 --cc debian/patches/0003-Add-long-and-ulong-types-to-generated-functions.patch
index c2253ee,0000000..7e7f81d
mode 100644,000000..100644
--- a/debian/patches/0003-Add-long-and-ulong-types-to-generated-functions.patch
+++ b/debian/patches/0003-Add-long-and-ulong-types-to-generated-functions.patch
@@@ -1,1722 -1,0 +1,1725 @@@
- 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"
++TYPES="long:8 ulong:8 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"
+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 {
++ long _long;
++ ulong _ulong;
+ 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;
+ };
+
++INLINE OVERLOADABLE ulong as_ulong(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._ulong;
++}
++
++INLINE OVERLOADABLE int2 as_int2(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE uint2 as_uint2(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._ushort4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE uchar8 as_uchar8(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._uchar8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(long v) {
++ union _type_cast_8_b u;
++ u._long = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE long as_long(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE int2 as_int2(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._int2;
++}
++
++INLINE OVERLOADABLE uint2 as_uint2(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._uint2;
++}
++
++INLINE OVERLOADABLE short4 as_short4(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._short4;
++}
++
++INLINE OVERLOADABLE ushort4 as_ushort4(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._ushort4;
++}
++
++INLINE OVERLOADABLE char8 as_char8(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._char8;
++}
++
++INLINE OVERLOADABLE uchar8 as_uchar8(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._uchar8;
++}
++
++INLINE OVERLOADABLE float2 as_float2(ulong v) {
++ union _type_cast_8_b u;
++ u._ulong = v;
++ return u._float2;
++}
++
++INLINE OVERLOADABLE long as_long(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE ulong as_ulong(int2 v) {
++ union _type_cast_8_b u;
++ u._int2 = v;
++ return u._ulong;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long as_long(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE ulong as_ulong(uint2 v) {
++ union _type_cast_8_b u;
++ u._uint2 = v;
++ return u._ulong;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long as_long(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE ulong as_ulong(short4 v) {
++ union _type_cast_8_b u;
++ u._short4 = v;
++ return u._ulong;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long as_long(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE ulong as_ulong(ushort4 v) {
++ union _type_cast_8_b u;
++ u._ushort4 = v;
++ return u._ulong;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long as_long(char8 v) {
++ union _type_cast_8_b u;
++ u._char8 = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE ulong as_ulong(char8 v) {
++ union _type_cast_8_b u;
++ u._char8 = v;
++ return u._ulong;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long as_long(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE ulong as_ulong(uchar8 v) {
++ union _type_cast_8_b u;
++ u._uchar8 = v;
++ return u._ulong;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long as_long(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._long;
++}
++
++INLINE OVERLOADABLE ulong as_ulong(float2 v) {
++ union _type_cast_8_b u;
++ u._float2 = v;
++ return u._ulong;
++}
++
+ 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 {
++ long2 _long2;
++ ulong2 _ulong2;
+ 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;
+ };
+
++INLINE OVERLOADABLE ulong2 as_ulong2(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._ulong2;
++}
++
++INLINE OVERLOADABLE int4 as_int4(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE uint4 as_uint4(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._uchar16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(long2 v) {
++ union _type_cast_16_b u;
++ u._long2 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE long2 as_long2(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE int4 as_int4(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._int4;
++}
++
++INLINE OVERLOADABLE uint4 as_uint4(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._uint4;
++}
++
++INLINE OVERLOADABLE short8 as_short8(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._short8;
++}
++
++INLINE OVERLOADABLE ushort8 as_ushort8(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._ushort8;
++}
++
++INLINE OVERLOADABLE char16 as_char16(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._char16;
++}
++
++INLINE OVERLOADABLE uchar16 as_uchar16(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._uchar16;
++}
++
++INLINE OVERLOADABLE float4 as_float4(ulong2 v) {
++ union _type_cast_16_b u;
++ u._ulong2 = v;
++ return u._float4;
++}
++
++INLINE OVERLOADABLE long2 as_long2(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE ulong2 as_ulong2(int4 v) {
++ union _type_cast_16_b u;
++ u._int4 = v;
++ return u._ulong2;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long2 as_long2(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE ulong2 as_ulong2(uint4 v) {
++ union _type_cast_16_b u;
++ u._uint4 = v;
++ return u._ulong2;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long2 as_long2(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE ulong2 as_ulong2(short8 v) {
++ union _type_cast_16_b u;
++ u._short8 = v;
++ return u._ulong2;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long2 as_long2(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE ulong2 as_ulong2(ushort8 v) {
++ union _type_cast_16_b u;
++ u._ushort8 = v;
++ return u._ulong2;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long2 as_long2(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE ulong2 as_ulong2(char16 v) {
++ union _type_cast_16_b u;
++ u._char16 = v;
++ return u._ulong2;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long2 as_long2(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE ulong2 as_ulong2(uchar16 v) {
++ union _type_cast_16_b u;
++ u._uchar16 = v;
++ return u._ulong2;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long2 as_long2(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._long2;
++}
++
++INLINE OVERLOADABLE ulong2 as_ulong2(float4 v) {
++ union _type_cast_16_b u;
++ u._float4 = v;
++ return u._ulong2;
++}
++
+ 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;
+ }
+
++union _type_cast_24_b {
++ long3 _long3;
++ ulong3 _ulong3;
++};
++
++INLINE OVERLOADABLE ulong3 as_ulong3(long3 v) {
++ union _type_cast_24_b u;
++ u._long3 = v;
++ return u._ulong3;
++}
++
++INLINE OVERLOADABLE long3 as_long3(ulong3 v) {
++ union _type_cast_24_b u;
++ u._ulong3 = v;
++ return u._long3;
++}
++
+ union _type_cast_32_b {
++ long4 _long4;
++ ulong4 _ulong4;
+ 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;
+ };
+
+-INLINE OVERLOADABLE uint8 as_uint8(int8 v) {
++INLINE OVERLOADABLE ulong4 as_ulong4(long4 v) {
+ union _type_cast_32_b u;
+- u._int8 = v;
++ u._long4 = v;
++ return u._ulong4;
++}
++
++INLINE OVERLOADABLE int8 as_int8(long4 v) {
++ union _type_cast_32_b u;
++ u._long4 = v;
++ return u._int8;
++}
++
++INLINE OVERLOADABLE uint8 as_uint8(long4 v) {
++ union _type_cast_32_b u;
++ u._long4 = v;
+ return u._uint8;
+ }
+
+-INLINE OVERLOADABLE short16 as_short16(int8 v) {
++INLINE OVERLOADABLE short16 as_short16(long4 v) {
+ union _type_cast_32_b u;
+- u._int8 = v;
++ u._long4 = v;
+ return u._short16;
+ }
+
+-INLINE OVERLOADABLE ushort16 as_ushort16(int8 v) {
++INLINE OVERLOADABLE ushort16 as_ushort16(long4 v) {
+ union _type_cast_32_b u;
+- u._int8 = v;
++ u._long4 = v;
+ return u._ushort16;
+ }
+
+-INLINE OVERLOADABLE float8 as_float8(int8 v) {
++INLINE OVERLOADABLE float8 as_float8(long4 v) {
++ union _type_cast_32_b u;
++ u._long4 = v;
++ return u._float8;
++}
++
++INLINE OVERLOADABLE long4 as_long4(ulong4 v) {
++ union _type_cast_32_b u;
++ u._ulong4 = v;
++ return u._long4;
++}
++
++INLINE OVERLOADABLE int8 as_int8(ulong4 v) {
++ union _type_cast_32_b u;
++ u._ulong4 = v;
++ return u._int8;
++}
++
++INLINE OVERLOADABLE uint8 as_uint8(ulong4 v) {
++ union _type_cast_32_b u;
++ u._ulong4 = v;
++ return u._uint8;
++}
++
++INLINE OVERLOADABLE short16 as_short16(ulong4 v) {
++ union _type_cast_32_b u;
++ u._ulong4 = v;
++ return u._short16;
++}
++
++INLINE OVERLOADABLE ushort16 as_ushort16(ulong4 v) {
++ union _type_cast_32_b u;
++ u._ulong4 = v;
++ return u._ushort16;
++}
++
++INLINE OVERLOADABLE float8 as_float8(ulong4 v) {
++ union _type_cast_32_b u;
++ u._ulong4 = v;
++ return u._float8;
++}
++
++INLINE OVERLOADABLE long4 as_long4(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._long4;
++}
++
++INLINE OVERLOADABLE ulong4 as_ulong4(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._ulong4;
++}
++
++INLINE OVERLOADABLE uint8 as_uint8(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._uint8;
++}
++
++INLINE OVERLOADABLE short16 as_short16(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._short16;
++}
++
++INLINE OVERLOADABLE ushort16 as_ushort16(int8 v) {
++ union _type_cast_32_b u;
++ u._int8 = v;
++ return u._ushort16;
++}
++
++INLINE OVERLOADABLE float8 as_float8(int8 v) {
+ union _type_cast_32_b u;
+ u._int8 = v;
+ return u._float8;
+ }
+
++INLINE OVERLOADABLE long4 as_long4(uint8 v) {
++ union _type_cast_32_b u;
++ u._uint8 = v;
++ return u._long4;
++}
++
++INLINE OVERLOADABLE ulong4 as_ulong4(uint8 v) {
++ union _type_cast_32_b u;
++ u._uint8 = v;
++ return u._ulong4;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long4 as_long4(short16 v) {
++ union _type_cast_32_b u;
++ u._short16 = v;
++ return u._long4;
++}
++
++INLINE OVERLOADABLE ulong4 as_ulong4(short16 v) {
++ union _type_cast_32_b u;
++ u._short16 = v;
++ return u._ulong4;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long4 as_long4(ushort16 v) {
++ union _type_cast_32_b u;
++ u._ushort16 = v;
++ return u._long4;
++}
++
++INLINE OVERLOADABLE ulong4 as_ulong4(ushort16 v) {
++ union _type_cast_32_b u;
++ u._ushort16 = v;
++ return u._ulong4;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long4 as_long4(float8 v) {
++ union _type_cast_32_b u;
++ u._float8 = v;
++ return u._long4;
++}
++
++INLINE OVERLOADABLE ulong4 as_ulong4(float8 v) {
++ union _type_cast_32_b u;
++ u._float8 = v;
++ return u._ulong4;
++}
++
+ 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 {
++ long8 _long8;
++ ulong8 _ulong8;
+ int16 _int16;
+ uint16 _uint16;
+ float16 _float16;
+ };
+
++INLINE OVERLOADABLE ulong8 as_ulong8(long8 v) {
++ union _type_cast_64_b u;
++ u._long8 = v;
++ return u._ulong8;
++}
++
++INLINE OVERLOADABLE int16 as_int16(long8 v) {
++ union _type_cast_64_b u;
++ u._long8 = v;
++ return u._int16;
++}
++
++INLINE OVERLOADABLE uint16 as_uint16(long8 v) {
++ union _type_cast_64_b u;
++ u._long8 = v;
++ return u._uint16;
++}
++
++INLINE OVERLOADABLE float16 as_float16(long8 v) {
++ union _type_cast_64_b u;
++ u._long8 = v;
++ return u._float16;
++}
++
++INLINE OVERLOADABLE long8 as_long8(ulong8 v) {
++ union _type_cast_64_b u;
++ u._ulong8 = v;
++ return u._long8;
++}
++
++INLINE OVERLOADABLE int16 as_int16(ulong8 v) {
++ union _type_cast_64_b u;
++ u._ulong8 = v;
++ return u._int16;
++}
++
++INLINE OVERLOADABLE uint16 as_uint16(ulong8 v) {
++ union _type_cast_64_b u;
++ u._ulong8 = v;
++ return u._uint16;
++}
++
++INLINE OVERLOADABLE float16 as_float16(ulong8 v) {
++ union _type_cast_64_b u;
++ u._ulong8 = v;
++ return u._float16;
++}
++
++INLINE OVERLOADABLE long8 as_long8(int16 v) {
++ union _type_cast_64_b u;
++ u._int16 = v;
++ return u._long8;
++}
++
++INLINE OVERLOADABLE ulong8 as_ulong8(int16 v) {
++ union _type_cast_64_b u;
++ u._int16 = v;
++ return u._ulong8;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long8 as_long8(uint16 v) {
++ union _type_cast_64_b u;
++ u._uint16 = v;
++ return u._long8;
++}
++
++INLINE OVERLOADABLE ulong8 as_ulong8(uint16 v) {
++ union _type_cast_64_b u;
++ u._uint16 = v;
++ return u._ulong8;
++}
++
+ 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;
+ }
+
++INLINE OVERLOADABLE long8 as_long8(float16 v) {
++ union _type_cast_64_b u;
++ u._float16 = v;
++ return u._long8;
++}
++
++INLINE OVERLOADABLE ulong8 as_ulong8(float16 v) {
++ union _type_cast_64_b u;
++ u._float16 = v;
++ return u._ulong8;
++}
++
+ 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;
+ }
+
++union _type_cast_128_b {
++ long16 _long16;
++ ulong16 _ulong16;
++};
++
++INLINE OVERLOADABLE ulong16 as_ulong16(long16 v) {
++ union _type_cast_128_b u;
++ u._long16 = v;
++ return u._ulong16;
++}
++
++INLINE OVERLOADABLE long16 as_long16(ulong16 v) {
++ union _type_cast_128_b u;
++ u._ulong16 = v;
++ return u._long16;
++}
++
++ // ##END_AS##
++
++ // ##BEGIN_CONVERT##
++INLINE OVERLOADABLE ulong2 convert_ulong2(long2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(long2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE uint2 convert_uint2(long2 v) {
++ return (uint2)((uint)(v.s0), (uint)(v.s1));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(long2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(long2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(long2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(long2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(long2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE long2 convert_long2(ulong2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE int2 convert_int2(ulong2 v) {
++ return (int2)((int)(v.s0), (int)(v.s1));
++}
++
++INLINE OVERLOADABLE uint2 convert_uint2(ulong2 v) {
++ return (uint2)((uint)(v.s0), (uint)(v.s1));
++}
++
++INLINE OVERLOADABLE short2 convert_short2(ulong2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(ulong2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(ulong2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(ulong2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE float2 convert_float2(ulong2 v) {
++ return (float2)((float)(v.s0), (float)(v.s1));
++}
++
++INLINE OVERLOADABLE long2 convert_long2(int2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong2 convert_ulong2(int2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long2 convert_long2(uint2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong2 convert_ulong2(uint2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long2 convert_long2(short2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong2 convert_ulong2(short2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long2 convert_long2(ushort2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong2 convert_ulong2(ushort2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long2 convert_long2(char2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong2 convert_ulong2(char2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long2 convert_long2(uchar2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong2 convert_ulong2(uchar2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long2 convert_long2(float2 v) {
++ return (long2)((long)(v.s0), (long)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong2 convert_ulong2(float2 v) {
++ return (ulong2)((ulong)(v.s0), (ulong)(v.s1));
++}
++
+ 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));
+ }
+
+-INLINE OVERLOADABLE short2 convert_short2(float2 v) {
+- return (short2)((short)(v.s0), (short)(v.s1));
++INLINE OVERLOADABLE short2 convert_short2(float2 v) {
++ return (short2)((short)(v.s0), (short)(v.s1));
++}
++
++INLINE OVERLOADABLE ushort2 convert_ushort2(float2 v) {
++ return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++}
++
++INLINE OVERLOADABLE char2 convert_char2(float2 v) {
++ return (char2)((char)(v.s0), (char)(v.s1));
++}
++
++INLINE OVERLOADABLE uchar2 convert_uchar2(float2 v) {
++ return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++}
++
++INLINE OVERLOADABLE ulong3 convert_ulong3(long3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(long3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(long3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(long3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(long3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(long3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(long3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
++}
++
++INLINE OVERLOADABLE float3 convert_float3(long3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
++}
++
++INLINE OVERLOADABLE long3 convert_long3(ulong3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
++}
++
++INLINE OVERLOADABLE int3 convert_int3(ulong3 v) {
++ return (int3)((int)(v.s0), (int)(v.s1), (int)(v.s2));
++}
++
++INLINE OVERLOADABLE uint3 convert_uint3(ulong3 v) {
++ return (uint3)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2));
++}
++
++INLINE OVERLOADABLE short3 convert_short3(ulong3 v) {
++ return (short3)((short)(v.s0), (short)(v.s1), (short)(v.s2));
++}
++
++INLINE OVERLOADABLE ushort3 convert_ushort3(ulong3 v) {
++ return (ushort3)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2));
++}
++
++INLINE OVERLOADABLE char3 convert_char3(ulong3 v) {
++ return (char3)((char)(v.s0), (char)(v.s1), (char)(v.s2));
++}
++
++INLINE OVERLOADABLE uchar3 convert_uchar3(ulong3 v) {
++ return (uchar3)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2));
+ }
+
+-INLINE OVERLOADABLE ushort2 convert_ushort2(float2 v) {
+- return (ushort2)((ushort)(v.s0), (ushort)(v.s1));
++INLINE OVERLOADABLE float3 convert_float3(ulong3 v) {
++ return (float3)((float)(v.s0), (float)(v.s1), (float)(v.s2));
+ }
+
+-INLINE OVERLOADABLE char2 convert_char2(float2 v) {
+- return (char2)((char)(v.s0), (char)(v.s1));
++INLINE OVERLOADABLE long3 convert_long3(int3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
+ }
+
+-INLINE OVERLOADABLE uchar2 convert_uchar2(float2 v) {
+- return (uchar2)((uchar)(v.s0), (uchar)(v.s1));
++INLINE OVERLOADABLE ulong3 convert_ulong3(int3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
+ }
+
+ 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));
+ }
+
++INLINE OVERLOADABLE long3 convert_long3(uint3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
++}
++
++INLINE OVERLOADABLE ulong3 convert_ulong3(uint3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long3 convert_long3(short3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
++}
++
++INLINE OVERLOADABLE ulong3 convert_ulong3(short3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long3 convert_long3(ushort3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
++}
++
++INLINE OVERLOADABLE ulong3 convert_ulong3(ushort3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long3 convert_long3(char3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
++}
++
++INLINE OVERLOADABLE ulong3 convert_ulong3(char3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long3 convert_long3(uchar3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
++}
++
++INLINE OVERLOADABLE ulong3 convert_ulong3(uchar3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long3 convert_long3(float3 v) {
++ return (long3)((long)(v.s0), (long)(v.s1), (long)(v.s2));
++}
++
++INLINE OVERLOADABLE ulong3 convert_ulong3(float3 v) {
++ return (ulong3)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE ulong4 convert_ulong4(long4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(long4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(long4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(long4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(long4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(long4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(long4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(long4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE long4 convert_long4(ulong4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE int4 convert_int4(ulong4 v) {
++ return (int4)((int)(v.s0), (int)(v.s1), (int)(v.s2), (int)(v.s3));
++}
++
++INLINE OVERLOADABLE uint4 convert_uint4(ulong4 v) {
++ return (uint4)((uint)(v.s0), (uint)(v.s1), (uint)(v.s2), (uint)(v.s3));
++}
++
++INLINE OVERLOADABLE short4 convert_short4(ulong4 v) {
++ return (short4)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3));
++}
++
++INLINE OVERLOADABLE ushort4 convert_ushort4(ulong4 v) {
++ return (ushort4)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3));
++}
++
++INLINE OVERLOADABLE char4 convert_char4(ulong4 v) {
++ return (char4)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3));
++}
++
++INLINE OVERLOADABLE uchar4 convert_uchar4(ulong4 v) {
++ return (uchar4)((uchar)(v.s0), (uchar)(v.s1), (uchar)(v.s2), (uchar)(v.s3));
++}
++
++INLINE OVERLOADABLE float4 convert_float4(ulong4 v) {
++ return (float4)((float)(v.s0), (float)(v.s1), (float)(v.s2), (float)(v.s3));
++}
++
++INLINE OVERLOADABLE long4 convert_long4(int4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE ulong4 convert_ulong4(int4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long4 convert_long4(uint4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE ulong4 convert_ulong4(uint4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long4 convert_long4(short4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE ulong4 convert_ulong4(short4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long4 convert_long4(ushort4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE ulong4 convert_ulong4(ushort4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long4 convert_long4(char4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE ulong4 convert_ulong4(char4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long4 convert_long4(uchar4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE ulong4 convert_ulong4(uchar4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long4 convert_long4(float4 v) {
++ return (long4)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3));
++}
++
++INLINE OVERLOADABLE ulong4 convert_ulong4(float4 v) {
++ return (ulong4)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE ulong8 convert_ulong8(long8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
++INLINE OVERLOADABLE int8 convert_int8(long8 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));
++}
++
++INLINE OVERLOADABLE uint8 convert_uint8(long8 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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(long8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(long8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(long8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++INLINE OVERLOADABLE uchar8 convert_uchar8(long8 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));
++}
++
++INLINE OVERLOADABLE float8 convert_float8(long8 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));
++}
++
++INLINE OVERLOADABLE long8 convert_long8(ulong8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE int8 convert_int8(ulong8 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));
++}
++
++INLINE OVERLOADABLE uint8 convert_uint8(ulong8 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));
++}
++
++INLINE OVERLOADABLE short8 convert_short8(ulong8 v) {
++ return (short8)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7));
++}
++
++INLINE OVERLOADABLE ushort8 convert_ushort8(ulong8 v) {
++ return (ushort8)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7));
++}
++
++INLINE OVERLOADABLE char8 convert_char8(ulong8 v) {
++ return (char8)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7));
++}
++
++INLINE OVERLOADABLE uchar8 convert_uchar8(ulong8 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));
++}
++
++INLINE OVERLOADABLE float8 convert_float8(ulong8 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));
++}
++
++INLINE OVERLOADABLE long8 convert_long8(int8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE ulong8 convert_ulong8(int8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long8 convert_long8(uint8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE ulong8 convert_ulong8(uint8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long8 convert_long8(short8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE ulong8 convert_ulong8(short8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long8 convert_long8(ushort8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE ulong8 convert_ulong8(ushort8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long8 convert_long8(char8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE ulong8 convert_ulong8(char8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long8 convert_long8(uchar8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE ulong8 convert_ulong8(uchar8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long8 convert_long8(float8 v) {
++ return (long8)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7));
++}
++
++INLINE OVERLOADABLE ulong8 convert_ulong8(float8 v) {
++ return (ulong8)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE ulong16 convert_ulong16(long16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
++INLINE OVERLOADABLE int16 convert_int16(long16 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));
++}
++
++INLINE OVERLOADABLE uint16 convert_uint16(long16 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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(long16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(long16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(long16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(long16 v) {
++ 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));
++}
++
++INLINE OVERLOADABLE float16 convert_float16(long16 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));
++}
++
++INLINE OVERLOADABLE long16 convert_long16(ulong16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE int16 convert_int16(ulong16 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));
++}
++
++INLINE OVERLOADABLE uint16 convert_uint16(ulong16 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));
++}
++
++INLINE OVERLOADABLE short16 convert_short16(ulong16 v) {
++ return (short16)((short)(v.s0), (short)(v.s1), (short)(v.s2), (short)(v.s3), (short)(v.s4), (short)(v.s5), (short)(v.s6), (short)(v.s7), (short)(v.s8), (short)(v.s9), (short)(v.sA), (short)(v.sB), (short)(v.sC), (short)(v.sD), (short)(v.sE), (short)(v.sF));
++}
++
++INLINE OVERLOADABLE ushort16 convert_ushort16(ulong16 v) {
++ return (ushort16)((ushort)(v.s0), (ushort)(v.s1), (ushort)(v.s2), (ushort)(v.s3), (ushort)(v.s4), (ushort)(v.s5), (ushort)(v.s6), (ushort)(v.s7), (ushort)(v.s8), (ushort)(v.s9), (ushort)(v.sA), (ushort)(v.sB), (ushort)(v.sC), (ushort)(v.sD), (ushort)(v.sE), (ushort)(v.sF));
++}
++
++INLINE OVERLOADABLE char16 convert_char16(ulong16 v) {
++ return (char16)((char)(v.s0), (char)(v.s1), (char)(v.s2), (char)(v.s3), (char)(v.s4), (char)(v.s5), (char)(v.s6), (char)(v.s7), (char)(v.s8), (char)(v.s9), (char)(v.sA), (char)(v.sB), (char)(v.sC), (char)(v.sD), (char)(v.sE), (char)(v.sF));
++}
++
++INLINE OVERLOADABLE uchar16 convert_uchar16(ulong16 v) {
++ 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));
++}
++
++INLINE OVERLOADABLE float16 convert_float16(ulong16 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));
++}
++
++INLINE OVERLOADABLE long16 convert_long16(int16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE ulong16 convert_ulong16(int16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long16 convert_long16(uint16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE ulong16 convert_ulong16(uint16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long16 convert_long16(short16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE ulong16 convert_ulong16(short16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long16 convert_long16(ushort16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE ulong16 convert_ulong16(ushort16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long16 convert_long16(char16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE ulong16 convert_ulong16(char16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long16 convert_long16(uchar16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE ulong16 convert_ulong16(uchar16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
+ 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));
+ }
+
++INLINE OVERLOADABLE long16 convert_long16(float16 v) {
++ return (long16)((long)(v.s0), (long)(v.s1), (long)(v.s2), (long)(v.s3), (long)(v.s4), (long)(v.s5), (long)(v.s6), (long)(v.s7), (long)(v.s8), (long)(v.s9), (long)(v.sA), (long)(v.sB), (long)(v.sC), (long)(v.sD), (long)(v.sE), (long)(v.sF));
++}
++
++INLINE OVERLOADABLE ulong16 convert_ulong16(float16 v) {
++ return (ulong16)((ulong)(v.s0), (ulong)(v.s1), (ulong)(v.s2), (ulong)(v.s3), (ulong)(v.s4), (ulong)(v.s5), (ulong)(v.s6), (ulong)(v.s7), (ulong)(v.s8), (ulong)(v.s9), (ulong)(v.sA), (ulong)(v.sB), (ulong)(v.sC), (ulong)(v.sD), (ulong)(v.sE), (ulong)(v.sF));
++}
++
+ INLINE OVERLOADABLE int16 convert_int16(float16 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));
+ }
+--
+1.7.10.4
+
diff --cc debian/patches/0004-Add-vector-argument-test-case.patch
index 4a1fbbe,0000000..cf43d98
mode 100644,000000..100644
--- a/debian/patches/0004-Add-vector-argument-test-case.patch
+++ b/debian/patches/0004-Add-vector-argument-test-case.patch
@@@ -1,74 -1,0 +1,74 @@@
- 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>
+---
+ kernels/compiler_function_argument2.cl | 6 ++++++
+ utests/CMakeLists.txt | 1 +
+ utests/compiler_function_argument2.cpp | 26 ++++++++++++++++++++++++++
+ 3 files changed, 33 insertions(+)
+ create mode 100644 kernels/compiler_function_argument2.cl
+ create mode 100644 utests/compiler_function_argument2.cpp
+
+diff --git a/kernels/compiler_function_argument2.cl b/kernels/compiler_function_argument2.cl
+new file mode 100644
+index 0000000..0985dbd
+--- /dev/null
++++ b/kernels/compiler_function_argument2.cl
+@@ -0,0 +1,6 @@
++__kernel void
++compiler_function_argument2(__global int *dst, int4 value)
++{
++ int id = (int)get_global_id(0);
++ dst[id] = value.w;
++}
+diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
- index 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
++ compiler_function_argument2.cpp
+ compiler_function_argument.cpp
+ compiler_function_constant0.cpp
+ compiler_function_constant1.cpp
+diff --git a/utests/compiler_function_argument2.cpp b/utests/compiler_function_argument2.cpp
+new file mode 100644
+index 0000000..1e398a9
+--- /dev/null
++++ b/utests/compiler_function_argument2.cpp
+@@ -0,0 +1,26 @@
++#include "utest_helper.hpp"
++
++struct int4 {int x,y,z,w;};
++void compiler_function_argument2(void)
++{
++ const size_t n = 2048;
++ const int4 value = {31, 32, 33, 34};
++
++ // Setup kernel and buffers
++ OCL_CREATE_KERNEL("compiler_function_argument2");
++ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
++ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
++ OCL_SET_ARG(1, sizeof(int4), &value);
++
++ // Run the kernel
++ globals[0] = n;
++ locals[0] = 16;
++ OCL_NDRANGE(1);
++ OCL_MAP_BUFFER(0);
++
++ // Check results
++ for (uint32_t i = 0; i < n; ++i)
++ OCL_ASSERT(((int*)buf_data[0])[i] == value.w);
++}
++
++MAKE_UTEST_FROM_FUNCTION(compiler_function_argument2);
+--
+1.7.10.4
+
diff --cc debian/patches/series
index f33ca92,0000000..9971719
mode 100644,000000..100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@@ -1,11 -1,0 +1,8 @@@
+debug
+flags
+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