gdb:

	* NEWS: Mention OpenCL C language support.
	* Makefile.in (SFILES): Add opencl-lang.c.
	(COMMON_OBS): Add opencl-lang.o.
	* opencl-lang.c: New File
	* defs.h (enum language): Add language_opencl.
	* dwarf2read.c (read_file_scope): Handle DW_AT_producer for the
	IBM XL C OpenCL compiler.
	* c-lang.h: Include "parser-defs.h".
	(evaluate_subexp_c): Declare.
	* c-lang.c (evaluate_subexp_c): Remove the static qualifier.
	(c_op_print_tab): Add declaration.
	* eval.c (binop_promote): Handle language_opencl.
	* c-exp.y: Lookup the primitive types instead of referring to the
	builtins.

gdb/testsuite:
	* Makefile.in (ALL_SUBDIRS): Add gdb.opencl.
	* configure.ac (AC_OUTPUT): Add gdb.opencl/Makefile.
	* configure: Regenerate.
	* gdb.opencl/Makefile.in: New File.
	* gdb.opencl/datatypes.exp: Likewise.
	* gdb.opencl/datatypes.cl: Likewise.
	* gdb.opencl/operators.exp: Likewise.
	* gdb.opencl/operators.cl: Likewise.
	* gdb.opencl/vec_comps.exp: Likewise.
	* gdb.opencl/vec_comps.cl: Likewise.
	* gdb.opencl/convs_casts.exp: Likewise.
	* gdb.opencl/convs_casts.cl: Likewise.
	* lib/opencl.exp: Likewise.
	* lib/opencl_hostapp.c: Likewise.
	* lib/opencl_kernel.cl: Likewise.
	* lib/cl_util.c: Likewise.
	* lib/cl_util.c: Likewise.
	* gdb.base/default.exp (set language): Add "opencl" to the list of
	languages.

gdb/doc:
	* gdb.texinfo: (Summary) Add mention about OpenCL C language support.
	(OpenCL C): New node.

31 files changed, 4550 insertions, 43 deletions

diff --git a/gdb/ChangeLog b/gdb/ChangeLog
index 9ea16e95ce9..793c0626f40 100644
--- a/gdb/ChangeLog
+++ b/gdb/ChangeLog
@@ -1,3 +1,20 @@
+2010-11-05  Ken Werner  <ken.werner@de.ibm.com>
+
+	* NEWS: Mention OpenCL C language support.
+	* Makefile.in (SFILES): Add opencl-lang.c.
+	(COMMON_OBS): Add opencl-lang.o.
+	* opencl-lang.c: New File
+	* defs.h (enum language): Add language_opencl.
+	* dwarf2read.c (read_file_scope): Handle DW_AT_producer for the
+	IBM XL C OpenCL compiler.
+	* c-lang.h: Include "parser-defs.h".
+	(evaluate_subexp_c): Declare.
+	* c-lang.c (evaluate_subexp_c): Remove the static qualifier.
+	(c_op_print_tab): Add declaration.
+	* eval.c (binop_promote): Handle language_opencl.
+	* c-exp.y: Lookup the primitive types instead of referring to the
+	builtins.
+
 2010-11-05  Jan Kratochvil  <jan.kratochvil@redhat.com>
 
 	Fix configure --enable-plugins --without-python.
diff --git a/gdb/Makefile.in b/gdb/Makefile.in
index 568fdb5eb8e..550badfdeb5 100644
--- a/gdb/Makefile.in
+++ b/gdb/Makefile.in
@@ -689,6 +689,7 @@ SFILES = ada-exp.y ada-lang.c ada-typeprint.c ada-valprint.c ada-tasks.c \
 	mi/mi-common.c \
 	objc-exp.y objc-lang.c \
 	objfiles.c osabi.c observer.c osdata.c \
+	opencl-lang.c \
 	p-exp.y p-lang.c p-typeprint.c p-valprint.c parse.c printcmd.c \
 	proc-service.list progspace.c \
 	prologue-value.c psymtab.c \
@@ -845,7 +846,7 @@ COMMON_OBS = $(DEPFILES) $(CONFIG_OBS) $(YYOBJ) \
 	ui-out.o cli-out.o \
 	varobj.o vec.o wrapper.o \
 	jv-lang.o jv-valprint.o jv-typeprint.o \
-	m2-lang.o p-lang.o p-typeprint.o p-valprint.o \
+	m2-lang.o opencl-lang.o p-lang.o p-typeprint.o p-valprint.o \
 	sentinel-frame.o \
 	complaints.o typeprint.o \
 	ada-typeprint.o c-typeprint.o f-typeprint.o m2-typeprint.o \
diff --git a/gdb/NEWS b/gdb/NEWS
index a71461cbf84..e883a8bd6e0 100644
--- a/gdb/NEWS
+++ b/gdb/NEWS
@@ -3,6 +3,10 @@
 
 *** Changes since GDB 7.2
 
+* OpenCL C
+  Initial support for the OpenCL C language (http://www.khronos.org/opencl)
+  has been integrated into GDB.
+
 * Python scripting
 
   ** GDB values in Python are now callable if the value represents a
diff --git a/gdb/c-exp.y b/gdb/c-exp.y
index 57e09b33b23..2e6c3712f50 100644
--- a/gdb/c-exp.y
+++ b/gdb/c-exp.y
@@ -612,7 +612,9 @@ exp	:	VARIABLE
 
 exp	:	SIZEOF '(' type ')'	%prec UNARY
 			{ write_exp_elt_opcode (OP_LONG);
-			  write_exp_elt_type (parse_type->builtin_int);
+			  write_exp_elt_type (lookup_signed_typename
+					      (parse_language, parse_gdbarch,
+					       "int"));
 			  CHECK_TYPEDEF ($3);
 			  write_exp_elt_longcst ((LONGEST) TYPE_LENGTH ($3));
 			  write_exp_elt_opcode (OP_LONG); }
@@ -980,61 +982,117 @@ typebase  /* Implements (approximately): (type-qualifier)* type-specifier */
 	:	TYPENAME
 			{ $$ = $1.type; }
 	|	INT_KEYWORD
-			{ $$ = parse_type->builtin_int; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "int"); }
 	|	LONG
-			{ $$ = parse_type->builtin_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long"); }
 	|	SHORT
-			{ $$ = parse_type->builtin_short; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "short"); }
 	|	LONG INT_KEYWORD
-			{ $$ = parse_type->builtin_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long"); }
 	|	LONG SIGNED_KEYWORD INT_KEYWORD
-			{ $$ = parse_type->builtin_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long"); }
 	|	LONG SIGNED_KEYWORD
-			{ $$ = parse_type->builtin_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long"); }
 	|	SIGNED_KEYWORD LONG INT_KEYWORD
-			{ $$ = parse_type->builtin_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long"); }
 	|	UNSIGNED LONG INT_KEYWORD
-			{ $$ = parse_type->builtin_unsigned_long; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "long"); }
 	|	LONG UNSIGNED INT_KEYWORD
-			{ $$ = parse_type->builtin_unsigned_long; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "long"); }
 	|	LONG UNSIGNED
-			{ $$ = parse_type->builtin_unsigned_long; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "long"); }
 	|	LONG LONG
-			{ $$ = parse_type->builtin_long_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long long"); }
 	|	LONG LONG INT_KEYWORD
-			{ $$ = parse_type->builtin_long_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long long"); }
 	|	LONG LONG SIGNED_KEYWORD INT_KEYWORD
-			{ $$ = parse_type->builtin_long_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long long"); }
 	|	LONG LONG SIGNED_KEYWORD
-			{ $$ = parse_type->builtin_long_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long long"); }
 	|	SIGNED_KEYWORD LONG LONG
-			{ $$ = parse_type->builtin_long_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long long"); }
 	|	SIGNED_KEYWORD LONG LONG INT_KEYWORD
-			{ $$ = parse_type->builtin_long_long; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "long long"); }
 	|	UNSIGNED LONG LONG
-			{ $$ = parse_type->builtin_unsigned_long_long; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "long long"); }
 	|	UNSIGNED LONG LONG INT_KEYWORD
-			{ $$ = parse_type->builtin_unsigned_long_long; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "long long"); }
 	|	LONG LONG UNSIGNED
-			{ $$ = parse_type->builtin_unsigned_long_long; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "long long"); }
 	|	LONG LONG UNSIGNED INT_KEYWORD
-			{ $$ = parse_type->builtin_unsigned_long_long; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "long long"); }
 	|	SHORT INT_KEYWORD
-			{ $$ = parse_type->builtin_short; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "short"); }
 	|	SHORT SIGNED_KEYWORD INT_KEYWORD
-			{ $$ = parse_type->builtin_short; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "short"); }
 	|	SHORT SIGNED_KEYWORD
-			{ $$ = parse_type->builtin_short; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "short"); }
 	|	UNSIGNED SHORT INT_KEYWORD
-			{ $$ = parse_type->builtin_unsigned_short; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "short"); }
 	|	SHORT UNSIGNED 
-			{ $$ = parse_type->builtin_unsigned_short; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "short"); }
 	|	SHORT UNSIGNED INT_KEYWORD
-			{ $$ = parse_type->builtin_unsigned_short; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "short"); }
 	|	DOUBLE_KEYWORD
-			{ $$ = parse_type->builtin_double; }
+			{ $$ = lookup_typename (parse_language, parse_gdbarch,
+						"double", (struct block *) NULL,
+						0); }
 	|	LONG DOUBLE_KEYWORD
-			{ $$ = parse_type->builtin_long_double; }
+			{ $$ = lookup_typename (parse_language, parse_gdbarch,
+						"long double",
+						(struct block *) NULL, 0); }
 	|	STRUCT name
 			{ $$ = lookup_struct (copy_name ($2),
 					      expression_context_block); }
@@ -1052,13 +1110,17 @@ typebase  /* Implements (approximately): (type-qualifier)* type-specifier */
 							 parse_gdbarch,
 							 TYPE_NAME($2.type)); }
 	|	UNSIGNED
-			{ $$ = parse_type->builtin_unsigned_int; }
+			{ $$ = lookup_unsigned_typename (parse_language,
+							 parse_gdbarch,
+							 "int"); }
 	|	SIGNED_KEYWORD typename
 			{ $$ = lookup_signed_typename (parse_language,
 						       parse_gdbarch,
 						       TYPE_NAME($2.type)); }
 	|	SIGNED_KEYWORD
-			{ $$ = parse_type->builtin_int; }
+			{ $$ = lookup_signed_typename (parse_language,
+						       parse_gdbarch,
+						       "int"); }
                 /* It appears that this rule for templates is never
                    reduced; template recognition happens by lookahead
                    in the token processing code in yylex. */         
@@ -1077,19 +1139,25 @@ typename:	TYPENAME
 		{
 		  $$.stoken.ptr = "int";
 		  $$.stoken.length = 3;
-		  $$.type = parse_type->builtin_int;
+		  $$.type = lookup_signed_typename (parse_language,
+						    parse_gdbarch,
+						    "int");
 		}
 	|	LONG
 		{
 		  $$.stoken.ptr = "long";
 		  $$.stoken.length = 4;
-		  $$.type = parse_type->builtin_long;
+		  $$.type = lookup_signed_typename (parse_language,
+						    parse_gdbarch,
+						    "long");
 		}
 	|	SHORT
 		{
 		  $$.stoken.ptr = "short";
 		  $$.stoken.length = 5;
-		  $$.type = parse_type->builtin_short;
+		  $$.type = lookup_signed_typename (parse_language,
+						    parse_gdbarch,
+						    "short");
 		}
 	;
 
diff --git a/gdb/c-lang.c b/gdb/c-lang.c
index 40c417256bf..015ba163089 100644
--- a/gdb/c-lang.c
+++ b/gdb/c-lang.c
@@ -933,7 +933,7 @@ parse_one_string (struct obstack *output, char *data, int len,
    are delegated to evaluate_subexp_standard; see that function for a
    description of the arguments.  */
 
-static struct value *
+struct value *
 evaluate_subexp_c (struct type *expect_type, struct expression *exp,
 		   int *pos, enum noside noside)
 {
diff --git a/gdb/c-lang.h b/gdb/c-lang.h
index a04fbb29438..dc571a41083 100644
--- a/gdb/c-lang.h
+++ b/gdb/c-lang.h
@@ -27,6 +27,7 @@ struct language_arch_info;
 
 #include "value.h"
 #include "macroexp.h"
+#include "parser-defs.h"
 
 
 /* The various kinds of C string and character.  Note that these
@@ -78,6 +79,10 @@ extern int c_value_print (struct value *, struct ui_file *,
 
 /* These are in c-lang.c: */
 
+extern struct value *evaluate_subexp_c (struct type *expect_type,
+					 struct expression *exp, int *pos,
+					 enum noside noside);
+
 extern void c_printchar (int, struct type *, struct ui_file *);
 
 extern void c_printstr (struct ui_file * stream, struct type *elttype,
@@ -93,6 +98,8 @@ extern const struct exp_descriptor exp_descriptor_c;
 extern void c_emit_char (int c, struct type *type,
 			 struct ui_file *stream, int quoter);
 
+extern const struct op_print c_op_print_tab[];
+
 /* These are in c-typeprint.c: */
 
 extern void c_type_print_base (struct type *, struct ui_file *, int, int);
diff --git a/gdb/defs.h b/gdb/defs.h
index 6cb0a1925b2..489de74e753 100644
--- a/gdb/defs.h
+++ b/gdb/defs.h
@@ -201,6 +201,7 @@ enum language
     language_asm,		/* Assembly language */
     language_pascal,		/* Pascal */
     language_ada,		/* Ada */
+    language_opencl,		/* OpenCL */
     language_minimal,		/* All other languages, minimal support only */
     nr_languages
   };
diff --git a/gdb/doc/ChangeLog b/gdb/doc/ChangeLog
index a8df0726966..63f97eced5c 100644
--- a/gdb/doc/ChangeLog
+++ b/gdb/doc/ChangeLog
@@ -1,3 +1,8 @@
+2010-11-05  Ken Werner  <ken.werner@de.ibm.com>
+
+	* gdb.texinfo: (Summary) Add mention about OpenCL C language support.
+	(OpenCL C): New node.
+
 2010-11-02  Doug Evans  <dje@google.com>
 
 	* gdb.texinfo (Pretty Printing): Expand into three sections,
diff --git a/gdb/doc/gdb.texinfo b/gdb/doc/gdb.texinfo
index 069dce4bb4f..993b0fbaa7e 100644
--- a/gdb/doc/gdb.texinfo
+++ b/gdb/doc/gdb.texinfo
@@ -221,6 +221,9 @@ Support for D is partial.  For information on D, see
 Support for Modula-2 is partial.  For information on Modula-2, see
 @ref{Modula-2,,Modula-2}.
 
+Support for OpenCL C is partial.  For information on OpenCL C, see
+@ref{OpenCL C,,OpenCL C}.
+
 @cindex Pascal
 Debugging Pascal programs which use sets, subranges, file variables, or
 nested functions does not currently work.  @value{GDBN} does not support
@@ -11611,7 +11614,7 @@ being set automatically by @value{GDBN}.
 @node Supported Languages
 @section Supported Languages
 
-@value{GDBN} supports C, C@t{++}, D, Objective-C, Fortran, Java, Pascal,
+@value{GDBN} supports C, C@t{++}, D, Objective-C, Fortran, Java, OpenCL C, Pascal,
 assembly, Modula-2, and Ada.
 @c This is false ...
 Some @value{GDBN} features may be used in expressions regardless of the
@@ -11632,6 +11635,7 @@ language reference or tutorial.
 * C::                           C and C@t{++}
 * D::                           D
 * Objective-C::                 Objective-C
+* OpenCL C::                    OpenCL C
 * Fortran::                     Fortran
 * Pascal::                      Pascal
 * Modula-2::                    Modula-2
@@ -12278,6 +12282,42 @@ the description of an object.  However, this command may only work
 with certain Objective-C libraries that have a particular hook
 function, @code{_NSPrintForDebugger}, defined.
 
+@node OpenCL C
+@subsection OpenCL C
+
+@cindex OpenCL C
+This section provides information about @value{GDBN}s OpenCL C support.
+
+@menu
+* OpenCL C Datatypes::
+* OpenCL C Expressions::
+* OpenCL C Operators::
+@end menu
+
+@node OpenCL C Datatypes
+@subsubsection OpenCL C Datatypes
+
+@cindex OpenCL C Datatypes
+@value{GDBN} supports the builtin scalar and vector datatypes specified
+by OpenCL 1.1.  In addition the half- and double-precision floating point
+data types of the @code{cl_khr_fp16} and @code{cl_khr_fp64} OpenCL
+extensions are also known to @value{GDBN}.
+
+@node OpenCL C Expressions
+@subsubsection OpenCL C Expressions
+
+@cindex OpenCL C Expressions
+@value{GDBN} supports accesses to vector components including the access as
+lvalue where possible.  Since OpenCL C is based on C99 most C expressions
+supported by @value{GDBN} can be used as well.
+
+@node OpenCL C Operators
+@subsubsection OpenCL C Operators
+
+@cindex OpenCL C Operators
+@value{GDBN} supports the operators specified by OpenCL 1.1 for scalar and
+vector data types.
+
 @node Fortran
 @subsection Fortran
 @cindex Fortran-specific support in @value{GDBN}
diff --git a/gdb/dwarf2read.c b/gdb/dwarf2read.c
index a91f14a4c93..404faf89e65 100644
--- a/gdb/dwarf2read.c
+++ b/gdb/dwarf2read.c
@@ -5089,6 +5089,12 @@ read_file_scope (struct die_info *die, struct dwarf2_cu *cu)
   if (attr)
     cu->producer = DW_STRING (attr);
 
+  /* The XLCL doesn't generate DW_LANG_OpenCL because this attribute is not
+     standardised yet.  As a workaround for the language detection we fall
+     back to the DW_AT_producer string.  */
+  if (cu->producer && strstr (cu->producer, "IBM XL C for OpenCL") != NULL)
+    cu->language = language_opencl;
+
   /* We assume that we're processing GCC output. */
   processing_gcc_compilation = 2;
 
diff --git a/gdb/eval.c b/gdb/eval.c
index 71c3ff8cae7..16e98aea5fa 100644
--- a/gdb/eval.c
+++ b/gdb/eval.c
@@ -603,6 +603,7 @@ binop_promote (const struct language_defn *language, struct gdbarch *gdbarch,
 	case language_cplus:
 	case language_asm:
 	case language_objc:
+	case language_opencl:
 	  /* No promotion required.  */
 	  break;
 
@@ -690,7 +691,24 @@ binop_promote (const struct language_defn *language, struct gdbarch *gdbarch,
 			       : builtin->builtin_long_long);
 	    }
 	  break;
-
+	case language_opencl:
+	  if (result_len <= TYPE_LENGTH (lookup_signed_typename
+					 (language, gdbarch, "int")))
+	    {
+	      promoted_type =
+		(unsigned_operation
+		 ? lookup_unsigned_typename (language, gdbarch, "int")
+		 : lookup_signed_typename (language, gdbarch, "int"));
+	    }
+	  else if (result_len <= TYPE_LENGTH (lookup_signed_typename
+					      (language, gdbarch, "long")))
+	    {
+	      promoted_type =
+		(unsigned_operation
+		 ? lookup_unsigned_typename (language, gdbarch, "long")
+		 : lookup_signed_typename (language, gdbarch,"long"));
+	    }
+	  break;
 	default:
 	  /* For other languages the result type is unchanged from gdb
 	     version 6.7 for backward compatibility.
diff --git a/gdb/opencl-lang.c b/gdb/opencl-lang.c
new file mode 100644
index 00000000000..088d49af32c
--- /dev/null
+++ b/gdb/opencl-lang.c
@@ -0,0 +1,1162 @@
+/* OpenCL language support for GDB, the GNU debugger.
+   Copyright (C) 2010 Free Software Foundation, Inc.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>.
+
+   This file is part of GDB.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+
+#include "defs.h"
+#include "gdb_string.h"
+#include "gdbtypes.h"
+#include "symtab.h"
+#include "expression.h"
+#include "parser-defs.h"
+#include "symtab.h"
+#include "language.h"
+#include "c-lang.h"
+#include "gdb_assert.h"
+
+extern void _initialize_opencl_language (void);
+
+/* This macro generates enum values from a given type.  */
+
+#define OCL_P_TYPE(TYPE)\
+  opencl_primitive_type_##TYPE,\
+  opencl_primitive_type_##TYPE##2,\
+  opencl_primitive_type_##TYPE##3,\
+  opencl_primitive_type_##TYPE##4,\
+  opencl_primitive_type_##TYPE##8,\
+  opencl_primitive_type_##TYPE##16
+
+enum opencl_primitive_types {
+  OCL_P_TYPE (char),
+  OCL_P_TYPE (uchar),
+  OCL_P_TYPE (short),
+  OCL_P_TYPE (ushort),
+  OCL_P_TYPE (int),
+  OCL_P_TYPE (uint),
+  OCL_P_TYPE (long),
+  OCL_P_TYPE (ulong),
+  OCL_P_TYPE (half),
+  OCL_P_TYPE (float),
+  OCL_P_TYPE (double),
+  opencl_primitive_type_bool,
+  opencl_primitive_type_unsigned_char,
+  opencl_primitive_type_unsigned_short,
+  opencl_primitive_type_unsigned_int,
+  opencl_primitive_type_unsigned_long,
+  opencl_primitive_type_size_t,
+  opencl_primitive_type_ptrdiff_t,
+  opencl_primitive_type_intptr_t,
+  opencl_primitive_type_uintptr_t,
+  opencl_primitive_type_void,
+  nr_opencl_primitive_types
+};
+
+/* This macro generates the type struct declarations from a given type.  */
+
+#define STRUCT_OCL_TYPE(TYPE)\
+  struct type *builtin_##TYPE;\
+  struct type *builtin_##TYPE##2;\
+  struct type *builtin_##TYPE##3;\
+  struct type *builtin_##TYPE##4;\
+  struct type *builtin_##TYPE##8;\
+  struct type *builtin_##TYPE##16
+
+struct builtin_opencl_type
+{
+  STRUCT_OCL_TYPE (char);
+  STRUCT_OCL_TYPE (uchar);
+  STRUCT_OCL_TYPE (short);
+  STRUCT_OCL_TYPE (ushort);
+  STRUCT_OCL_TYPE (int);
+  STRUCT_OCL_TYPE (uint);
+  STRUCT_OCL_TYPE (long);
+  STRUCT_OCL_TYPE (ulong);
+  STRUCT_OCL_TYPE (half);
+  STRUCT_OCL_TYPE (float);
+  STRUCT_OCL_TYPE (double);
+  struct type *builtin_bool;
+  struct type *builtin_unsigned_char;
+  struct type *builtin_unsigned_short;
+  struct type *builtin_unsigned_int;
+  struct type *builtin_unsigned_long;
+  struct type *builtin_size_t;
+  struct type *builtin_ptrdiff_t;
+  struct type *builtin_intptr_t;
+  struct type *builtin_uintptr_t;
+  struct type *builtin_void;
+};
+
+static struct gdbarch_data *opencl_type_data;
+
+const struct builtin_opencl_type *
+builtin_opencl_type (struct gdbarch *gdbarch)
+{
+  return gdbarch_data (gdbarch, opencl_type_data);
+}
+
+/* Returns the corresponding OpenCL vector type from the given type code,
+   the length of the element type, the unsigned flag and the amount of
+   elements (N).  */
+
+static struct type *
+lookup_opencl_vector_type (struct gdbarch *gdbarch, enum type_code code,
+			   unsigned int el_length, unsigned int flag_unsigned,
+			   int n)
+{
+  int i;
+  unsigned int length;
+  struct type *type = NULL;
+  struct type **types = (struct type **) builtin_opencl_type (gdbarch);
+
+  /* Check if n describes a valid OpenCL vector size (2, 3, 4, 8, 16).  */
+  if (n != 2 && n != 3 && n != 4 && n != 8 && n != 16)
+    error (_("Invalid OpenCL vector size: %d"), n);
+
+  /* Triple vectors have the size of a quad vector.  */
+  length = (n == 3) ?  el_length * 4 : el_length * n;
+
+  for (i = 0; i < nr_opencl_primitive_types; i++)
+    {
+      LONGEST lowb, highb;
+
+      if (TYPE_CODE (types[i]) == TYPE_CODE_ARRAY && TYPE_VECTOR (types[i])
+	  && get_array_bounds (types[i], &lowb, &highb)
+	  && TYPE_CODE (TYPE_TARGET_TYPE (types[i])) == code
+	  && TYPE_UNSIGNED (TYPE_TARGET_TYPE (types[i])) == flag_unsigned
+	  && TYPE_LENGTH (TYPE_TARGET_TYPE (types[i])) == el_length
+	  && TYPE_LENGTH (types[i]) == length
+	  && highb - lowb + 1 == n)
+	{
+	  type = types[i];
+	  break;
+	}
+    }
+
+  return type;
+}
+
+/* Returns nonzero if the array ARR contains duplicates within
+     the first N elements.  */
+
+static int
+array_has_dups (int *arr, int n)
+{
+  int i, j;
+
+  for (i = 0; i < n; i++)
+    {
+      for (j = i + 1; j < n; j++)
+        {
+          if (arr[i] == arr[j])
+            return 1;
+        }
+    }
+
+  return 0;
+}
+
+/* The OpenCL component access syntax allows to create lvalues referring to
+   selected elements of an original OpenCL vector in arbitrary order.  This
+   structure holds the information to describe such lvalues.  */
+
+struct lval_closure
+{
+  /* Reference count.  */
+  int refc;
+  /* The number of indices.  */
+  int n;
+  /* The element indices themselves.  */
+  int *indices;
+  /* A pointer to the original value.  */
+  struct value *val;
+};
+
+/* Allocates an instance of struct lval_closure.  */
+
+static struct lval_closure *
+allocate_lval_closure (int *indices, int n, struct value *val)
+{
+  struct lval_closure *c = XZALLOC (struct lval_closure);
+
+  c->refc = 1;
+  c->n = n;
+  c->indices = XCALLOC (n, int);
+  memcpy (c->indices, indices, n * sizeof (int));
+  value_incref (val); /* Increment the reference counter of the value.  */
+  c->val = val;
+
+  return c;
+}
+
+static void
+lval_func_read (struct value *v)
+{
+  struct lval_closure *c = (struct lval_closure *) value_computed_closure (v);
+  struct type *type = check_typedef (value_type (v));
+  struct type *eltype = TYPE_TARGET_TYPE (check_typedef (value_type (c->val)));
+  int offset = value_offset (v);
+  int elsize = TYPE_LENGTH (eltype);
+  int n, i, j = 0;
+  LONGEST lowb = 0;
+  LONGEST highb = 0;
+
+  if (TYPE_CODE (type) == TYPE_CODE_ARRAY
+      && !get_array_bounds (type, &lowb, &highb))
+    error (_("Could not determine the vector bounds"));
+
+  /* Assume elsize aligned offset.  */
+  gdb_assert (offset % elsize == 0);
+  offset /= elsize;
+  n = offset + highb - lowb + 1;
+  gdb_assert (n <= c->n);
+
+  for (i = offset; i < n; i++)
+    memcpy (value_contents_raw (v) + j++ * elsize,
+	    value_contents (c->val) + c->indices[i] * elsize,
+	    elsize);
+}
+
+static void
+lval_func_write (struct value *v, struct value *fromval)
+{
+  struct value *mark = value_mark ();
+  struct lval_closure *c = (struct lval_closure *) value_computed_closure (v);
+  struct type *type = check_typedef (value_type (v));
+  struct type *eltype = TYPE_TARGET_TYPE (check_typedef (value_type (c->val)));
+  int offset = value_offset (v);
+  int elsize = TYPE_LENGTH (eltype);
+  int n, i, j = 0;
+  LONGEST lowb = 0;
+  LONGEST highb = 0;
+
+  if (TYPE_CODE (type) == TYPE_CODE_ARRAY
+      && !get_array_bounds (type, &lowb, &highb))
+    error (_("Could not determine the vector bounds"));
+
+  /* Assume elsize aligned offset.  */
+  gdb_assert (offset % elsize == 0);
+  offset /= elsize;
+  n = offset + highb - lowb + 1;
+
+  /* Since accesses to the fourth component of a triple vector is undefined we
+     just skip writes to the fourth element.  Imagine something like this:
+       int3 i3 = (int3)(0, 1, 2);
+       i3.hi.hi = 5;
+     In this case n would be 4 (offset=12/4 + 1) while c->n would be 3.  */
+  if (n > c->n)
+    n = c->n;
+
+  for (i = offset; i < n; i++)
+    {
+      struct value *from_elm_val = allocate_value (eltype);
+      struct value *to_elm_val = value_subscript (c->val, c->indices[i]);
+
+      memcpy (value_contents_writeable (from_elm_val),
+	      value_contents (fromval) + j++ * elsize,
+	      elsize);
+      value_assign (to_elm_val, from_elm_val);
+    }
+
+  value_free_to_mark (mark);
+}
+
+/* Return nonzero if all bits in V within OFFSET and LENGTH are valid.  */
+
+static int
+lval_func_check_validity (const struct value *v, int offset, int length)
+{
+  struct lval_closure *c = (struct lval_closure *) value_computed_closure (v);
+  /* Size of the target type in bits.  */
+  int elsize =
+      TYPE_LENGTH (TYPE_TARGET_TYPE (check_typedef (value_type (c->val)))) * 8;
+  int startrest = offset % elsize;
+  int start = offset / elsize;
+  int endrest = (offset + length) % elsize;
+  int end = (offset + length) / elsize;
+  int i;
+
+  if (endrest)
+    end++;
+
+  if (end > c->n)
+    return 0;
+
+  for (i = start; i < end; i++)
+    {
+      int startoffset = (i == start) ? startrest : 0;
+      int length = (i == end) ? endrest : elsize;
+
+      if (!value_bits_valid (c->val, c->indices[i] * elsize + startoffset,
+			     length))
+	return 0;
+    }
+
+  return 1;
+}
+
+/* Return nonzero if any bit in V is valid.  */
+
+static int
+lval_func_check_any_valid (const struct value *v)
+{
+  struct lval_closure *c = (struct lval_closure *) value_computed_closure (v);
+  /* Size of the target type in bits.  */
+  int elsize =
+      TYPE_LENGTH (TYPE_TARGET_TYPE (check_typedef (value_type (c->val)))) * 8;
+  int i;
+
+  for (i = 0; i < c->n; i++)
+    if (value_bits_valid (c->val, c->indices[i] * elsize, elsize))
+      return 1;
+
+  return 0;
+}
+
+static void *
+lval_func_copy_closure (const struct value *v)
+{
+  struct lval_closure *c = (struct lval_closure *) value_computed_closure (v);
+
+  ++c->refc;
+
+  return c;
+}
+
+static void
+lval_func_free_closure (struct value *v)
+{
+  struct lval_closure *c = (struct lval_closure *) value_computed_closure (v);
+
+  --c->refc;
+
+  if (c->refc == 0)
+    {
+      xfree (c->indices);
+      xfree (c);
+      value_free (c->val); /* Decrement the reference counter of the value.  */
+    }
+}
+
+static struct lval_funcs opencl_value_funcs =
+  {
+    lval_func_read,
+    lval_func_write,
+    lval_func_check_validity,
+    lval_func_check_any_valid,
+    lval_func_copy_closure,
+    lval_func_free_closure
+  };
+
+/* Creates a sub-vector from VAL.  The elements are selected by the indices of
+   an array with the length of N.  Supported values for NOSIDE are
+   EVAL_NORMAL and EVAL_AVOID_SIDE_EFFECTS.  */
+
+static struct value *
+create_value (struct gdbarch *gdbarch, struct value *val, enum noside noside,
+	      int *indices, int n)
+{
+  struct type *type = check_typedef (value_type (val));
+  struct type *elm_type = TYPE_TARGET_TYPE (type);
+  struct value *ret;
+
+  /* Check if a single component of a vector is requested which means
+     the resulting type is a (primitive) scalar type.  */
+  if (n == 1)
+    {
+      if (noside == EVAL_AVOID_SIDE_EFFECTS)
+        ret = value_zero (elm_type, not_lval);
+      else
+        ret = value_subscript (val, indices[0]);
+    }
+  else
+    {
+      /* Multiple components of the vector are requested which means the
+	 resulting type is a vector as well.  */
+      struct type *dst_type =
+	lookup_opencl_vector_type (gdbarch, TYPE_CODE (elm_type),
+				   TYPE_LENGTH (elm_type),
+				   TYPE_UNSIGNED (elm_type), n);
+
+      if (dst_type == NULL)
+	dst_type = init_vector_type (elm_type, n);
+
+      make_cv_type (TYPE_CONST (type), TYPE_VOLATILE (type), dst_type, NULL);
+
+      if (noside == EVAL_AVOID_SIDE_EFFECTS)
+	ret = allocate_value (dst_type);
+      else
+	{
+	  /* Check whether to create a lvalue or not.  */
+	  if (VALUE_LVAL (val) != not_lval && !array_has_dups (indices, n))
+	    {
+	      struct lval_closure *c = allocate_lval_closure (indices, n, val);
+	      ret = allocate_computed_value (dst_type, &opencl_value_funcs, c);
+	    }
+	  else
+	    {
+	      int i;
+
+	      ret = allocate_value (dst_type);
+
+	      /* Copy src val contents into the destination value.  */
+	      for (i = 0; i < n; i++)
+		memcpy (value_contents_writeable (ret)
+			+ (i * TYPE_LENGTH (elm_type)),
+			value_contents (val)
+			+ (indices[i] * TYPE_LENGTH (elm_type)),
+			TYPE_LENGTH (elm_type));
+	    }
+	}
+    }
+  return ret;
+}
+
+/* OpenCL vector component access.  */
+
+static struct value *
+opencl_component_ref (struct expression *exp, struct value *val, char *comps,
+		      enum noside noside)
+{
+  LONGEST lowb, highb;
+  int src_len;
+  struct value *v;
+  int indices[16], i;
+  int dst_len;
+
+  if (!get_array_bounds (check_typedef (value_type (val)), &lowb, &highb))
+    error (_("Could not determine the vector bounds"));
+
+  src_len = highb - lowb + 1;
+
+  /* Throw an error if the amount of array elements does not fit a
+     valid OpenCL vector size (2, 3, 4, 8, 16).  */
+  if (src_len != 2 && src_len != 3 && src_len != 4 && src_len != 8
+      && src_len != 16)
+    error (_("Invalid OpenCL vector size"));
+
+  if (strcmp (comps, "lo") == 0 )
+    {
+      dst_len = (src_len == 3) ? 2 : src_len / 2;
+
+      for (i = 0; i < dst_len; i++)
+	indices[i] = i;
+    }
+  else if (strcmp (comps, "hi") == 0)
+    {
+      dst_len = (src_len == 3) ? 2 : src_len / 2;
+
+      for (i = 0; i < dst_len; i++)
+	indices[i] = dst_len + i;
+    }
+  else if (strcmp (comps, "even") == 0)
+    {
+      dst_len = (src_len == 3) ? 2 : src_len / 2;
+
+      for (i = 0; i < dst_len; i++)
+	indices[i] = i*2;
+    }
+  else if (strcmp (comps, "odd") == 0)
+    {
+      dst_len = (src_len == 3) ? 2 : src_len / 2;
+
+      for (i = 0; i < dst_len; i++)
+        indices[i] = i*2+1;
+    }
+  else if (strncasecmp (comps, "s", 1) == 0)
+    {
+#define HEXCHAR_TO_INT(C) ((C >= '0' && C <= '9') ? \
+                           C-'0' : ((C >= 'A' && C <= 'F') ? \
+                           C-'A'+10 : ((C >= 'a' && C <= 'f') ? \
+                           C-'a'+10 : -1)))
+
+      dst_len = strlen (comps);
+      /* Skip the s/S-prefix.  */
+      dst_len--;
+
+      for (i = 0; i < dst_len; i++)
+	{
+	  indices[i] = HEXCHAR_TO_INT(comps[i+1]);
+	  /* Check if the requested component is invalid or exceeds
+	     the vector.  */
+	  if (indices[i] < 0 || indices[i] >= src_len)
+	    error (_("Invalid OpenCL vector component accessor %s"), comps);
+	}
+    }
+  else
+    {
+      dst_len = strlen (comps);
+
+      for (i = 0; i < dst_len; i++)
+	{
+	  /* x, y, z, w */
+	  switch (comps[i])
+	  {
+	  case 'x':
+	    indices[i] = 0;
+	    break;
+	  case 'y':
+	    indices[i] = 1;
+	    break;
+	  case 'z':
+	    if (src_len < 3)
+	      error (_("Invalid OpenCL vector component accessor %s"), comps);
+	    indices[i] = 2;
+	    break;
+	  case 'w':
+	    if (src_len < 4)
+	      error (_("Invalid OpenCL vector component accessor %s"), comps);
+	    indices[i] = 3;
+	    break;
+	  default:
+	    error (_("Invalid OpenCL vector component accessor %s"), comps);
+	    break;
+	  }
+	}
+    }
+
+  /* Throw an error if the amount of requested components does not
+     result in a valid length (1, 2, 3, 4, 8, 16).  */
+  if (dst_len != 1 && dst_len != 2 && dst_len != 3 && dst_len != 4
+      && dst_len != 8 && dst_len != 16)
+    error (_("Invalid OpenCL vector component accessor %s"), comps);
+
+  v = create_value (exp->gdbarch, val, noside, indices, dst_len);
+
+  return v;
+}
+
+/* Perform the unary logical not (!) operation.  */
+
+static struct value *
+opencl_logical_not (struct expression *exp, struct value *arg)
+{
+  struct type *type = check_typedef (value_type (arg));
+  struct type *rettype;
+  struct value *ret;
+
+  if (TYPE_CODE (type) == TYPE_CODE_ARRAY && TYPE_VECTOR (type))
+    {
+      struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type));
+      LONGEST lowb, highb;
+      int i;
+
+      if (!get_array_bounds (type, &lowb, &highb))
+	error (_("Could not determine the vector bounds"));
+
+      /* Determine the resulting type of the operation and allocate the
+	 value.  */
+      rettype = lookup_opencl_vector_type (exp->gdbarch, TYPE_CODE_INT,
+					   TYPE_LENGTH (eltype), 0,
+					   highb - lowb + 1);
+      ret = allocate_value (rettype);
+
+      for (i = 0; i < highb - lowb + 1; i++)
+	{
+	  /* For vector types, the unary operator shall return a 0 if the
+	  value of its operand compares unequal to 0, and -1 (i.e. all bits
+	  set) if the value of its operand compares equal to 0.  */
+	  int tmp = value_logical_not (value_subscript (arg, i)) ? -1 : 0;
+	  memset (value_contents_writeable (ret) + i * TYPE_LENGTH (eltype),
+		  tmp, TYPE_LENGTH (eltype));
+	}
+    }
+  else
+    {
+      rettype = language_bool_type (exp->language_defn, exp->gdbarch);
+      ret = value_from_longest (rettype, value_logical_not (arg));
+    }
+
+  return ret;
+}
+
+/* Perform a relational operation on two scalar operands.  */
+
+static int
+scalar_relop (struct value *val1, struct value *val2, enum exp_opcode op)
+{
+  int ret;
+
+  switch (op)
+    {
+    case BINOP_EQUAL:
+      ret = value_equal (val1, val2);
+      break;
+    case BINOP_NOTEQUAL:
+      ret = !value_equal (val1, val2);
+      break;
+    case BINOP_LESS:
+      ret = value_less (val1, val2);
+      break;
+    case BINOP_GTR:
+      ret = value_less (val2, val1);
+      break;
+    case BINOP_GEQ:
+      ret = value_less (val2, val1) || value_equal (val1, val2);
+      break;
+    case BINOP_LEQ:
+      ret = value_less (val1, val2) || value_equal (val1, val2);
+      break;
+    case BINOP_LOGICAL_AND:
+      ret = !value_logical_not (val1) && !value_logical_not (val2);
+      break;
+    case BINOP_LOGICAL_OR:
+      ret = !value_logical_not (val1) || !value_logical_not (val2);
+      break;
+    default:
+      error (_("Attempt to perform an unsupported operation"));
+      break;
+    }
+  return ret;
+}
+
+/* Perform a relational operation on two vector operands.  */
+
+static struct value *
+vector_relop (struct expression *exp, struct value *val1, struct value *val2,
+	      enum exp_opcode op)
+{
+  struct value *ret;
+  struct type *type1, *type2, *eltype1, *eltype2, *rettype;
+  int t1_is_vec, t2_is_vec, i;
+  LONGEST lowb1, lowb2, highb1, highb2;
+
+  type1 = check_typedef (value_type (val1));
+  type2 = check_typedef (value_type (val2));
+
+  t1_is_vec = (TYPE_CODE (type1) == TYPE_CODE_ARRAY && TYPE_VECTOR (type1));
+  t2_is_vec = (TYPE_CODE (type2) == TYPE_CODE_ARRAY && TYPE_VECTOR (type2));
+
+  if (!t1_is_vec || !t2_is_vec)
+    error (_("Vector operations are not supported on scalar types"));
+
+  eltype1 = check_typedef (TYPE_TARGET_TYPE (type1));
+  eltype2 = check_typedef (TYPE_TARGET_TYPE (type2));
+
+  if (!get_array_bounds (type1,&lowb1, &highb1)
+      || !get_array_bounds (type2, &lowb2, &highb2))
+    error (_("Could not determine the vector bounds"));
+
+  /* Check whether the vector types are compatible.  */
+  if (TYPE_CODE (eltype1) != TYPE_CODE (eltype2)
+      || TYPE_LENGTH (eltype1) != TYPE_LENGTH (eltype2)
+      || TYPE_UNSIGNED (eltype1) != TYPE_UNSIGNED (eltype2)
+      || lowb1 != lowb2 || highb1 != highb2)
+    error (_("Cannot perform operation on vectors with different types"));
+
+  /* Determine the resulting type of the operation and allocate the value.  */
+  rettype = lookup_opencl_vector_type (exp->gdbarch, TYPE_CODE_INT,
+				       TYPE_LENGTH (eltype1), 0,
+				       highb1 - lowb1 + 1);
+  ret = allocate_value (rettype);
+
+  for (i = 0; i < highb1 - lowb1 + 1; i++)
+    {
+      /* For vector types, the relational, equality and logical operators shall
+	 return 0 if the specified relation is false and -1 (i.e. all bits set)
+	 if the specified relation is true.  */
+      int tmp = scalar_relop (value_subscript (val1, i),
+			      value_subscript (val2, i), op) ? -1 : 0;
+      memset (value_contents_writeable (ret) + i * TYPE_LENGTH (eltype1),
+	      tmp, TYPE_LENGTH (eltype1));
+     }
+
+  return ret;
+}
+
+/* Perform a relational operation on two operands.  */
+
+static struct value *
+opencl_relop (struct expression *exp, struct value *arg1, struct value *arg2,
+	      enum exp_opcode op)
+{
+  struct value *val;
+  struct type *type1 = check_typedef (value_type (arg1));
+  struct type *type2 = check_typedef (value_type (arg2));
+  int t1_is_vec = (TYPE_CODE (type1) == TYPE_CODE_ARRAY
+		   && TYPE_VECTOR (type1));
+  int t2_is_vec = (TYPE_CODE (type2) == TYPE_CODE_ARRAY
+		   && TYPE_VECTOR (type2));
+
+  if (!t1_is_vec && !t2_is_vec)
+    {
+      int tmp = scalar_relop (arg1, arg2, op);
+      struct type *type =
+	language_bool_type (exp->language_defn, exp->gdbarch);
+
+      val = value_from_longest (type, tmp);
+    }
+  else if (t1_is_vec && t2_is_vec)
+    {
+      val = vector_relop (exp, arg1, arg2, op);
+    }
+  else
+    {
+      /* Widen the scalar operand to a vector.  */
+      struct value **v = t1_is_vec ? &arg2 : &arg1;
+      struct type *t = t1_is_vec ? type2 : type1;
+
+      if (TYPE_CODE (t) != TYPE_CODE_FLT && !is_integral_type (t))
+	error (_("Argument to operation not a number or boolean."));
+
+      *v = value_cast (t1_is_vec ? type1 : type2, *v);
+      val = vector_relop (exp, arg1, arg2, op);
+    }
+
+  return val;
+}
+
+/* Expression evaluator for the OpenCL.  Most operations are delegated to
+   evaluate_subexp_standard; see that function for a description of the
+   arguments.  */
+
+static struct value *
+evaluate_subexp_opencl (struct type *expect_type, struct expression *exp,
+		   int *pos, enum noside noside)
+{
+  enum exp_opcode op = exp->elts[*pos].opcode;
+  struct value *arg1 = NULL;
+  struct value *arg2 = NULL;
+  struct type *type1, *type2;
+
+  switch (op)
+    {
+    /* Handle binary relational and equality operators that are either not
+       or differently defined for GNU vectors.  */
+    case BINOP_EQUAL:
+    case BINOP_NOTEQUAL:
+    case BINOP_LESS:
+    case BINOP_GTR:
+    case BINOP_GEQ:
+    case BINOP_LEQ:
+      (*pos)++;
+      arg1 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+      arg2 = evaluate_subexp (value_type (arg1), exp, pos, noside);
+
+      if (noside == EVAL_SKIP)
+	return value_from_longest (builtin_type (exp->gdbarch)->
+				   builtin_int, 1);
+
+      return opencl_relop (exp, arg1, arg2, op);
+
+    /* Handle the logical unary operator not(!).  */
+    case UNOP_LOGICAL_NOT:
+      (*pos)++;
+      arg1 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+
+      if (noside == EVAL_SKIP)
+	return value_from_longest (builtin_type (exp->gdbarch)->
+				   builtin_int, 1);
+
+      return opencl_logical_not (exp, arg1);
+
+    /* Handle the logical operator and(&&) and or(||).  */
+    case BINOP_LOGICAL_AND:
+    case BINOP_LOGICAL_OR:
+      (*pos)++;
+      arg1 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+
+      if (noside == EVAL_SKIP)
+	{
+	  arg2 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+
+	  return value_from_longest (builtin_type (exp->gdbarch)->
+				     builtin_int, 1);
+	}
+      else
+	{
+	  /* For scalar operations we need to avoid evaluating operands
+	     unecessarily.  However, for vector operations we always need to
+	     evaluate both operands.  Unfortunately we only know which of the
+	     two cases apply after we know the type of the second operand.
+	     Therefore we evaluate it once using EVAL_AVOID_SIDE_EFFECTS.  */
+	  int oldpos = *pos;
+
+	  arg2 = evaluate_subexp (NULL_TYPE, exp, pos, EVAL_AVOID_SIDE_EFFECTS);
+	  *pos = oldpos;
+	  type1 = check_typedef (value_type (arg1));
+	  type2 = check_typedef (value_type (arg2));
+
+	  if ((TYPE_CODE (type1) == TYPE_CODE_ARRAY && TYPE_VECTOR (type1))
+	      || (TYPE_CODE (type2) == TYPE_CODE_ARRAY && TYPE_VECTOR (type2)))
+	    {
+	      arg2 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+
+	      return opencl_relop (exp, arg1, arg2, op);
+	    }
+	  else
+	    {
+	      /* For scalar built-in types, only evaluate the right
+		 hand operand if the left hand operand compares
+		 unequal(&&)/equal(||) to 0.  */
+	      int res;
+	      int tmp = value_logical_not (arg1);
+
+	      if (op == BINOP_LOGICAL_OR)
+		tmp = !tmp;
+
+	      arg2 = evaluate_subexp (NULL_TYPE, exp, pos,
+				      tmp ? EVAL_SKIP : noside);
+	      type1 = language_bool_type (exp->language_defn, exp->gdbarch);
+
+	      if (op == BINOP_LOGICAL_AND)
+		res = !tmp && !value_logical_not (arg2);
+	      else /* BINOP_LOGICAL_OR */
+		res = tmp || !value_logical_not (arg2);
+
+	      return value_from_longest (type1, res);
+	    }
+	}
+
+    /* Handle the ternary selection operator.  */
+    case TERNOP_COND:
+      (*pos)++;
+      arg1 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+      type1 = check_typedef (value_type (arg1));
+      if (TYPE_CODE (type1) == TYPE_CODE_ARRAY && TYPE_VECTOR (type1))
+	{
+	  struct value *arg3, *tmp, *ret;
+	  struct type *eltype2, *type3, *eltype3;
+	  int t2_is_vec, t3_is_vec, i;
+	  LONGEST lowb1, lowb2, lowb3, highb1, highb2, highb3;
+
+	  arg2 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+	  arg3 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+	  type2 = check_typedef (value_type (arg2));
+	  type3 = check_typedef (value_type (arg3));
+	  t2_is_vec
+	    = TYPE_CODE (type2) == TYPE_CODE_ARRAY && TYPE_VECTOR (type2);
+	  t3_is_vec
+	    = TYPE_CODE (type3) == TYPE_CODE_ARRAY && TYPE_VECTOR (type3);
+
+	  /* Widen the scalar operand to a vector if necessary.  */
+	  if (t2_is_vec || !t3_is_vec)
+	    {
+	      arg3 = value_cast (type2, arg3);
+	      type3 = value_type (arg3);
+	    }
+	  else if (!t2_is_vec || t3_is_vec)
+	    {
+	      arg2 = value_cast (type3, arg2);
+	      type2 = value_type (arg2);
+	    }
+	  else if (!t2_is_vec || !t3_is_vec)
+	    {
+	      /* Throw an error if arg2 or arg3 aren't vectors.  */
+	      error (_("\
+Cannot perform conditional operation on incompatible types"));
+	    }
+
+	  eltype2 = check_typedef (TYPE_TARGET_TYPE (type2));
+	  eltype3 = check_typedef (TYPE_TARGET_TYPE (type3));
+
+	  if (!get_array_bounds (type1, &lowb1, &highb1)
+	      || !get_array_bounds (type2, &lowb2, &highb2)
+	      || !get_array_bounds (type3, &lowb3, &highb3))
+	    error (_("Could not determine the vector bounds"));
+
+	  /* Throw an error if the types of arg2 or arg3 are incompatible.  */
+	  if (TYPE_CODE (eltype2) != TYPE_CODE (eltype3)
+	      || TYPE_LENGTH (eltype2) != TYPE_LENGTH (eltype3)
+	      || TYPE_UNSIGNED (eltype2) != TYPE_UNSIGNED (eltype3)
+	      || lowb2 != lowb3 || highb2 != highb3)
+	    error (_("\
+Cannot perform operation on vectors with different types"));
+
+	  /* Throw an error if the sizes of arg1 and arg2/arg3 differ.  */
+	  if (lowb1 != lowb2 || lowb1 != lowb3
+	      || highb1 != highb2 || highb1 != highb3)
+	    error (_("\
+Cannot perform conditional operation on vectors with different sizes"));
+
+	  ret = allocate_value (type2);
+
+	  for (i = 0; i < highb1 - lowb1 + 1; i++)
+	    {
+	      tmp = value_logical_not (value_subscript (arg1, i)) ?
+		    value_subscript (arg3, i) : value_subscript (arg2, i);
+	      memcpy (value_contents_writeable (ret) +
+		      i * TYPE_LENGTH (eltype2), value_contents_all (tmp),
+		      TYPE_LENGTH (eltype2));
+	    }
+
+	  return ret;
+	}
+      else
+	{
+	  if (value_logical_not (arg1))
+	    {
+	      /* Skip the second operand.  */
+	      evaluate_subexp (NULL_TYPE, exp, pos, EVAL_SKIP);
+
+	      return evaluate_subexp (NULL_TYPE, exp, pos, noside);
+	    }
+	  else
+	    {
+	      /* Skip the third operand.  */
+	      arg2 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+	      evaluate_subexp (NULL_TYPE, exp, pos, EVAL_SKIP);
+
+	      return arg2;
+	    }
+	}
+
+    /* Handle STRUCTOP_STRUCT to allow component access on OpenCL vectors.  */
+    case STRUCTOP_STRUCT:
+      {
+	int pc = (*pos)++;
+	int tem = longest_to_int (exp->elts[pc + 1].longconst);
+
+	(*pos) += 3 + BYTES_TO_EXP_ELEM (tem + 1);
+	arg1 = evaluate_subexp (NULL_TYPE, exp, pos, noside);
+	type1 = check_typedef (value_type (arg1));
+
+	if (noside == EVAL_SKIP)
+	  {
+	    return value_from_longest (builtin_type (exp->gdbarch)->
+				       builtin_int, 1);
+	  }
+	else if (TYPE_CODE (type1) == TYPE_CODE_ARRAY && TYPE_VECTOR (type1))
+	  {
+	    return opencl_component_ref (exp, arg1, &exp->elts[pc + 2].string,
+					 noside);
+	  }
+	else
+	  {
+	    if (noside == EVAL_AVOID_SIDE_EFFECTS)
+	      return
+		  value_zero (lookup_struct_elt_type
+			      (value_type (arg1),&exp->elts[pc + 2].string, 0),
+			      lval_memory);
+	    else
+	      return value_struct_elt (&arg1, NULL,
+				       &exp->elts[pc + 2].string, NULL,
+				       "structure");
+	  }
+      }
+    default:
+      break;
+    }
+
+  return evaluate_subexp_c (expect_type, exp, pos, noside);
+}
+
+void
+opencl_language_arch_info (struct gdbarch *gdbarch,
+		      struct language_arch_info *lai)
+{
+  const struct builtin_opencl_type *builtin = builtin_opencl_type (gdbarch);
+
+  lai->string_char_type = builtin->builtin_char;
+  lai->primitive_type_vector
+    = GDBARCH_OBSTACK_CALLOC (gdbarch, nr_opencl_primitive_types + 1,
+			      struct type *);
+
+/* This macro fills the primitive_type_vector from a given type.  */
+#define FILL_TYPE_VECTOR(LAI, TYPE)\
+  LAI->primitive_type_vector [opencl_primitive_type_##TYPE]\
+    = builtin->builtin_##TYPE;\
+  LAI->primitive_type_vector [opencl_primitive_type_##TYPE##2]\
+    = builtin->builtin_##TYPE##2;\
+  LAI->primitive_type_vector [opencl_primitive_type_##TYPE##3]\
+    = builtin->builtin_##TYPE##3;\
+  LAI->primitive_type_vector [opencl_primitive_type_##TYPE##4]\
+    = builtin->builtin_##TYPE##4;\
+  LAI->primitive_type_vector [opencl_primitive_type_##TYPE##8]\
+    = builtin->builtin_##TYPE##8;\
+  LAI->primitive_type_vector [opencl_primitive_type_##TYPE##16]\
+    = builtin->builtin_##TYPE##16
+
+  FILL_TYPE_VECTOR (lai, char);
+  FILL_TYPE_VECTOR (lai, uchar);
+  FILL_TYPE_VECTOR (lai, short);
+  FILL_TYPE_VECTOR (lai, ushort);
+  FILL_TYPE_VECTOR (lai, int);
+  FILL_TYPE_VECTOR (lai, uint);
+  FILL_TYPE_VECTOR (lai, long);
+  FILL_TYPE_VECTOR (lai, ulong);
+  FILL_TYPE_VECTOR (lai, half);
+  FILL_TYPE_VECTOR (lai, float);
+  FILL_TYPE_VECTOR (lai, double);
+  lai->primitive_type_vector [opencl_primitive_type_bool]
+    = builtin->builtin_bool;
+  lai->primitive_type_vector [opencl_primitive_type_unsigned_char]
+    = builtin->builtin_unsigned_char;
+  lai->primitive_type_vector [opencl_primitive_type_unsigned_short]
+    = builtin->builtin_unsigned_short;
+  lai->primitive_type_vector [opencl_primitive_type_unsigned_int]
+    = builtin->builtin_unsigned_int;
+  lai->primitive_type_vector [opencl_primitive_type_unsigned_long]
+    = builtin->builtin_unsigned_long;
+  lai->primitive_type_vector [opencl_primitive_type_half]
+    = builtin->builtin_half;
+  lai->primitive_type_vector [opencl_primitive_type_size_t]
+    = builtin->builtin_size_t;
+  lai->primitive_type_vector [opencl_primitive_type_ptrdiff_t]
+    = builtin->builtin_ptrdiff_t;
+  lai->primitive_type_vector [opencl_primitive_type_intptr_t]
+    = builtin->builtin_intptr_t;
+  lai->primitive_type_vector [opencl_primitive_type_uintptr_t]
+    = builtin->builtin_uintptr_t;
+  lai->primitive_type_vector [opencl_primitive_type_void]
+    = builtin->builtin_void;
+
+  /* Specifies the return type of logical and relational operations.  */
+  lai->bool_type_symbol = "int";
+  lai->bool_type_default = builtin->builtin_int;
+}
+
+const struct exp_descriptor exp_descriptor_opencl =
+{
+  print_subexp_standard,
+  operator_length_standard,
+  operator_check_standard,
+  op_name_standard,
+  dump_subexp_body_standard,
+  evaluate_subexp_opencl
+};
+
+const struct language_defn opencl_language_defn =
+{
+  "opencl",			/* Language name */
+  language_opencl,
+  range_check_off,
+  type_check_off,
+  case_sensitive_on,
+  array_row_major,
+  macro_expansion_c,
+  &exp_descriptor_opencl,
+  c_parse,
+  c_error,
+  null_post_parser,
+  c_printchar,			/* Print a character constant */
+  c_printstr,			/* Function to print string constant */
+  c_emit_char,			/* Print a single char */
+  c_print_type,			/* Print a type using appropriate syntax */
+  c_print_typedef,		/* Print a typedef using appropriate syntax */
+  c_val_print,			/* Print a value using appropriate syntax */
+  c_value_print,		/* Print a top-level value */
+  NULL,				/* Language specific skip_trampoline */
+  NULL,                         /* name_of_this */
+  basic_lookup_symbol_nonlocal,	/* lookup_symbol_nonlocal */
+  basic_lookup_transparent_type,/* lookup_transparent_type */
+  NULL,				/* Language specific symbol demangler */
+  NULL,				/* Language specific class_name_from_physname */
+  c_op_print_tab,		/* expression operators for printing */
+  1,				/* c-style arrays */
+  0,				/* String lower bound */
+  default_word_break_characters,
+  default_make_symbol_completion_list,
+  opencl_language_arch_info,
+  default_print_array_index,
+  default_pass_by_reference,
+  c_get_string,
+  LANG_MAGIC
+};
+
+static void *
+build_opencl_types (struct gdbarch *gdbarch)
+{
+  struct builtin_opencl_type *builtin_opencl_type
+    = GDBARCH_OBSTACK_ZALLOC (gdbarch, struct builtin_opencl_type);
+
+/* Helper macro to create strings.  */
+#define STRINGIFY(S) #S
+/* This macro allocates and assigns the type struct pointers
+   for the vector types.  */
+#define BUILD_OCL_VTYPES(TYPE)\
+  builtin_opencl_type->builtin_##TYPE##2\
+    = init_vector_type (builtin_opencl_type->builtin_##TYPE, 2);\
+  TYPE_NAME (builtin_opencl_type->builtin_##TYPE##2) = STRINGIFY(TYPE ## 2);\
+  builtin_opencl_type->builtin_##TYPE##3\
+    = init_vector_type (builtin_opencl_type->builtin_##TYPE, 3);\
+  TYPE_NAME (builtin_opencl_type->builtin_##TYPE##3) = STRINGIFY(TYPE ## 3);\
+  TYPE_LENGTH (builtin_opencl_type->builtin_##TYPE##3)\
+    = 4 * TYPE_LENGTH (builtin_opencl_type->builtin_##TYPE);\
+  builtin_opencl_type->builtin_##TYPE##4\
+    = init_vector_type (builtin_opencl_type->builtin_##TYPE, 4);\
+  TYPE_NAME (builtin_opencl_type->builtin_##TYPE##4) = STRINGIFY(TYPE ## 4);\
+  builtin_opencl_type->builtin_##TYPE##8\
+    = init_vector_type (builtin_opencl_type->builtin_##TYPE, 8);\
+  TYPE_NAME (builtin_opencl_type->builtin_##TYPE##8) = STRINGIFY(TYPE ## 8);\
+  builtin_opencl_type->builtin_##TYPE##16\
+    = init_vector_type (builtin_opencl_type->builtin_##TYPE, 16);\
+  TYPE_NAME (builtin_opencl_type->builtin_##TYPE##16) = STRINGIFY(TYPE ## 16)
+
+  builtin_opencl_type->builtin_char
+    = arch_integer_type (gdbarch, 8, 0, "char");
+  BUILD_OCL_VTYPES (char);
+  builtin_opencl_type->builtin_uchar
+    = arch_integer_type (gdbarch, 8, 1, "uchar");
+  BUILD_OCL_VTYPES (uchar);
+  builtin_opencl_type->builtin_short
+    = arch_integer_type (gdbarch, 16, 0, "short");
+  BUILD_OCL_VTYPES (short);
+  builtin_opencl_type->builtin_ushort
+    = arch_integer_type (gdbarch, 16, 1, "ushort");
+  BUILD_OCL_VTYPES (ushort);
+  builtin_opencl_type->builtin_int
+    = arch_integer_type (gdbarch, 32, 0, "int");
+  BUILD_OCL_VTYPES (int);
+  builtin_opencl_type->builtin_uint
+    = arch_integer_type (gdbarch, 32, 1, "uint");
+  BUILD_OCL_VTYPES (uint);
+  builtin_opencl_type->builtin_long
+    = arch_integer_type (gdbarch, 64, 0, "long");
+  BUILD_OCL_VTYPES (long);
+  builtin_opencl_type->builtin_ulong
+    = arch_integer_type (gdbarch, 64, 1, "ulong");
+  BUILD_OCL_VTYPES (ulong);
+  builtin_opencl_type->builtin_half
+    = arch_float_type (gdbarch, 16, "half", floatformats_ieee_half);
+  BUILD_OCL_VTYPES (half);
+  builtin_opencl_type->builtin_float
+    = arch_float_type (gdbarch, 32, "float", floatformats_ieee_single);
+  BUILD_OCL_VTYPES (float);
+  builtin_opencl_type->builtin_double
+    = arch_float_type (gdbarch, 64, "double", floatformats_ieee_double);
+  BUILD_OCL_VTYPES (double);
+  builtin_opencl_type->builtin_bool
+    = arch_boolean_type (gdbarch, 32, 1, "bool");
+  builtin_opencl_type->builtin_unsigned_char
+    = arch_integer_type (gdbarch, 8, 1, "unsigned char");
+  builtin_opencl_type->builtin_unsigned_short
+    = arch_integer_type (gdbarch, 16, 1, "unsigned short");
+  builtin_opencl_type->builtin_unsigned_int
+    = arch_integer_type (gdbarch, 32, 1, "unsigned int");
+  builtin_opencl_type->builtin_unsigned_long
+    = arch_integer_type (gdbarch, 64, 1, "unsigned long");
+  builtin_opencl_type->builtin_size_t
+    = arch_integer_type (gdbarch, gdbarch_ptr_bit (gdbarch), 1, "size_t");
+  builtin_opencl_type->builtin_ptrdiff_t
+    = arch_integer_type (gdbarch, gdbarch_ptr_bit (gdbarch), 0, "ptrdiff_t");
+  builtin_opencl_type->builtin_intptr_t
+    = arch_integer_type (gdbarch, gdbarch_ptr_bit (gdbarch), 0, "intptr_t");
+  builtin_opencl_type->builtin_uintptr_t
+    = arch_integer_type (gdbarch, gdbarch_ptr_bit (gdbarch), 1, "uintptr_t");
+  builtin_opencl_type->builtin_void
+    = arch_type (gdbarch, TYPE_CODE_VOID, 1, "void");
+
+  return builtin_opencl_type;
+}
+
+void
+_initialize_opencl_language (void)
+{
+  opencl_type_data = gdbarch_data_register_post_init (build_opencl_types);
+  add_language (&opencl_language_defn);
+}
diff --git a/gdb/testsuite/ChangeLog b/gdb/testsuite/ChangeLog
index d68e8509b0e..1a2dde552ac 100644
--- a/gdb/testsuite/ChangeLog
+++ b/gdb/testsuite/ChangeLog
@@ -1,3 +1,25 @@
+2010-11-05  Ken Werner  <ken.werner@de.ibm.com>
+
+	* Makefile.in (ALL_SUBDIRS): Add gdb.opencl.
+	* configure.ac (AC_OUTPUT): Add gdb.opencl/Makefile.
+	* configure: Regenerate.
+	* gdb.opencl/Makefile.in: New File.
+	* gdb.opencl/datatypes.exp: Likewise.
+	* gdb.opencl/datatypes.cl: Likewise.
+	* gdb.opencl/operators.exp: Likewise.
+	* gdb.opencl/operators.cl: Likewise.
+	* gdb.opencl/vec_comps.exp: Likewise.
+	* gdb.opencl/vec_comps.cl: Likewise.
+	* gdb.opencl/convs_casts.exp: Likewise.
+	* gdb.opencl/convs_casts.cl: Likewise.
+	* lib/opencl.exp: Likewise.
+	* lib/opencl_hostapp.c: Likewise.
+	* lib/opencl_kernel.cl: Likewise.
+	* lib/cl_util.c: Likewise.
+	* lib/cl_util.c: Likewise.
+	* gdb.base/default.exp (set language): Add "opencl" to the list of
+	languages.
+
 2010-11-04  Sami Wagiaalla  <swagiaal@redhat.com>
 
 	* gdb.cp/overload.exp: Added test for inheritance overload.
diff --git a/gdb/testsuite/Makefile.in b/gdb/testsuite/Makefile.in
index 8d8d7049491..d02689b0d9e 100644
--- a/gdb/testsuite/Makefile.in
+++ b/gdb/testsuite/Makefile.in
@@ -36,8 +36,8 @@ RPATH_ENVVAR = @RPATH_ENVVAR@
 ALL_SUBDIRS = gdb.ada gdb.arch gdb.asm gdb.base gdb.cp gdb.disasm \
 	gdb.dwarf2 \
 	gdb.fortran gdb.server gdb.java gdb.mi gdb.multi \
-	gdb.objc gdb.opt gdb.pascal gdb.python gdb.threads gdb.trace \
-	gdb.xml \
+	gdb.objc gdb.opencl gdb.opt gdb.pascal gdb.python gdb.threads \
+	gdb.trace gdb.xml \
 	$(SUBDIRS)
 
 EXPECT = `if [ -f $${rootme}/../../expect/expect ] ; then \
diff --git a/gdb/testsuite/configure b/gdb/testsuite/configure
index 7b1248a2464..b523d1b782c 100755
--- a/gdb/testsuite/configure
+++ b/gdb/testsuite/configure
@@ -3515,7 +3515,7 @@ done
 
 
 
-ac_config_files="$ac_config_files Makefile gdb.ada/Makefile gdb.arch/Makefile gdb.asm/Makefile gdb.base/Makefile gdb.cp/Makefile gdb.disasm/Makefile gdb.dwarf2/Makefile gdb.fortran/Makefile gdb.server/Makefile gdb.java/Makefile gdb.mi/Makefile gdb.modula2/Makefile gdb.multi/Makefile gdb.objc/Makefile gdb.opt/Makefile gdb.pascal/Makefile gdb.python/Makefile gdb.reverse/Makefile gdb.threads/Makefile gdb.trace/Makefile gdb.xml/Makefile"
+ac_config_files="$ac_config_files Makefile gdb.ada/Makefile gdb.arch/Makefile gdb.asm/Makefile gdb.base/Makefile gdb.cp/Makefile gdb.disasm/Makefile gdb.dwarf2/Makefile gdb.fortran/Makefile gdb.server/Makefile gdb.java/Makefile gdb.mi/Makefile gdb.modula2/Makefile gdb.multi/Makefile gdb.objc/Makefile gdb.opt/Makefile gdb.pascal/Makefile gdb.python/Makefile gdb.reverse/Makefile gdb.threads/Makefile gdb.trace/Makefile gdb.xml/Makefile gdb.opencl/Makefile"
 
 cat >confcache <<\_ACEOF
 # This file is a shell script that caches the results of configure
@@ -4237,6 +4237,7 @@ do
     "gdb.threads/Makefile") CONFIG_FILES="$CONFIG_FILES gdb.threads/Makefile" ;;
     "gdb.trace/Makefile") CONFIG_FILES="$CONFIG_FILES gdb.trace/Makefile" ;;
     "gdb.xml/Makefile") CONFIG_FILES="$CONFIG_FILES gdb.xml/Makefile" ;;
+    "gdb.opencl/Makefile") CONFIG_FILES="$CONFIG_FILES gdb.opencl/Makefile" ;;
 
   *) as_fn_error "invalid argument: \`$ac_config_target'" "$LINENO" 5;;
   esac
diff --git a/gdb/testsuite/configure.ac b/gdb/testsuite/configure.ac
index c8668e52a9a..2748108ff69 100644
--- a/gdb/testsuite/configure.ac
+++ b/gdb/testsuite/configure.ac
@@ -144,6 +144,6 @@ AC_OUTPUT([Makefile \
   gdb.cp/Makefile gdb.disasm/Makefile gdb.dwarf2/Makefile \
   gdb.fortran/Makefile gdb.server/Makefile gdb.java/Makefile \
   gdb.mi/Makefile gdb.modula2/Makefile gdb.multi/Makefile \
-  gdb.objc/Makefile gdb.opt/Makefile gdb.pascal/Makefile \
+  gdb.objc/Makefile gdb.opencl/Makefile gdb.opt/Makefile gdb.pascal/Makefile \
   gdb.python/Makefile gdb.reverse/Makefile \
   gdb.threads/Makefile gdb.trace/Makefile gdb.xml/Makefile])
diff --git a/gdb/testsuite/gdb.base/default.exp b/gdb/testsuite/gdb.base/default.exp
index 7afa865d97a..b6ecdcb86fe 100644
--- a/gdb/testsuite/gdb.base/default.exp
+++ b/gdb/testsuite/gdb.base/default.exp
@@ -527,7 +527,7 @@ gdb_test "set history size" "Argument required .integer to set it to.*" "set his
 #test set history
 gdb_test "set history" "\"set history\" must be followed by the name of a history subcommand.(\[^\r\n\]*\[\r\n\])+List of set history subcommands:(\[^\r\n\]*\[\r\n\])+set history expansion -- Set history expansion on command input(\[^\r\n\]*\[\r\n\])+set history filename -- Set the filename in which to record the command history(\[^\r\n\]*\[\r\n\])+set history save -- Set saving of the history record on exit(\[^\r\n\]*\[\r\n\])+set history size -- Set the size of the command history(\[^\r\n\]*\[\r\n\])+Type \"help set history\" followed by set history subcommand name for full documentation.(\[^\r\n\]*\[\r\n\])+Command name abbreviations are allowed if unambiguous." "set history"
 #test set language
-gdb_test "set language" "Requires an argument. Valid arguments are auto, local, unknown, ada, c, c.., asm, minimal, d, fortran, objective-c, java, modula-2, pascal." "set language"
+gdb_test "set language" "Requires an argument. Valid arguments are auto, local, unknown, ada, c, c.., asm, minimal, d, fortran, objective-c, java, modula-2, opencl, pascal." "set language"
 #test set listsize
 gdb_test "set listsize" "Argument required .integer to set it to.*" "set listsize"
 #test set print "p" abbreviation
diff --git a/gdb/testsuite/gdb.opencl/Makefile.in b/gdb/testsuite/gdb.opencl/Makefile.in
new file mode 100644
index 00000000000..c12aef3e818
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/Makefile.in
@@ -0,0 +1,17 @@
+VPATH = @srcdir@
+srcdir = @srcdir@
+
+EXECUTABLES = datatypes vec_comps convs_casts operators
+
+all info install-info dvi install uninstall installcheck check:
+	@echo "Nothing to be done for $@..."
+
+clean mostlyclean:
+	-rm -f *~ *.o a.out core corefile gcore.test
+	-rm -f $(EXECUTABLES)
+
+distclean maintainer-clean realclean: clean
+	-rm -f *~ core
+	-rm -f Makefile config.status config.log
+	-rm -f *-init.exp
+	-rm -fr *.log summary detail *.plog *.sum *.psum site.*
diff --git a/gdb/testsuite/gdb.opencl/convs_casts.cl b/gdb/testsuite/gdb.opencl/convs_casts.cl
new file mode 100644
index 00000000000..a024c5140f3
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/convs_casts.cl
@@ -0,0 +1,55 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2010 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>  */
+
+int opencl_version = __OPENCL_VERSION__;
+
+#ifdef HAVE_cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+int have_cl_khr_fp64 = 1;
+#else
+int have_cl_khr_fp64 = 0;
+#endif
+
+#ifdef HAVE_cl_khr_fp16
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+int have_cl_khr_fp16 = 1;
+#else
+int have_cl_khr_fp16 = 0;
+#endif
+
+char c = 123;
+uchar uc = 123;
+short s = 123;
+ushort us = 123;
+int i = 123;
+uint ui = 123;
+long l = 123;
+ulong ul = 123;
+#ifdef cl_khr_fp16
+half h = 123.0;
+#endif
+float f = 123.0;
+#ifdef cl_khr_fp64
+double d = 123.0;
+#endif
+
+__kernel void testkernel (__global int *data)
+{
+  data[get_global_id(0)] = 1;
+}
diff --git a/gdb/testsuite/gdb.opencl/convs_casts.exp b/gdb/testsuite/gdb.opencl/convs_casts.exp
new file mode 100644
index 00000000000..34ea635c7c1
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/convs_casts.exp
@@ -0,0 +1,95 @@
+# Copyright 2010 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+#
+# Contributed by Ken Werner <ken.werner@de.ibm.com>.
+#
+# Tests GDBs support for OpenCL type conversions and casts.
+
+if $tracelevel {
+    strace $tracelevel
+}
+
+load_lib opencl.exp
+
+if { [skip_opencl_tests] } {
+    return 0
+}
+
+set testfile "convs_casts"
+set clprogram [remote_download target ${srcdir}/${subdir}/${testfile}.cl]
+
+# Compile the generic OpenCL host app
+if { [gdb_compile_opencl_hostapp "${clprogram}" "${testfile}" "" ] != "" } {
+    untested ${testfile}.exp
+    return -1
+}
+
+# Load the OpenCL app
+clean_restart ${testfile}
+
+# Set breakpoint at the OpenCL kernel
+gdb_test_multiple "break testkernel" "set pending breakpoint" {
+     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
+            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
+     }
+}
+
+gdb_run_cmd
+gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
+
+# Retrieve some information about availability of OpenCL extensions
+set have_cl_khr_fp64 [get_integer_valueof "have_cl_khr_fp64" 0]
+set have_cl_khr_fp16 [get_integer_valueof "have_cl_khr_fp16" 0]
+
+proc vec_casts { name } {
+  global have_cl_khr_fp16 have_cl_khr_fp64
+  set types {"char" "uchar" "short" "ushort" "int" "uint" "long" "ulong" "half" "float" "double"}
+  set len [llength ${types}]
+
+  for {set i 0} {$i < ${len}} {incr i} {
+    set type [lindex ${types} $i]
+
+    gdb_test "print/d (${type}2)${name}" " = \\{123, 123\\}"
+    gdb_test "print/d (${type}3)${name}" " = \\{123, 123, 123\\}"
+    gdb_test "print/d (${type}4)${name}" " = \\{123, 123, 123, 123\\}"
+    gdb_test "print/d (${type}8)${name}" " = \\{123, 123, 123, 123, 123, 123, 123, 123\\}"
+    gdb_test "print/d (${type}16)${name}" " = \\{123 <repeats 16 times>\\}"
+
+    gdb_test "ptype (${type}2)${name}" "${type} \\\[2\\\]"
+    gdb_test "ptype (${type}3)${name}" "${type} \\\[3\\\]"
+    gdb_test "ptype (${type}4)${name}" "${type} \\\[4\\\]"
+    gdb_test "ptype (${type}8)${name}" "${type} \\\[8\\\]"
+    gdb_test "ptype (${type}16)${name}" "${type} \\\[16\\\]"
+  }
+}
+
+vec_casts "c"
+vec_casts "uc"
+vec_casts "s"
+vec_casts "us"
+vec_casts "i"
+vec_casts "ui"
+vec_casts "l"
+vec_casts "ul"
+if { ${have_cl_khr_fp16} } {
+  vec_casts "h"
+}
+vec_casts "f"
+if { ${have_cl_khr_fp64} } {
+  vec_casts "d"
+}
+
+# Delete the OpenCL program source
+remote_file target delete ${clprogram}
diff --git a/gdb/testsuite/gdb.opencl/datatypes.cl b/gdb/testsuite/gdb.opencl/datatypes.cl
new file mode 100644
index 00000000000..c0d2a1eb337
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/datatypes.cl
@@ -0,0 +1,145 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2010 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>  */
+
+int opencl_version = __OPENCL_VERSION__;
+
+#ifdef HAVE_cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+int have_cl_khr_fp64 = 1;
+#else
+int have_cl_khr_fp64 = 0;
+#endif
+
+#ifdef HAVE_cl_khr_fp16
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+int have_cl_khr_fp16 = 1;
+#else
+int have_cl_khr_fp16 = 0;
+#endif
+
+bool b = 0;
+
+char   c   = 1;
+char2  c2  = (char2) (1, 2);
+#ifdef CL_VERSION_1_1
+char3  c3  = (char3) (1, 2, 3);
+#endif
+char4  c4  = (char4) (1, 2, 3, 4);
+char8  c8  = (char8) (1, 2, 3, 4, 5, 6, 7, 8);
+char16 c16 = (char16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+
+uchar   uc   = 1;
+uchar2  uc2  = (uchar2) (1, 2);
+#ifdef CL_VERSION_1_1
+uchar3  uc3  = (uchar3) (1, 2, 3);
+#endif
+uchar4  uc4  = (uchar4) (1, 2, 3, 4);
+uchar8  uc8  = (uchar8) (1, 2, 3, 4, 5, 6, 7, 8);
+uchar16 uc16 = (uchar16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+
+short   s   = -1;
+short2  s2  = (short2) (-1, -2);
+#ifdef CL_VERSION_1_1
+short3  s3  = (short3) (-1, -2, -3);
+#endif
+short4  s4  = (short4) (-1, -2, -3, -4);
+short8  s8  = (short8) (-1, -2, -3, -4, -5, -6, -7, -8);
+short16 s16 = (short16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
+
+ushort   us   = 1;
+ushort2  us2  = (ushort2) (1, 2);
+#ifdef CL_VERSION_1_1
+ushort3  us3  = (ushort3) (1, 2, 3);
+#endif
+ushort4  us4  = (ushort4) (1, 2, 3, 4);
+ushort8  us8  = (ushort8) (1, 2, 3, 4, 5, 6, 7, 8);
+ushort16 us16 = (ushort16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+
+int   i   = -1;
+int2  i2  = (int2) (-1, -2);
+#ifdef CL_VERSION_1_1
+int3  i3  = (int3) (-1, -2, -3);
+#endif
+int4  i4  = (int4) (-1, -2, -3, -4);
+int8  i8  = (int8) (-1, -2, -3, -4, -5, -6, -7, -8);
+int16 i16 = (int16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
+
+uint   ui   = 1;
+uint2  ui2  = (uint2) (1, 2);
+#ifdef CL_VERSION_1_1
+uint3  ui3  = (uint3) (1, 2, 3);
+#endif
+uint4  ui4  = (uint4) (1, 2, 3, 4);
+uint8  ui8  = (uint8) (1, 2, 3, 4, 5, 6, 7, 8);
+uint16 ui16 = (uint16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+
+long   l   = -1;
+long2  l2  = (long2) (-1, -2);
+#ifdef CL_VERSION_1_1
+long3  l3  = (long3) (-1, -2, -3);
+#endif
+long4  l4  = (long4) (-1, -2, -3, -4);
+long8  l8  = (long8) (-1, -2, -3, -4, -5, -6, -7, -8);
+long16 l16 = (long16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
+
+ulong   ul   = 1;
+ulong2  ul2  = (ulong2) (1, 2);
+#ifdef CL_VERSION_1_1
+ulong3  ul3  = (ulong3) (1, 2, 3);
+#endif
+ulong4  ul4  = (ulong4) (1, 2, 3, 4);
+ulong8  ul8  = (ulong8) (1, 2, 3, 4, 5, 6, 7, 8);
+ulong16 ul16 = (ulong16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+
+half *ph;
+#ifdef cl_khr_fp16
+half   h   = 1.0;
+half2  h2  = (half2) (1.0, 2.0);
+#ifdef CL_VERSION_1_1
+half3  h3  = (half3) (1.0, 2.0, 3.0);
+#endif
+half4  h4  = (half4) (1.0, 2.0, 3.0, 4.0);
+half8  h8  = (half8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
+half16 h16 = (half16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
+#endif
+
+float   f   = 1.0;
+float2  f2  = (float2) (1.0, 2.0);
+#ifdef CL_VERSION_1_1
+float3  f3  = (float3) (1.0, 2.0, 3.0);
+#endif
+float4  f4  = (float4) (1.0, 2.0, 3.0, 4.0);
+float8  f8  = (float8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
+float16 f16 = (float16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
+
+#ifdef cl_khr_fp64
+double   d   = 1.0;
+double2  d2  = (double2) (1.0, 2.0);
+#ifdef CL_VERSION_1_1
+double3  d3  = (double3) (1.0, 2.0, 3.0);
+#endif
+double4  d4  = (double4) (1.0, 2.0, 3.0, 4.0);
+double8  d8  = (double8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
+double16 d16 = (double16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
+#endif
+
+__kernel void testkernel (__global int *data)
+{
+  data[get_global_id(0)] = 1;
+}
diff --git a/gdb/testsuite/gdb.opencl/datatypes.exp b/gdb/testsuite/gdb.opencl/datatypes.exp
new file mode 100644
index 00000000000..45c9e5262fb
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/datatypes.exp
@@ -0,0 +1,471 @@
+# Copyright 2010 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+#
+# Contributed by Ken Werner <ken.werner@de.ibm.com>.
+#
+# Tests OpenCL data types.
+
+if $tracelevel {
+    strace $tracelevel
+}
+
+load_lib opencl.exp
+
+if { [skip_opencl_tests] } {
+    return 0
+}
+
+set testfile "datatypes"
+set clprogram [remote_download target ${srcdir}/${subdir}/${testfile}.cl]
+
+# Compile the generic OpenCL host app
+if { [gdb_compile_opencl_hostapp "${clprogram}" "${testfile}" "" ] != "" } {
+    untested ${testfile}.exp
+    return -1
+}
+
+gdb_exit
+gdb_start
+
+# Manually switch the language to opencl
+gdb_test_no_output "set language opencl" "No prompt when setting the language to opencl"
+
+# Check OpenCL data types (GDB)
+gdb_test "whatis bool" "type = bool"
+gdb_test "p sizeof(bool)" " = 4"
+
+gdb_test "whatis char" "type = char"
+gdb_test "p sizeof(char)" " = 1"
+gdb_test "whatis char2" "type = char2"
+gdb_test "p sizeof(char2)" " = 2"
+gdb_test "whatis char3" "type = char3"
+gdb_test "p sizeof(char3)" " = 4"
+gdb_test "whatis char4" "type = char4"
+gdb_test "p sizeof(char4)" " = 4"
+gdb_test "whatis char8" "type = char8"
+gdb_test "p sizeof(char8)" " = 8"
+gdb_test "whatis char16" "type = char16"
+gdb_test "p sizeof(char16)" " = 16"
+
+gdb_test "whatis unsigned char" "type = unsigned char"
+gdb_test "p sizeof(unsigned char)" " = 1"
+gdb_test "whatis uchar" "type = uchar"
+gdb_test "p sizeof(uchar)" " = 1"
+gdb_test "whatis uchar2" "type = uchar2"
+gdb_test "p sizeof(uchar2)" " = 2"
+gdb_test "whatis uchar3" "type = uchar3"
+gdb_test "p sizeof(uchar3)" " = 4"
+gdb_test "whatis uchar4" "type = uchar4"
+gdb_test "p sizeof(uchar4)" " = 4"
+gdb_test "whatis uchar8" "type = uchar8"
+gdb_test "p sizeof(uchar8)" " = 8"
+gdb_test "whatis uchar16" "type = uchar16"
+gdb_test "p sizeof(uchar16)" " = 16"
+
+gdb_test "whatis short" "type = short"
+gdb_test "p sizeof(short)" " = 2"
+gdb_test "whatis short2" "type = short2"
+gdb_test "p sizeof(short2)" " = 4"
+gdb_test "whatis short3" "type = short3"
+gdb_test "p sizeof(short3)" " = 8"
+gdb_test "whatis short4" "type = short4"
+gdb_test "p sizeof(short4)" " = 8"
+gdb_test "whatis short8" "type = short8"
+gdb_test "p sizeof(short8)" " = 16"
+gdb_test "whatis short16" "type = short16"
+gdb_test "p sizeof(short16)" " = 32"
+
+gdb_test "whatis unsigned short" "type = unsigned short"
+gdb_test "p sizeof(unsigned short)" " = 2"
+gdb_test "whatis ushort" "type = ushort"
+gdb_test "p sizeof(ushort)" " = 2"
+gdb_test "whatis ushort2" "type = ushort2"
+gdb_test "p sizeof(ushort2)" " = 4"
+gdb_test "whatis ushort3" "type = ushort3"
+gdb_test "p sizeof(ushort3)" " = 8"
+gdb_test "whatis ushort4" "type = ushort4"
+gdb_test "p sizeof(ushort4)" " = 8"
+gdb_test "whatis ushort8" "type = ushort8"
+gdb_test "p sizeof(ushort8)" " = 16"
+gdb_test "whatis ushort16" "type = ushort16"
+gdb_test "p sizeof(ushort16)" " = 32"
+
+gdb_test "whatis int" "type = int"
+gdb_test "p sizeof(int)" " = 4"
+gdb_test "whatis int2" "type = int2"
+gdb_test "p sizeof(int2)" " = 8"
+gdb_test "whatis int3" "type = int3"
+gdb_test "p sizeof(int3)" " = 16"
+gdb_test "whatis int4" "type = int4"
+gdb_test "p sizeof(int4)" " = 16"
+gdb_test "whatis int8" "type = int8"
+gdb_test "p sizeof(int8)" " = 32"
+gdb_test "whatis int16" "type = int16"
+gdb_test "p sizeof(int16)" " = 64"
+
+gdb_test "whatis unsigned int" "type = unsigned int"
+gdb_test "p sizeof(unsigned int)" " = 4"
+gdb_test "whatis uint" "type = uint"
+gdb_test "p sizeof(uint)" " = 4"
+gdb_test "whatis uint2" "type = uint2"
+gdb_test "p sizeof(uint2)" " = 8"
+gdb_test "whatis uint3" "type = uint3"
+gdb_test "p sizeof(uint3)" " = 16"
+gdb_test "whatis uint4" "type = uint4"
+gdb_test "p sizeof(uint4)" " = 16"
+gdb_test "whatis uint8" "type = uint8"
+gdb_test "p sizeof(uint8)" " = 32"
+gdb_test "whatis uint16" "type = uint16"
+gdb_test "p sizeof(uint16)" " = 64"
+
+gdb_test "whatis long" "type = long"
+gdb_test "p sizeof(long)" " = 8"
+gdb_test "whatis long2" "type = long2"
+gdb_test "p sizeof(long2)" " = 16"
+gdb_test "whatis long3" "type = long3"
+gdb_test "p sizeof(long3)" " = 32"
+gdb_test "whatis long4" "type = long4"
+gdb_test "p sizeof(long4)" " = 32"
+gdb_test "whatis long8" "type = long8"
+gdb_test "p sizeof(long8)" " = 64"
+gdb_test "whatis long16" "type = long16"
+gdb_test "p sizeof(long16)" " = 128"
+
+gdb_test "whatis unsigned long" "type = unsigned long"
+gdb_test "p sizeof(unsigned long)" " = 8"
+gdb_test "whatis ulong" "type = ulong"
+gdb_test "p sizeof(ulong)" " = 8"
+gdb_test "whatis ulong2" "type = ulong2"
+gdb_test "p sizeof(ulong2)" " = 16"
+gdb_test "whatis ulong3" "type = ulong3"
+gdb_test "p sizeof(ulong3)" " = 32"
+gdb_test "whatis ulong4" "type = ulong4"
+gdb_test "p sizeof(ulong4)" " = 32"
+gdb_test "whatis ulong8" "type = ulong8"
+gdb_test "p sizeof(ulong8)" " = 64"
+gdb_test "whatis ulong16" "type = ulong16"
+gdb_test "p sizeof(ulong16)" " = 128"
+
+gdb_test "whatis half" "type = half"
+gdb_test "p sizeof(half)" " = 2"
+gdb_test "whatis half2" "type = half2"
+gdb_test "p sizeof(half2)" " = 4"
+gdb_test "whatis half3" "type = half3"
+gdb_test "p sizeof(half3)" " = 8"
+gdb_test "whatis half4" "type = half4"
+gdb_test "p sizeof(half4)" " = 8"
+gdb_test "whatis half8" "type = half8"
+gdb_test "p sizeof(half8)" " = 16"
+gdb_test "whatis half16" "type = half16"
+gdb_test "p sizeof(half16)" " = 32"
+
+gdb_test "whatis float" "type = float"
+gdb_test "p sizeof(float)" " = 4"
+gdb_test "whatis float2" "type = float2"
+gdb_test "p sizeof(float2)" " = 8"
+gdb_test "whatis float3" "type = float3"
+gdb_test "p sizeof(float3)" " = 16"
+gdb_test "whatis float4" "type = float4"
+gdb_test "p sizeof(float4)" " = 16"
+gdb_test "whatis float8" "type = float8"
+gdb_test "p sizeof(float8)" " = 32"
+gdb_test "whatis float16" "type = float16"
+gdb_test "p sizeof(float16)" " = 64"
+
+gdb_test "whatis double" "type = double"
+gdb_test "p sizeof(double)" " = 8"
+gdb_test "whatis double2" "type = double2"
+gdb_test "p sizeof(double2)" " = 16"
+gdb_test "whatis double3" "type = double3"
+gdb_test "p sizeof(double3)" " = 32"
+gdb_test "whatis double4" "type = double4"
+gdb_test "p sizeof(double4)" " = 32"
+gdb_test "whatis double8" "type = double8"
+gdb_test "p sizeof(double8)" " = 64"
+gdb_test "whatis double16" "type = double16"
+gdb_test "p sizeof(double16)" " = 128"
+
+# Set the language back to the default: "auto; currently c"
+gdb_test_no_output "set language c" "No prompt when setting the language to c"
+gdb_test_no_output "set language auto" "No prompt when setting the language to auto"
+
+# Load the OpenCL app
+gdb_reinitialize_dir $srcdir/$subdir
+gdb_load ${objdir}/${subdir}/${testfile}
+
+# Set breakpoint at the OpenCL kernel
+gdb_test_multiple "break testkernel" "set pending breakpoint" {
+     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
+            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
+     }
+}
+
+gdb_run_cmd
+gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
+
+# Check if the language was switched to opencl
+gdb_test "show language" "The current source language is \"auto; currently opencl\"\."
+
+# Retrieve some information about the OpenCL version and the availability of extensions
+set opencl_version [get_integer_valueof "opencl_version" 0]
+set have_cl_khr_fp64 [get_integer_valueof "have_cl_khr_fp64" 0]
+set have_cl_khr_fp16 [get_integer_valueof "have_cl_khr_fp16" 0]
+
+# Check OpenCL data types (DWARF)
+gdb_test "whatis b" "type = bool"
+gdb_test "p sizeof(b)" " = 4"
+gdb_test "print b" " = 0"
+
+gdb_test "whatis c" "type = char"
+gdb_test "p sizeof(c)" " = 1"
+gdb_test "print/d c" " = 1"
+gdb_test "whatis c2" "type = char \\\[2\\\]"
+gdb_test "p sizeof(c2)" " = 2"
+gdb_test "print c2" " = \\{1, 2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis c3" "type = char \\\[3\\\]"
+  gdb_test "p sizeof(c3)" " = 4"
+  gdb_test "print c3" " = \\{1, 2, 3\\}"
+}
+gdb_test "whatis c4" "type = char \\\[4\\\]"
+gdb_test "p sizeof(c4)" " = 4"
+gdb_test "print c4" " = \\{1, 2, 3, 4\\}"
+gdb_test "whatis c8" "type = char \\\[8\\\]"
+gdb_test "p sizeof(c8)" " = 8"
+gdb_test "print c8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+gdb_test "whatis c16" "type = char \\\[16\\\]"
+gdb_test "p sizeof(c16)" " = 16"
+gdb_test "print c16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+
+gdb_test "whatis uc" "type = (uchar|unsigned char)"
+gdb_test "p sizeof(uc)" " = 1"
+gdb_test "print/d uc" " = 1"
+gdb_test "whatis uc2" "type = (uchar|unsigned char) \\\[2\\\]"
+gdb_test "p sizeof(uc2)" " = 2"
+gdb_test "print uc2" " = \\{1, 2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis uc3" "type = (uchar|unsigned char) \\\[3\\\]"
+  gdb_test "p sizeof(uchar3)" " = 4"
+  gdb_test "print uc3" " = \\{1, 2, 3\\}"
+}
+gdb_test "whatis uc4" "type = (uchar|unsigned char) \\\[4\\\]"
+gdb_test "p sizeof(uc4)" " = 4"
+gdb_test "print uc4" " = \\{1, 2, 3, 4\\}"
+gdb_test "whatis uc8" "type = (uchar|unsigned char) \\\[8\\\]"
+gdb_test "p sizeof(uc8)" " = 8"
+gdb_test "print uc8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+gdb_test "whatis uc16" "type = (uchar|unsigned char) \\\[16\\\]"
+gdb_test "p sizeof(uc16)" " = 16"
+gdb_test "print uc16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+
+gdb_test "whatis s" "type = short"
+gdb_test "p sizeof(s)" " = 2"
+gdb_test "print s" " = -1"
+gdb_test "whatis s2" "type = short \\\[2\\\]"
+gdb_test "p sizeof(s2)" " = 4"
+gdb_test "print s2" " = \\{-1, -2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis s3" "type = short \\\[3\\\]"
+  gdb_test "p sizeof(s3)" " = 8"
+  gdb_test "print s3" " = \\{-1, -2, -3\\}"
+}
+gdb_test "whatis s4" "type = short \\\[4\\\]"
+gdb_test "p sizeof(s4)" " = 8"
+gdb_test "print s4" " = \\{-1, -2, -3, -4\\}"
+gdb_test "whatis s8" "type = short \\\[8\\\]"
+gdb_test "p sizeof(s8)" " = 16"
+gdb_test "print s8" " = \\{-1, -2, -3, -4, -5, -6, -7, -8\\}"
+gdb_test "whatis s16" "type = short \\\[16\\\]"
+gdb_test "p sizeof(s16)" " = 32"
+gdb_test "print s16" " = \\{-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16\\}"
+
+gdb_test "whatis us" "type = (ushort|unsigned short)"
+gdb_test "p sizeof(us)" " = 2"
+gdb_test "print us" " = 1"
+gdb_test "whatis us2" "type = (ushort|unsigned short) \\\[2\\\]"
+gdb_test "p sizeof(us2)" " = 4"
+gdb_test "print us2" " = \\{1, 2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis us3" "type = (ushort|unsigned short) \\\[3\\\]"
+  gdb_test "p sizeof(us3)" " = 8"
+  gdb_test "print us3" " = \\{1, 2, 3\\}"
+}
+gdb_test "whatis us4" "type = (ushort|unsigned short) \\\[4\\\]"
+gdb_test "p sizeof(us4)" " = 8"
+gdb_test "print us4" " = \\{1, 2, 3, 4\\}"
+gdb_test "whatis us8" "type = (ushort|unsigned short) \\\[8\\\]"
+gdb_test "p sizeof(us8)" " = 16"
+gdb_test "print us8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+gdb_test "whatis us16" "type = (ushort|unsigned short) \\\[16\\\]"
+gdb_test "p sizeof(us16)" " = 32"
+gdb_test "print us16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+
+gdb_test "whatis i" "type = int"
+gdb_test "p sizeof(i)" " = 4"
+gdb_test "print i" " = -1"
+gdb_test "whatis i2" "type = int \\\[2\\\]"
+gdb_test "p sizeof(i2)" " = 8"
+gdb_test "print i2" " = \\{-1, -2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis i3" "type = int \\\[3\\\]"
+  gdb_test "p sizeof(i3)" " = 16"
+  gdb_test "print i3" " = \\{-1, -2, -3\\}"
+}
+gdb_test "whatis i4" "type = int \\\[4\\\]"
+gdb_test "p sizeof(i4)" " = 16"
+gdb_test "print i4" " = \\{-1, -2, -3, -4\\}"
+gdb_test "whatis i8" "type = int \\\[8\\\]"
+gdb_test "p sizeof(i8)" " = 32"
+gdb_test "print i8" " = \\{-1, -2, -3, -4, -5, -6, -7, -8\\}"
+gdb_test "whatis i16" "type = int \\\[16\\\]"
+gdb_test "p sizeof(i16)" " = 64"
+gdb_test "print i16" " = \\{-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16\\}"
+
+gdb_test "whatis ui" "type = (uint|unsigned int)"
+gdb_test "p sizeof(ui)" " = 4"
+gdb_test "print ui" " = 1"
+gdb_test "whatis ui2" "type = (uint|unsigned int) \\\[2\\\]"
+gdb_test "p sizeof(ui2)" " = 8"
+gdb_test "print ui2" " = \\{1, 2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis ui3" "type = (uint|unsigned int) \\\[3\\\]"
+  gdb_test "p sizeof(ui3)" " = 16"
+  gdb_test "print ui3" " = \\{1, 2, 3\\}"
+}
+gdb_test "whatis ui4" "type = (uint|unsigned int) \\\[4\\\]"
+gdb_test "p sizeof(ui4)" " = 16"
+gdb_test "print ui4" " = \\{1, 2, 3, 4\\}"
+gdb_test "whatis ui8" "type = (uint|unsigned int) \\\[8\\\]"
+gdb_test "p sizeof(ui8)" " = 32"
+gdb_test "print ui8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+gdb_test "whatis ui16" "type = (uint|unsigned int) \\\[16\\\]"
+gdb_test "p sizeof(ui16)" " = 64"
+gdb_test "print ui16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+
+gdb_test "whatis l" "type = long"
+gdb_test "p sizeof(l)" " = 8"
+gdb_test "print l" " = -1"
+gdb_test "whatis l2" "type = long \\\[2\\\]"
+gdb_test "p sizeof(l2)" " = 16"
+gdb_test "print l2" " = \\{-1, -2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis l3" "type = long \\\[3\\\]"
+  gdb_test "p sizeof(l3)" " = 32"
+  gdb_test "print l3" " = \\{-1, -2, -3\\}"
+}
+gdb_test "whatis l4" "type = long \\\[4\\\]"
+gdb_test "p sizeof(l4)" " = 32"
+gdb_test "print l4" " = \\{-1, -2, -3, -4\\}"
+gdb_test "whatis l8" "type = long \\\[8\\\]"
+gdb_test "p sizeof(l8)" " = 64"
+gdb_test "print l8" " = \\{-1, -2, -3, -4, -5, -6, -7, -8\\}"
+gdb_test "whatis l16" "type = long \\\[16\\\]"
+gdb_test "p sizeof(l16)" " = 128"
+gdb_test "print l16" " = \\{-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16\\}"
+
+gdb_test "whatis ul" "type = (ulong|unsigned long)"
+gdb_test "p sizeof(ul)" " = 8"
+gdb_test "print ul" " = 1"
+gdb_test "whatis ul2" "type = (ulong|unsigned long) \\\[2\\\]"
+gdb_test "p sizeof(ul2)" " = 16"
+gdb_test "print ul2" " = \\{1, 2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis ul3" "type = (ulong|unsigned long) \\\[3\\\]"
+  gdb_test "p sizeof(ul3)" " = 32"
+  gdb_test "print ul3" " = \\{1, 2, 3\\}"
+}
+gdb_test "whatis ul4" "type = (ulong|unsigned long) \\\[4\\\]"
+gdb_test "p sizeof(ul4)" " = 32"
+gdb_test "print ul4" " = \\{1, 2, 3, 4\\}"
+gdb_test "whatis ul8" "type = (ulong|unsigned long) \\\[8\\\]"
+gdb_test "p sizeof(ul8)" " = 64"
+gdb_test "print ul8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+gdb_test "whatis ul16" "type = (ulong|unsigned long) \\\[16\\\]"
+gdb_test "p sizeof(ul16)" " = 128"
+gdb_test "print ul16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+
+gdb_test "whatis ph" "type = half *"
+gdb_test "whatis *ph" "type = half"
+gdb_test "p sizeof(*ph)" " = 2"
+
+if { ${have_cl_khr_fp16} } {
+  gdb_test "whatis h" "type = half"
+  gdb_test "p sizeof(h)" " = 2"
+  gdb_test "print h" " = 1"
+  gdb_test "whatis h2" "type = half \\\[2\\\]"
+  gdb_test "p sizeof(h2)" " = 4"
+  gdb_test "print h2" " = \\{1, 2\\}"
+  if { ${opencl_version} >= 110 } {
+    gdb_test "whatis h3" "type = half \\\[3\\\]"
+    gdb_test "p sizeof(h3)" " = 8"
+    gdb_test "print h3" " = \\{1, 2, 3\\}"
+  }
+  gdb_test "whatis h4" "type = half \\\[4\\\]"
+  gdb_test "p sizeof(h4)" " = 8"
+  gdb_test "print h4" " = \\{1, 2, 3, 4\\}"
+  gdb_test "whatis h8" "type = half \\\[8\\\]"
+  gdb_test "p sizeof(h8)" " = 16"
+  gdb_test "print h8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+  gdb_test "whatis h16" "type = half \\\[16\\\]"
+  gdb_test "p sizeof(h16)" " = 16"
+  gdb_test "print h16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+}
+
+gdb_test "whatis f" "type = float"
+gdb_test "p sizeof(f)" " = 4"
+gdb_test "print f" " = 1"
+gdb_test "whatis f2" "type = float \\\[2\\\]"
+gdb_test "p sizeof(f2)" " = 8"
+gdb_test "print f2" " = \\{1, 2\\}"
+if { ${opencl_version} >= 110 } {
+  gdb_test "whatis f3" "type = float \\\[3\\\]"
+  gdb_test "p sizeof(f3)" " = 16"
+  gdb_test "print f3" " = \\{1, 2, 3\\}"
+}
+gdb_test "whatis f4" "type = float \\\[4\\\]"
+gdb_test "p sizeof(f4)" " = 16"
+gdb_test "print f4" " = \\{1, 2, 3, 4\\}"
+gdb_test "whatis f8" "type = float \\\[8\\\]"
+gdb_test "p sizeof(f8)" " = 32"
+gdb_test "print f8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+gdb_test "whatis f16" "type = float \\\[16\\\]"
+gdb_test "p sizeof(f16)" " = 64"
+gdb_test "print f16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+
+if { ${have_cl_khr_fp64} } {
+  gdb_test "whatis d" "type = double"
+  gdb_test "p sizeof(d)" " = 8"
+  gdb_test "print d" " = 1"
+  gdb_test "whatis d2" "type = double \\\[2\\\]"
+  gdb_test "p sizeof(d2)" " = 16"
+  gdb_test "print d2" " = \\{1, 2\\}"
+  if { ${opencl_version} >= 110 } {
+    gdb_test "whatis d3" "type = double \\\[3\\\]"
+    gdb_test "p sizeof(d3)" " = 32"
+    gdb_test "print d3" " = \\{1, 2, 3\\}"
+  }
+  gdb_test "whatis d4" "type = double \\\[4\\\]"
+  gdb_test "p sizeof(d4)" " = 32"
+  gdb_test "print d4" " = \\{1, 2, 3, 4\\}"
+  gdb_test "whatis d8" "type = double \\\[8\\\]"
+  gdb_test "p sizeof(d8)" " = 64"
+  gdb_test "print d8" " = \\{1, 2, 3, 4, 5, 6, 7, 8\\}"
+  gdb_test "whatis d16" "type = double \\\[16\\\]"
+  gdb_test "p sizeof(d16)" " = 128"
+  gdb_test "print d16" " = \\{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16\\}"
+}
+
+# Delete the OpenCL program source
+remote_file target delete ${clprogram}
diff --git a/gdb/testsuite/gdb.opencl/operators.cl b/gdb/testsuite/gdb.opencl/operators.cl
new file mode 100644
index 00000000000..0974c040d39
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/operators.cl
@@ -0,0 +1,105 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2010 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>  */
+
+int opencl_version = __OPENCL_VERSION__;
+
+#ifdef HAVE_cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+int have_cl_khr_fp64 = 1;
+#else
+int have_cl_khr_fp64 = 0;
+#endif
+
+#ifdef HAVE_cl_khr_fp16
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+int have_cl_khr_fp16 = 1;
+#else
+int have_cl_khr_fp16 = 0;
+#endif
+
+char ca = 2;
+char cb = 1;
+uchar uca = 2;
+uchar ucb = 1;
+char4 c4a = (char4) (2, 4, 8, 16);
+char4 c4b = (char4) (1, 2, 8, 4);
+uchar4 uc4a = (uchar4) (2, 4, 8, 16);
+uchar4 uc4b = (uchar4) (1, 2, 8, 4);
+
+short sa = 2;
+short sb = 1;
+ushort usa = 2;
+ushort usb = 1;
+short4 s4a = (short4) (2, 4, 8, 16);
+short4 s4b = (short4) (1, 2, 8, 4);
+ushort4 us4a = (ushort4) (2, 4, 8, 16);
+ushort4 us4b = (ushort4) (1, 2, 8, 4);
+
+int ia = 2;
+int ib = 1;
+uint uia = 2;
+uint uib = 1;
+int4 i4a = (int4) (2, 4, 8, 16);
+int4 i4b = (int4) (1, 2, 8, 4);
+uint4 ui4a = (uint4) (2, 4, 8, 16);
+uint4 ui4b = (uint4) (1, 2, 8, 4);
+
+long la = 2;
+long lb = 1;
+ulong ula = 2;
+ulong ulb = 1;
+long4 l4a = (long4) (2, 4, 8, 16);
+long4 l4b = (long4) (1, 2, 8, 4);
+ulong4 ul4a = (ulong4) (2, 4, 8, 16);
+ulong4 ul4b = (ulong4) (1, 2, 8, 4);
+
+#ifdef cl_khr_fp16
+half ha = 2;
+half hb = 1;
+half4 h4a = (half4) (2, 4, 8, 16);
+half4 h4b = (half4) (1, 2, 8, 4);
+#endif
+
+float fa = 2;
+float fb = 1;
+float4 f4a = (float4) (2, 4, 8, 16);
+float4 f4b = (float4) (1, 2, 8, 4);
+
+#ifdef cl_khr_fp64
+double da = 2;
+double db = 1;
+double4 d4a = (double4) (2, 4, 8, 16);
+double4 d4b = (double4) (1, 2, 8, 4);
+#endif
+
+uint4 ui4 = (uint4) (2, 4, 8, 16);
+int2 i2 = (int2) (1, 2);
+long2 l2 = (long2) (1, 2);
+#ifdef cl_khr_fp16
+half2 h2 = (half2) (1, 2);
+#endif
+float2 f2 = (float2) (1, 2);
+#ifdef cl_khr_fp64
+double2 d2 = (double2) (1, 2);
+#endif
+
+__kernel void testkernel (__global int *data)
+{
+  data[get_global_id(0)] = 1;
+}
diff --git a/gdb/testsuite/gdb.opencl/operators.exp b/gdb/testsuite/gdb.opencl/operators.exp
new file mode 100644
index 00000000000..b60c65c9d3c
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/operators.exp
@@ -0,0 +1,955 @@
+# Copyright 2010 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+#
+# Contributed by Ken Werner <ken.werner@de.ibm.com>.
+#
+# Tests GDBs support for OpenCL operators.
+
+if $tracelevel {
+    strace $tracelevel
+}
+
+load_lib opencl.exp
+
+if { [skip_opencl_tests] } {
+    return 0
+}
+
+set testfile "operators"
+set clprogram [remote_download target ${srcdir}/${subdir}/${testfile}.cl]
+
+# Compile the generic OpenCL host app
+if { [gdb_compile_opencl_hostapp "${clprogram}" "${testfile}" "" ] != "" } {
+    untested ${testfile}.exp
+    return -1
+}
+
+# Load the OpenCL app
+clean_restart ${testfile}
+
+# Set breakpoint at the OpenCL kernel
+gdb_test_multiple "break testkernel" "set pending breakpoint" {
+     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
+            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
+     }
+}
+
+gdb_run_cmd
+gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
+
+# Retrieve some information about availability of OpenCL extensions
+set have_cl_khr_fp64 [get_integer_valueof "have_cl_khr_fp64" 0]
+set have_cl_khr_fp16 [get_integer_valueof "have_cl_khr_fp16" 0]
+
+proc check_basic { name type isfloat } {
+  gdb_test "print/d ${name}a" " = 2"
+  gdb_test "print/d ${name}b" " = 1"
+  gdb_test "print/d ${name}4a" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d ${name}4b" " = \\{1, 2, 8, 4\\}"
+
+  gdb_test "ptype ${name}a" "type = ${type}"
+  gdb_test "ptype ${name}b" "type = ${type}"
+  gdb_test "ptype ${name}4a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4b" "type = ${type} \\\[4\\\]"
+
+  if { ! ${isfloat} } {
+    gdb_test "print/d u${name}a" " = 2"
+    gdb_test "print/d u${name}b" " = 1"
+    gdb_test "print/d u${name}4a" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4b" " = \\{1, 2, 8, 4\\}"
+    gdb_test "ptype u${name}a" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  }
+}
+
+# Arithmetic operators
+proc check_arithmetic_ops { name type isfloat size } {
+  # scalar with scalar
+  gdb_test "print/d ${name}a + ${name}b" " = 3"
+  gdb_test "print/d ${name}a - ${name}b" " = 1"
+  gdb_test "print/d ${name}a * ${name}b" " = 2"
+  gdb_test "print/d ${name}a / ${name}b" " = 2"
+  # scalar with vector
+  gdb_test "print/d ${name}a + ${name}4b" " = \\{3, 4, 10, 6\\}"
+  gdb_test "print/d ${name}4a - ${name}b" " = \\{1, 3, 7, 15\\}"
+  gdb_test "print/d ${name}4a * ${name}b" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d ${name}a / ${name}4b" " = \\{2, 1, 0, 0\\}"
+  # vector with vector
+  gdb_test "print/d ${name}4a + ${name}4b" " = \\{3, 6, 16, 20\\}"
+  gdb_test "print/d ${name}4a - ${name}4b" " = \\{1, 2, 0, 12\\}"
+  gdb_test "print/d ${name}4a * ${name}4b" " = \\{2, 8, 64, 64\\}"
+  gdb_test "print/d ${name}4a / ${name}4b" " = \\{2, 2, 1, 4\\}"
+
+  # scalar
+  gdb_test "print/d ${name}a++" " = 2"
+  gdb_test "print/d ++${name}a" " = 4"
+  gdb_test "print/d ${name}a--" " = 4"
+  gdb_test "print/d --${name}a" " = 2"
+  gdb_test "print/d +${name}a" " = 2"
+  gdb_test "print/d -${name}a" " = -2"
+  # vector
+  gdb_test "print/d ${name}4a++" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d ++${name}4a" " = \\{4, 6, 10, 18\\}"
+  gdb_test "print/d ${name}4a--" " = \\{4, 6, 10, 18\\}"
+  gdb_test "print/d --${name}4a" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d +${name}4a" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d -${name}4a" " = \\{-2, -4, -8, -16\\}"
+
+  # scalar with vector
+  gdb_test "ptype ${name}a + ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a - ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}a * ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a / ${name}b" "type = ${type} \\\[4\\\]"
+  # vector with vector
+  gdb_test "ptype ${name}4a + ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a - ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a * ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a / ${name}4b" "type = ${type} \\\[4\\\]"
+
+  # scalar
+  gdb_test "ptype ${name}a++" "type = ${type}"
+  gdb_test "ptype ++${name}a" "type = ${type}"
+  gdb_test "ptype ${name}a--" "type = ${type}"
+  gdb_test "ptype --${name}a" "type = ${type}"
+  # vector
+  gdb_test "ptype ${name}4a++" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ++${name}4a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a--" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype --${name}4a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype +${name}4a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype -${name}4a" "type = ${type} \\\[4\\\]"
+
+  if { ${isfloat} } {
+    # scalar with scalar
+    gdb_test "ptype ${name}a + ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a - ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a * ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a / ${name}b" "type = ${type}"
+    # scalar
+    gdb_test "ptype +${name}a" "type = ${type}"
+    gdb_test "ptype -${name}a" "type = ${type}"
+  } else {
+    # scalar with scalar
+    gdb_test "print/d ${name}a % ${name}b" " = 0"
+    # scalar with vector
+    gdb_test "print/d ${name}4a % ${name}b" " = \\{0, 0, 0, 0\\}"
+    # vector with vector
+    gdb_test "print/d ${name}4a % ${name}b" " = \\{0, 0, 0, 0\\}"
+
+    # scalar with scalar
+    gdb_test "print/d u${name}a + u${name}b" " = 3"
+    gdb_test "print/d u${name}a - u${name}b" " = 1"
+    gdb_test "print/d u${name}a * u${name}b" " = 2"
+    gdb_test "print/d u${name}a / u${name}b" " = 2"
+    gdb_test "print/d u${name}a % u${name}b" " = 0"
+    # scalar with vector
+    gdb_test "print/d u${name}a + u${name}4b" " = \\{3, 4, 10, 6\\}"
+    gdb_test "print/d u${name}4a - u${name}b" " = \\{1, 3, 7, 15\\}"
+    gdb_test "print/d u${name}4a * u${name}b" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}a / u${name}4b" " = \\{2, 1, 0, 0\\}"
+    gdb_test "print/d u${name}4a % u${name}b" " = \\{0, 0, 0, 0\\}"
+    # vector with vector
+    gdb_test "print/d u${name}4a + u${name}4b" " = \\{3, 6, 16, 20\\}"
+    gdb_test "print/d u${name}4a - u${name}4b" " = \\{1, 2, 0, 12\\}"
+    gdb_test "print/d u${name}4a * u${name}4b" " = \\{2, 8, 64, 64\\}"
+    gdb_test "print/d u${name}4a / u${name}4b" " = \\{2, 2, 1, 4\\}"
+    gdb_test "print/d u${name}4a % u${name}4b" " = \\{0, 0, 0, 0\\}"
+
+    # scalar
+    gdb_test "print/d u${name}a++" " = 2"
+    gdb_test "print/d ++u${name}a" " = 4"
+    gdb_test "print/d u${name}a--" " = 4"
+    gdb_test "print/d --u${name}a" " = 2"
+    gdb_test "print/d +u${name}a" " = 2"
+    gdb_test "print/x -u${name}a" " = 0x.*fe"
+    # vector
+    gdb_test "print/d u${name}4a++" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d ++u${name}4a" " = \\{4, 6, 10, 18\\}"
+    gdb_test "print/d u${name}4a--" " = \\{4, 6, 10, 18\\}"
+    gdb_test "print/d --u${name}4a" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d +u${name}4a" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/x -u${name}4a" " = \\{0x.*fe, 0x.*fc, 0x.*f8, 0x.*f0\\}"
+
+    # scalar with scalar
+    if { ${size} < 4 } {
+      gdb_test "ptype ${name}a + ${name}b" "type = int"
+      gdb_test "ptype ${name}a - ${name}b" "type = int"
+      gdb_test "ptype ${name}a * ${name}b" "type = int"
+      gdb_test "ptype ${name}a / ${name}b" "type = int"
+      gdb_test "ptype ${name}a % ${name}b" "type = int"
+      gdb_test "ptype +${name}a" "type = int"
+      gdb_test "ptype -${name}a" "type = int"
+      gdb_test "ptype u${name}a + u${name}b" "type = int"
+      gdb_test "ptype u${name}a - u${name}b" "type = int"
+      gdb_test "ptype u${name}a * u${name}b" "type = int"
+      gdb_test "ptype u${name}a / u${name}b" "type = int"
+      gdb_test "ptype u${name}a % u${name}b" "type = int"
+      gdb_test "ptype +u${name}a" "type = int"
+      gdb_test "ptype -u${name}a" "type = int"
+    } elseif { ${size} == 4 } {
+      gdb_test "ptype ${name}a + ${name}b" "type = int"
+      gdb_test "ptype ${name}a - ${name}b" "type = int"
+      gdb_test "ptype ${name}a * ${name}b" "type = int"
+      gdb_test "ptype ${name}a / ${name}b" "type = int"
+      gdb_test "ptype ${name}a % ${name}b" "type = int"
+      gdb_test "ptype +${name}a" "type = int"
+      gdb_test "ptype -${name}a" "type = int"
+      gdb_test "ptype u${name}a + u${name}b" "type = (unsigned int|uint)"
+      gdb_test "ptype u${name}a - u${name}b" "type = (unsigned int|uint)"
+      gdb_test "ptype u${name}a * u${name}b" "type = (unsigned int|uint)"
+      gdb_test "ptype u${name}a / u${name}b" "type = (unsigned int|uint)"
+      gdb_test "ptype u${name}a % u${name}b" "type = (unsigned int|uint)"
+      gdb_test "ptype +u${name}a" "type = (unsigned int|uint)"
+      gdb_test "ptype -u${name}a" "type = (unsigned int|uint)"
+    } else { # ${size} == 8
+      gdb_test "ptype ${name}a + ${name}b" "type = long"
+      gdb_test "ptype ${name}a - ${name}b" "type = long"
+      gdb_test "ptype ${name}a * ${name}b" "type = long"
+      gdb_test "ptype ${name}a / ${name}b" "type = long"
+      gdb_test "ptype ${name}a % ${name}b" "type = long"
+      gdb_test "ptype +${name}a" "type = long"
+      gdb_test "ptype -${name}a" "type = long"
+      gdb_test "ptype u${name}a + u${name}b" "type = (unsigned long|ulong)"
+      gdb_test "ptype u${name}a - u${name}b" "type = (unsigned long|ulong)"
+      gdb_test "ptype u${name}a * u${name}b" "type = (unsigned long|ulong)"
+      gdb_test "ptype u${name}a / u${name}b" "type = (unsigned long|ulong)"
+      gdb_test "ptype u${name}a % u${name}b" "type = (unsigned long|ulong)"
+      # scalar
+      gdb_test "ptype +u${name}a" "type = (unsigned long|ulong)"
+      gdb_test "ptype -u${name}a" "type = (unsigned long|ulong)"
+    }
+    gdb_test "ptype u${name}a++" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype ++u${name}a" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a--" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype --u${name}a" "type = (unsigned ${type}|u${type})"
+    # scalar with vector
+    gdb_test "ptype ${name}a % ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}a + u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a - u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}a * u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a / u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}a % u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    # vector with vector
+    gdb_test "ptype ${name}4a % ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}4a + u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a - u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a * u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a / u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a % u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a++" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype ++u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a--" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype --u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype +u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype -u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  }
+}
+
+# Relational operators
+proc check_relational_ops { name type isfloat size } {
+  # scalar with scalar
+  gdb_test "print/d ${name}a > ${name}b" " = 1"
+  gdb_test "print/d ${name}b < ${name}a" " = 1"
+  gdb_test "print/d ${name}b >= ${name}a" " = 0"
+  gdb_test "print/d ${name}a <= ${name}b" " = 0"
+  # scalar with vector
+  gdb_test "print/d ${name}4a > ${name}b" " = \\{-1, -1, -1, -1\\}"
+  gdb_test "print/d ${name}a < ${name}4b" " = \\{0, 0, -1, -1\\}"
+  gdb_test "print/d ${name}4a >= ${name}b" " = \\{-1, -1, -1, -1\\}"
+  gdb_test "print/d ${name}a <= ${name}4b" " = \\{0, -1, -1, -1\\}"
+  # vector with vector
+  gdb_test "print/d ${name}4a > ${name}4b" " = \\{-1, -1, 0, -1\\}"
+  gdb_test "print/d ${name}4b < ${name}4a" " = \\{-1, -1, 0, -1\\}"
+  gdb_test "print/d ${name}4b >= ${name}4a" " = \\{0, 0, -1, 0\\}"
+  gdb_test "print/d ${name}4a <= ${name}4b" " = \\{0, 0, -1, 0\\}"
+
+  # result type should be int for scalars
+  gdb_test "ptype ${name}a < ${name}b" "type = int"
+  gdb_test "ptype ${name}a > ${name}b" "type = int"
+  gdb_test "ptype ${name}a <= ${name}b" "type = int"
+  gdb_test "ptype ${name}a >= ${name}b" "type = int"
+
+  if { ${isfloat} } {
+    if { ${size} == 2 } {
+      # result type should be short for half precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a > ${name}b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}a < ${name}4b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}4a >= ${name}b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}a <= ${name}4b" "type = short \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype ${name}4a > ${name}4b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}4a < ${name}4b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}4a >= ${name}4b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}4a <= ${name}4b" "type = short \\\[4\\\]"
+    } elseif { ${size} == 4 } {
+      # result type should be int for single precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a > ${name}b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}a < ${name}4b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}4a >= ${name}b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}a <= ${name}4b" "type = int \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype ${name}4a > ${name}4b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}4a < ${name}4b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}4a >= ${name}4b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}4a <= ${name}4b" "type = int \\\[4\\\]"
+    } else { # ${size} == 8
+      # result type should be long for double precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a > ${name}b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}a < ${name}4b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}4a >= ${name}b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}a <= ${name}4b" "type = long \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype ${name}4a > ${name}4b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}4a < ${name}4b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}4a >= ${name}4b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}4a <= ${name}4b" "type = long \\\[4\\\]"
+    }
+  } else {
+    # scalar with scalar
+    gdb_test "print/d u${name}a > u${name}b" " = 1"
+    gdb_test "print/d u${name}b < u${name}a" " = 1"
+    gdb_test "print/d u${name}b >= u${name}a" " = 0"
+    gdb_test "print/d u${name}a <= u${name}b" " = 0"
+    # scalar with vector
+    gdb_test "print/d u${name}4a > u${name}b" " = \\{-1, -1, -1, -1\\}"
+    gdb_test "print/d u${name}a < u${name}4b" " = \\{0, 0, -1, -1\\}"
+    gdb_test "print/d u${name}4a >= u${name}b" " = \\{-1, -1, -1, -1\\}"
+    gdb_test "print/d u${name}a <= u${name}4b" " = \\{0, -1, -1, -1\\}"
+    # vector with vector
+    gdb_test "print/d u${name}4a > u${name}4b" " = \\{-1, -1, 0, -1\\}"
+    gdb_test "print/d u${name}4b < u${name}4a" " = \\{-1, -1, 0, -1\\}"
+    gdb_test "print/d u${name}4b >= u${name}4a" " = \\{0, 0, -1, 0\\}"
+    gdb_test "print/d u${name}4a <= u${name}4b" " = \\{0, 0, -1, 0\\}"
+
+    # result type for unsigned operands is signed
+    # scalar with scalar
+    gdb_test "ptype u${name}a < u${name}b" "type = int"
+    gdb_test "ptype u${name}a > u${name}b" "type = int"
+    gdb_test "ptype u${name}a <= u${name}b" "type = int"
+    gdb_test "ptype u${name}a >= u${name}b" "type = int"
+    # scalar with vector
+    gdb_test "ptype u${name}4a > u${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}a < u${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}4a >= u${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}a <= u${name}4b" "type = ${type} \\\[4\\\]"
+    # vector with vector
+    gdb_test "ptype u${name}4a > u${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}4a < u${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}4a >= u${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}4a <= u${name}4b" "type = ${type} \\\[4\\\]"
+  }
+}
+
+# Equality operators
+proc check_equality_ops { name type isfloat size } {
+  # scalar with scalar
+  gdb_test "print/d ${name}a == ${name}b" " = 0"
+  gdb_test "print/d ${name}a != ${name}b" " = 1"
+  # scalar with vector
+  gdb_test "print/d ${name}4a == ${name}b" " = \\{0, 0, 0, 0\\}"
+  gdb_test "print/d ${name}a != ${name}4b" " = \\{-1, 0, -1, -1\\}"
+  # vector with vector
+  gdb_test "print/d ${name}4a == ${name}4b" " = \\{0, 0, -1, 0\\}"
+  gdb_test "print/d ${name}4a != ${name}4b" " = \\{-1, -1, 0, -1\\}"
+
+  # scalar with scalar
+  gdb_test "ptype ${name}a == ${name}b" "type = int"
+  gdb_test "ptype ${name}a != ${name}b" "type = int"
+
+  if { ${isfloat} } {
+    if { ${size} == 2 } {
+      # result type should be short for half precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a == ${name}b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}a != ${name}4b" "type = short \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype ${name}4a == ${name}4b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}4a != ${name}4b" "type = short \\\[4\\\]"
+    } elseif { ${size} == 4 } {
+      # result type should be int for single precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a == ${name}b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}a != ${name}4b" "type = int \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype ${name}4a == ${name}4b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}4a != ${name}4b" "type = int \\\[4\\\]"
+    } else { # ${size} == 8
+      # result type should be long for double precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a == ${name}b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}a != ${name}4b" "type = long \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype ${name}4a == ${name}4b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}4a != ${name}4b" "type = long \\\[4\\\]"
+    }
+  } else {
+    # scalar with scalar
+    gdb_test "print/d u${name}a == u${name}b" " = 0"
+    gdb_test "print/d u${name}a != u${name}b" " = 1"
+    # scalar with vector
+    gdb_test "print/d u${name}4a == u${name}b" " = \\{0, 0, 0, 0\\}"
+    gdb_test "print/d u${name}a != u${name}4b" " = \\{-1, 0, -1, -1\\}"
+    # vector with vector
+    gdb_test "print/d u${name}4a == u${name}4b" " = \\{0, 0, -1, 0\\}"
+    gdb_test "print/d u${name}4b != u${name}4a" " = \\{-1, -1, 0, -1\\}"
+
+    # result type for unsigned operands is signed
+    # scalar with scalar
+    gdb_test "ptype u${name}a == u${name}b" "type = int"
+    gdb_test "ptype u${name}a != u${name}b" "type = int"
+    # scalar with vector
+    gdb_test "ptype u${name}4a == u${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}a != u${name}4b" "type = ${type} \\\[4\\\]"
+    # vector with vector
+    gdb_test "ptype u${name}4a == u${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}4a != u${name}4b" "type = ${type} \\\[4\\\]"
+  }
+}
+
+# Shift operators
+proc check_shift_ops { name type size } {
+  # scalar with scalar
+  gdb_test "print/d ${name}a << ${name}b" " = 4"
+  gdb_test "print/d ${name}a >> ${name}b" " = 1"
+  gdb_test "print/d u${name}a << u${name}b" " = 4"
+  gdb_test "print/d u${name}a >> u${name}b" " = 1"
+  # scalar with vector
+  gdb_test "print/d ${name}4a << ${name}b" " = \\{4, 8, 16, 32\\}"
+  gdb_test "print/d ${name}4a >> ${name}b" " = \\{1, 2, 4, 8\\}"
+  gdb_test "print/d u${name}4a << u${name}b" " = \\{4, 8, 16, 32\\}"
+  gdb_test "print/d u${name}4a >> u${name}b" " = \\{1, 2, 4, 8\\}"
+  # vector with vector
+  if { ${size} == 1 } {
+    gdb_test "print/d ${name}4a << ${name}4b" " = \\{4, 16, 0, 0\\}"
+    gdb_test "print/d u${name}4a << u${name}4b" " = \\{4, 16, 0, 0\\}"
+  } else {
+    gdb_test "print/d ${name}4a << ${name}4b" " = \\{4, 16, 2048, 256\\}"
+    gdb_test "print/d u${name}4a << u${name}4b" " = \\{4, 16, 2048, 256\\}"
+  }
+  gdb_test "print/d ${name}4a >> ${name}4b" " = \\{1, 1, 0, 1\\}"
+  gdb_test "print/d u${name}4a >> u${name}4b" " = \\{1, 1, 0, 1\\}"
+
+  # scalar with scalar
+  if { ${size} < 4 } {
+    gdb_test "ptype ${name}a << ${name}b" "type = int"
+    gdb_test "ptype ${name}a >> ${name}b" "type = int"
+    gdb_test "ptype u${name}a << u${name}b" "type = int"
+    gdb_test "ptype u${name}a >> u${name}b" "type = int"
+  } elseif { ${size} == 4 } {
+    gdb_test "ptype ${name}a << ${name}b" "type = int"
+    gdb_test "ptype ${name}a >> ${name}b" "type = int"
+    gdb_test "ptype u${name}a << u${name}b" "type = (unsigned int|uint)"
+    gdb_test "ptype u${name}a >> u${name}b" "type = (unsigned int|uint)"
+  } else { # ${size} == 8
+    gdb_test "ptype ${name}a << ${name}b" "type = long"
+    gdb_test "ptype ${name}a >> ${name}b" "type = long"
+    gdb_test "ptype u${name}a << u${name}b" "type = (unsigned long|ulong)"
+    gdb_test "ptype u${name}a >> u${name}b" "type = (unsigned long|ulong)"
+  }
+  # scalar with vector
+  gdb_test "ptype ${name}4a << ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a >> ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype u${name}4a << u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  gdb_test "ptype u${name}4a >> u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  # vector with vector
+  gdb_test "ptype ${name}4a << ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a >> ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype u${name}4a << u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  gdb_test "ptype u${name}4a >> u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+}
+
+# Bitwise operators
+proc check_bitwise_ops { name type size } {
+  # scalar with scalar
+  gdb_test "print/d ${name}a & ${name}b" " = 0"
+  gdb_test "print/d ${name}a | ${name}b" " = 3"
+  gdb_test "print/d ${name}a ^ ${name}b" " = 3"
+  gdb_test "print/d u${name}a & u${name}b" " = 0"
+  gdb_test "print/d u${name}a | u${name}b" " = 3"
+  gdb_test "print/d u${name}a ^ u${name}b" " = 3"
+  # scalar with vector
+  gdb_test "print/d ${name}4a & ${name}b" " = \\{0, 0, 0, 0\\}"
+  gdb_test "print/d ${name}a | ${name}4b" " = \\{3, 2, 10, 6\\}"
+  gdb_test "print/d ${name}4a ^ ${name}b" " = \\{3, 5, 9, 17\\}"
+  gdb_test "print/d u${name}4a & u${name}b" " = \\{0, 0, 0, 0\\}"
+  gdb_test "print/d u${name}a | u${name}4b" " = \\{3, 2, 10, 6\\}"
+  gdb_test "print/d u${name}4a ^ u${name}b" " = \\{3, 5, 9, 17\\}"
+  # vector with vector
+  gdb_test "print/d ${name}4a & ${name}4b" " = \\{0, 0, 8, 0\\}"
+  gdb_test "print/d ${name}4a | ${name}4b" " = \\{3, 6, 8, 20\\}"
+  gdb_test "print/d ${name}4a ^ ${name}4b" " = \\{3, 6, 0, 20\\}"
+  gdb_test "print/d u${name}4a & u${name}4b" " = \\{0, 0, 8, 0\\}"
+  gdb_test "print/d u${name}4a | u${name}4b" " = \\{3, 6, 8, 20\\}"
+  gdb_test "print/d u${name}4a ^ u${name}4b" " = \\{3, 6, 0, 20\\}"
+
+  # scalar with scalar
+  if { ${size} < 4 } {
+    gdb_test "ptype ${name}a & ${name}b" "type = int"
+    gdb_test "ptype ${name}a | ${name}b" "type = int"
+    gdb_test "ptype ${name}a ^ ${name}b" "type = int"
+    gdb_test "ptype u${name}a & u${name}b" "type = int"
+    gdb_test "ptype u${name}a | u${name}b" "type = int"
+    gdb_test "ptype u${name}a ^ u${name}b" "type = int"
+  } elseif { ${size} == 4 } {
+    gdb_test "ptype ${name}a & ${name}b" "type = int"
+    gdb_test "ptype ${name}a | ${name}b" "type = int"
+    gdb_test "ptype ${name}a ^ ${name}b" "type = int"
+    gdb_test "ptype u${name}a & u${name}b" "type = (unsigned int|uint)"
+    gdb_test "ptype u${name}a | u${name}b" "type = (unsigned int|uint)"
+    gdb_test "ptype u${name}a ^ u${name}b" "type = (unsigned int|uint)"
+  } else { # ${size} == 8
+    gdb_test "ptype ${name}a & ${name}b" "type = long"
+    gdb_test "ptype ${name}a | ${name}b" "type = long"
+    gdb_test "ptype ${name}a ^ ${name}b" "type = long"
+    gdb_test "ptype u${name}a & u${name}b" "type = (unsigned long|ulong)"
+    gdb_test "ptype u${name}a | u${name}b" "type = (unsigned long|ulong)"
+    gdb_test "ptype u${name}a ^ u${name}b" "type = (unsigned long|ulong)"
+  }
+  # scalar with vector
+  gdb_test "ptype ${name}4a & ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}a | ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a ^ ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype u${name}4a & u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  gdb_test "ptype u${name}a | u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  gdb_test "ptype u${name}4a ^ u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  # vector with vector
+  gdb_test "ptype ${name}4a & ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a | ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a ^ ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype u${name}4a & u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  gdb_test "ptype u${name}4a | u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  gdb_test "ptype u${name}4a ^ u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+
+  # scalar
+  if { ${size} < 8 } {
+    gdb_test "print/x ~${name}a" " = 0xfffffffd"
+    gdb_test "print/x ~u${name}a" " = 0xfffffffd"
+  } else {
+    gdb_test "print/x ~${name}a" " = 0xfffffffffffffffd"
+    gdb_test "print/x ~u${name}a" " = 0xfffffffffffffffd"
+  }
+  # vector
+  if { ${size} == 1 } {
+    gdb_test "print/x ~${name}4a" " = \\{0xfd, 0xfb, 0xf7, 0xef\\}"
+    gdb_test "print/x ~u${name}4a" " = \\{0xfd, 0xfb, 0xf7, 0xef\\}"
+  } elseif { ${size} == 2 } {
+    gdb_test "print/x ~${name}4a" " = \\{0xfffd, 0xfffb, 0xfff7, 0xffef\\}"
+    gdb_test "print/x ~u${name}4a" " = \\{0xfffd, 0xfffb, 0xfff7, 0xffef\\}"
+  } elseif { ${size} == 4 } {
+    gdb_test "print/x ~${name}4a" " = \\{0xfffffffd, 0xfffffffb, 0xfffffff7, 0xffffffef\\}"
+    gdb_test "print/x ~u${name}4a" " = \\{0xfffffffd, 0xfffffffb, 0xfffffff7, 0xffffffef\\}"
+  } else { # ${size} == 8
+    gdb_test "print/x ~${name}4a" " = \\{0xfffffffffffffffd, 0xfffffffffffffffb, 0xfffffffffffffff7, 0xffffffffffffffef\\}"
+    gdb_test "print/x ~u${name}4a" " = \\{0xfffffffffffffffd, 0xfffffffffffffffb, 0xfffffffffffffff7, 0xffffffffffffffef\\}"
+  }
+  # scalar
+  if { ${size} < 4 } {
+    gdb_test "ptype ~${name}a" "type = int"
+    gdb_test "ptype ~u${name}a" "type = int"
+  } elseif { ${size} == 4 } {
+    gdb_test "ptype ~${name}a" "type = int"
+    gdb_test "ptype ~u${name}a" "type = (unsigned int|uint)"
+  } else { # ${size} == 8
+    gdb_test "ptype ~${name}a" "type = long"
+    gdb_test "ptype ~u${name}a" "type = (unsigned long|ulong)"
+  }
+  # vector
+  gdb_test "ptype ~${name}4a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ~u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+}
+
+# Logical operators
+proc check_logical_ops { name type isfloat size } {
+  # scalar
+  gdb_test "print/d !${name}a " " = 0"
+  gdb_test "print/d !!${name}a " " = 1"
+  # vector
+  gdb_test "print/d !${name}4a " " = \\{0, 0, 0, 0\\}"
+  gdb_test "print/d !!${name}4a " " = \\{-1, -1, -1, -1\\}"
+
+  # scalar with scalar
+  gdb_test "print/d ${name}a && ${name}b" " = 1"
+  gdb_test "print/d ${name}a && !${name}b" " = 0"
+  gdb_test "print/d ${name}a || ${name}b" " = 1"
+  gdb_test "print/d ${name}a || !${name}b" " = 1"
+  gdb_test "print/d !${name}a || !${name}b" " = 0"
+
+  # scalar with vector
+  gdb_test "print/d ${name}4a && ${name}b" " = \\{-1, -1, -1, -1\\}"
+  gdb_test "print/d ${name}4a && !${name}b" " = \\{0, 0, 0, 0\\}"
+  gdb_test "print/d ${name}a || ${name}4b" " = \\{-1, -1, -1, -1\\}"
+  gdb_test "print/d ${name}a || !${name}4b" " = \\{-1, -1, -1, -1\\}"
+  gdb_test "print/d !${name}4a || !${name}b" " = \\{0, 0, 0, 0\\}"
+  # vector with vector
+  gdb_test "print/d ${name}4a && ${name}4b" " = \\{-1, -1, -1, -1\\}"
+  gdb_test "print/d ${name}4a || ${name}4b" " = \\{-1, -1, -1, -1\\}"
+
+  # result type should be int for scalars
+  gdb_test "ptype !${name}a" "type = int"
+  gdb_test "ptype ${name}a && ${name}b" "type = int"
+  gdb_test "ptype ${name}a || ${name}b" "type = int"
+
+  if { ${isfloat} } {
+    if { ${size} == 2 } {
+      # result type should be short for half precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a && ${name}b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}a || ${name}4b" "type = short \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype !${name}4a" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}4a && ${name}4b" "type = short \\\[4\\\]"
+      gdb_test "ptype ${name}4a || ${name}4b" "type = short \\\[4\\\]"
+    } elseif { ${size} == 4 } {
+      # result type should be int for single precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a && ${name}b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}a || ${name}4b" "type = int \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype !${name}4a" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}4a && ${name}4b" "type = int \\\[4\\\]"
+      gdb_test "ptype ${name}4a || ${name}4b" "type = int \\\[4\\\]"
+    } else { # ${size} == 8
+      # result type should be long for double precision floating point vectors
+      # scalar with vector
+      gdb_test "ptype ${name}4a && ${name}b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}a || ${name}4b" "type = long \\\[4\\\]"
+      # vector with vector
+      gdb_test "ptype !${name}4a" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}4a && ${name}4b" "type = long \\\[4\\\]"
+      gdb_test "ptype ${name}4a || ${name}4b" "type = long \\\[4\\\]"
+    }
+  } else {
+    # unsigned scalar
+    gdb_test "print/d !u${name}a " " = 0"
+    gdb_test "print/d !!u${name}a " " = 1"
+    # unsigned vector
+    gdb_test "print/d !u${name}4a " " = \\{0, 0, 0, 0\\}"
+    gdb_test "print/d !!u${name}4a " " = \\{-1, -1, -1, -1\\}"
+
+    # scalar with scalar
+    gdb_test "print/d u${name}a && u${name}b" " = 1"
+    gdb_test "print/d u${name}a || u${name}b" " = 1"
+    # scalar with vector
+    gdb_test "print/d u${name}4a && u${name}b" " = \\{-1, -1, -1, -1\\}"
+    gdb_test "print/d u${name}a || u${name}4b" " = \\{-1, -1, -1, -1\\}"
+    # vector with vector
+    gdb_test "print/d u${name}4a && u${name}4b" " = \\{-1, -1, -1, -1\\}"
+    gdb_test "print/d u${name}4a || u${name}4b" " = \\{-1, -1, -1, -1\\}"
+
+    # scalar
+    gdb_test "ptype !u${name}a" "type = int"
+    # vector
+    gdb_test "ptype !${name}4a" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype !u${name}4a" "type = ${type} \\\[4\\\]"
+
+    # scalar with vector
+    gdb_test "ptype ${name}4a && ${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}a || ${name}4b" "type = ${type} \\\[4\\\]"
+    # result type for unsigned vector operand is signed
+    gdb_test "ptype u${name}4a && u${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}a || u${name}4b" "type = ${type} \\\[4\\\]"
+    # vector with vector
+    gdb_test "ptype ${name}4a && ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a || ${name}4b" "type = ${type} \\\[4\\\]"
+    # result type for unsigned vector operand is signed
+    gdb_test "ptype u${name}4a && u${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype u${name}4a || u${name}4b" "type = ${type} \\\[4\\\]"
+  }
+}
+
+# Conditional operator
+proc check_conditional_op { name type isfloat } {
+  # scalar with scalar
+  gdb_test "print/d ${name}a ? ${name}b : ${name}a" " = 1"
+  gdb_test "print/d !${name}a ? ${name}b : ${name}a" " = 2"
+  # scalar with vector
+  gdb_test "print/d ${name}4a ? ${name}4b : ${name}a" " = \\{1, 2, 8, 4\\}"
+  gdb_test "print/d ${name}4a ? ${name}b : ${name}4a" " = \\{1, 1, 1, 1\\}"
+  gdb_test "print/d ${name}4a > 4 ? 1 : ${name}4a" " = \\{2, 4, 1, 1\\}"
+  gdb_test "print/d ${name}4a > 4 ? ${name}4b : ${name}a" " = \\{2, 2, 8, 4\\}"
+  # vector with vector
+  gdb_test "print/d ${name}4a ? ${name}4b : ${name}4a" " = \\{1, 2, 8, 4\\}"
+  gdb_test "print/d ${name}4a > 4 ? ${name}4b : ${name}4a" " = \\{2, 4, 8, 4\\}"
+
+  # scalar with scalar
+  gdb_test "ptype ${name}a ? ${name}b : ${name}a" "type = ${type}"
+  # scalar with vector
+  gdb_test "ptype ${name}4a ? ${name}4b : ${name}a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a ? ${name}b : ${name}4a" "type = ${type} \\\[4\\\]"
+  # vector with vector
+  gdb_test "ptype ${name}4a ? ${name}4b : ${name}4a" "type = ${type} \\\[4\\\]"
+
+  if { !${isfloat} } {
+    # scalar with scalar
+    gdb_test "print/d u${name}a ? u${name}b : u${name}a" " = 1"
+    gdb_test "print/d !u${name}a ? u${name}b : u${name}a" " = 2"
+    # scalar with vector
+    gdb_test "print/d u${name}4a ? u${name}4b : u${name}a" " = \\{1, 2, 8, 4\\}"
+    gdb_test "print/d u${name}4a ? u${name}b : u${name}4a" " = \\{1, 1, 1, 1\\}"
+    gdb_test "print/d u${name}4a > 4 ? 1 : u${name}4a" " = \\{2, 4, 1, 1\\}"
+    gdb_test "print/d u${name}4a > 4 ? u${name}4b : u${name}a" " = \\{2, 2, 8, 4\\}"
+    # vector with vector
+    gdb_test "print/d u${name}4a ? u${name}4b : u${name}4a" " = \\{1, 2, 8, 4\\}"
+    gdb_test "print/d u${name}4a > 4 ? u${name}4b : u${name}4a" " = \\{2, 4, 8, 4\\}"
+
+    # scalar with scalar
+    gdb_test "ptype u${name}a ? u${name}b : u${name}a" "type = (unsigned ${type}|u${type})"
+    # scalar with vector
+    gdb_test "ptype u${name}4a ? u${name}4b : u${name}a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a ? u${name}b : u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    # vector with vector
+    gdb_test "ptype u${name}4a ? u${name}4b : u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  }
+}
+
+# Assignment operators
+proc check_assignment_ops { name type isfloat size } {
+  # scalar with scalar
+  gdb_test "print/d ${name}a = ${name}b" " = 1"
+  gdb_test "print/d ${name}a = 2" " = 2"
+  gdb_test "print/d ${name}a += ${name}b" " = 3"
+  gdb_test "print/d ${name}a -= ${name}b" " = 2"
+  gdb_test "print/d ${name}b *= ${name}a" " = 2"
+  gdb_test "print/d ${name}b /= ${name}a" " = 1"
+  # scalar with vector
+  gdb_test "print/d ${name}4a = ${name}b" " = \\{1, 1, 1, 1\\}"
+  gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d ${name}4a += ${name}b" " = \\{3, 5, 9, 17\\}"
+  gdb_test "print/d ${name}4a -= ${name}b" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d ${name}4b *= ${name}a" " = \\{2, 4, 16, 8\\}"
+  gdb_test "print/d ${name}4b /= ${name}a" " = \\{1, 2, 8, 4\\}"
+  # vector with vector
+  gdb_test "print/d ${name}4a = ${name}4b" " = \\{1, 2, 8, 4\\}"
+  gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d ${name}4a += ${name}4b" " = \\{3, 6, 16, 20\\}"
+  gdb_test "print/d ${name}4a -= ${name}4b" " = \\{2, 4, 8, 16\\}"
+  gdb_test "print/d ${name}4b *= ${name}4a" " = \\{2, 8, 64, 64\\}"
+  gdb_test "print/d ${name}4b /= ${name}4a" " = \\{1, 2, 8, 4\\}"
+
+  # scalar with scalar
+  gdb_test "ptype ${name}a = ${name}b" "type = ${type}"
+  gdb_test "ptype ${name}a += ${name}b" "type = ${type}"
+  gdb_test "ptype ${name}a -= ${name}b" "type = ${type}"
+  gdb_test "ptype ${name}a *= ${name}b" "type = ${type}"
+  gdb_test "ptype ${name}a /= ${name}b" "type = ${type}"
+  # scalar with vector
+  gdb_test "ptype ${name}4a = ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a += ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a -= ${name}b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4b *= ${name}a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4b /= ${name}a" "type = ${type} \\\[4\\\]"
+  # vector with vector
+  gdb_test "ptype ${name}4a = ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a += ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4a -= ${name}4b" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4b *= ${name}4a" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}4b /= ${name}4a" "type = ${type} \\\[4\\\]"
+
+  if { !${isfloat} } {
+    # scalar with scalar
+    gdb_test "print/d ${name}a %= ${name}b" " = 0"
+    gdb_test "print/d ${name}a = 2" " = 2"
+    gdb_test "print/d ${name}a <<= ${name}b" " = 4"
+    gdb_test "print/d ${name}a = 2" " = 2"
+    gdb_test "print/d ${name}a >>= ${name}b" " = 1"
+    gdb_test "print/d ${name}a = 2" " = 2"
+    gdb_test "print/d ${name}a &= ${name}b" " = 0"
+    gdb_test "print/d ${name}a = 2" " = 2"
+    gdb_test "print/d ${name}a |= ${name}b" " = 3"
+    gdb_test "print/d ${name}a = 2" " = 2"
+    gdb_test "print/d ${name}a ^= ${name}b" " = 3"
+    gdb_test "print/d ${name}a = 2" " = 2"
+    # scalar with vector
+    gdb_test "print/d ${name}4b %= ${name}a" " = \\{1, 0, 0, 0\\}"
+    gdb_test "print/d ${name}4b = \{1, 2, 8, 4\}" " = \\{1, 2, 8, 4\\}"
+    gdb_test "print/d ${name}4a <<= ${name}b" " = \\{4, 8, 16, 32\\}"
+    gdb_test "print/d ${name}4a >>= ${name}b" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d ${name}4a &= ${name}b" " = \\{0, 0, 0, 0\\}"
+    gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d ${name}4a |= ${name}b" " = \\{3, 5, 9, 17\\}"
+    gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d ${name}4a ^= ${name}b" " = \\{3, 5, 9, 17\\}"
+    gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    # vector with vector
+    gdb_test "print/d ${name}4b %= ${name}4a" " = \\{1, 2, 0, 4\\}"
+    gdb_test "print/d ${name}4b = \{1, 2, 8, 4\}" " = \\{1, 2, 8, 4\\}"
+    if { ${size} == 1 } {
+      gdb_test "print/d ${name}4a <<= ${name}4b" " = \\{4, 16, 0, 0\\}"
+      gdb_test "print/d ${name}4a >>= ${name}4b" " = \\{2, 4, 0, 0\\}"
+      gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    } else {
+      gdb_test "print/d ${name}4a <<= ${name}4b" " = \\{4, 16, 2048, 256\\}"
+      gdb_test "print/d ${name}4a >>= ${name}4b" " = \\{2, 4, 8, 16\\}"
+    }
+    gdb_test "print/d ${name}4a &= ${name}4b" " = \\{0, 0, 8, 0\\}"
+    gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d ${name}4a |= ${name}4b" " = \\{3, 6, 8, 20\\}"
+    gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d ${name}4a ^= ${name}4b" " = \\{3, 6, 0, 20\\}"
+    gdb_test "print/d ${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+
+    # scalar with scalar
+    gdb_test "ptype ${name}a %= ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a <<= ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a >>= ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a &= ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a |= ${name}b" "type = ${type}"
+    gdb_test "ptype ${name}a ^= ${name}b" "type = ${type}"
+    # scalar with vector
+    gdb_test "ptype ${name}4a %= ${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a <<= ${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a >>= ${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a &= ${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a |= ${name}b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a ^= ${name}b" "type = ${type} \\\[4\\\]"
+    # vector with vector
+    gdb_test "ptype ${name}4a %= ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a <<= ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a >>= ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a &= ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a |= ${name}4b" "type = ${type} \\\[4\\\]"
+    gdb_test "ptype ${name}4a ^= ${name}4b" "type = ${type} \\\[4\\\]"
+
+    # scalar with scalar
+    gdb_test "print/d u${name}a = u${name}b" " = 1"
+    gdb_test "print/d u${name}a = 2" " = 2"
+    gdb_test "print/d u${name}a += u${name}b" " = 3"
+    gdb_test "print/d u${name}a -= u${name}b" " = 2"
+    gdb_test "print/d u${name}b *= u${name}a" " = 2"
+    gdb_test "print/d u${name}b /= u${name}a" " = 1"
+    gdb_test "print/d u${name}a %= u${name}b" " = 0"
+    gdb_test "print/d u${name}a = 2" " = 2"
+    gdb_test "print/d u${name}a <<= u${name}b" " = 4"
+    gdb_test "print/d u${name}a = 2" " = 2"
+    gdb_test "print/d u${name}a >>= u${name}b" " = 1"
+    gdb_test "print/d u${name}a = 2" " = 2"
+    gdb_test "print/d u${name}a &= u${name}b" " = 0"
+    gdb_test "print/d u${name}a = 2" " = 2"
+    gdb_test "print/d u${name}a |= u${name}b" " = 3"
+    gdb_test "print/d u${name}a = 2" " = 2"
+    gdb_test "print/d u${name}a ^= u${name}b" " = 3"
+    gdb_test "print/d u${name}a = 2" " = 2"
+    # scalar with vector
+    gdb_test "print/d u${name}4a = u${name}b" " = \\{1, 1, 1, 1\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4a += u${name}b" " = \\{3, 5, 9, 17\\}"
+    gdb_test "print/d u${name}4a -= u${name}b" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4b *= u${name}a" " = \\{2, 4, 16, 8\\}"
+    gdb_test "print/d u${name}4b /= u${name}a" " = \\{1, 2, 8, 4\\}"
+    gdb_test "print/d u${name}4b %= u${name}a" " = \\{1, 0, 0, 0\\}"
+    gdb_test "print/d u${name}4b = \{1, 2, 8, 4\}" " = \\{1, 2, 8, 4\\}"
+    gdb_test "print/d u${name}4a <<= u${name}b" " = \\{4, 8, 16, 32\\}"
+    gdb_test "print/d u${name}4a >>= u${name}b" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4a &= u${name}b" " = \\{0, 0, 0, 0\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4a |= u${name}b" " = \\{3, 5, 9, 17\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4a ^= u${name}b" " = \\{3, 5, 9, 17\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    # vector with vector
+    gdb_test "print/d u${name}4a = u${name}4b" " = \\{1, 2, 8, 4\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4a += u${name}4b" " = \\{3, 6, 16, 20\\}"
+    gdb_test "print/d u${name}4a -= u${name}4b" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4b *= u${name}4a" " = \\{2, 8, 64, 64\\}"
+    gdb_test "print/d u${name}4b /= u${name}4a" " = \\{1, 2, 8, 4\\}"
+    gdb_test "print/d u${name}4b %= u${name}4a" " = \\{1, 2, 0, 4\\}"
+    gdb_test "print/d u${name}4b = \{1, 2, 8, 4\}" " = \\{1, 2, 8, 4\\}"
+    if { ${size} == 1 } {
+      gdb_test "print/d u${name}4a <<= u${name}4b" " = \\{4, 16, 0, 0\\}"
+      gdb_test "print/d u${name}4a >>= u${name}4b" " = \\{2, 4, 0, 0\\}"
+      gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    } else {
+      gdb_test "print/d u${name}4a <<= u${name}4b" " = \\{4, 16, 2048, 256\\}"
+      gdb_test "print/d u${name}4a >>= u${name}4b" " = \\{2, 4, 8, 16\\}"
+    }
+    gdb_test "print/d u${name}4a &= u${name}4b" " = \\{0, 0, 8, 0\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4a |= u${name}4b" " = \\{3, 6, 8, 20\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+    gdb_test "print/d u${name}4a ^= u${name}4b" " = \\{3, 6, 0, 20\\}"
+    gdb_test "print/d u${name}4a = \{2, 4, 8, 16\}" " = \\{2, 4, 8, 16\\}"
+
+    # scalar with scalar
+    gdb_test "ptype u${name}a = u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a += u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a -= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a *= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a /= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a %= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a <<= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a >>= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a &= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a |= u${name}b" "type = (unsigned ${type}|u${type})"
+    gdb_test "ptype u${name}a ^= u${name}b" "type = (unsigned ${type}|u${type})"
+    # scalar with vector
+    gdb_test "ptype u${name}4a = u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a += u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a -= u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4b *= u${name}a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4b /= u${name}a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a %= u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a <<= u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a >>= u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a &= u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a |= u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a ^= u${name}b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    # vector with vector
+    gdb_test "ptype u${name}4a = u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a += u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a -= u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4b *= u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4b /= u${name}4a" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a %= u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a <<= u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a >>= u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a &= u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a |= u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+    gdb_test "ptype u${name}4a ^= u${name}4b" "type = (unsigned ${type}|u${type}) \\\[4\\\]"
+  }
+}
+
+proc do_check { name type isfloat size } {
+  check_basic ${name} ${type} ${isfloat}
+  check_arithmetic_ops ${name} ${type} ${isfloat} ${size}
+  check_relational_ops ${name} ${type} ${isfloat} ${size}
+  check_equality_ops ${name} ${type} ${isfloat} ${size}
+  if { !${isfloat} } {
+    check_shift_ops ${name} ${type} ${size}
+    check_bitwise_ops ${name} ${type} ${size}
+  }
+  check_logical_ops ${name} ${type} ${isfloat} ${size}
+  check_conditional_op ${name} ${type} ${isfloat}
+  check_assignment_ops ${name} ${type} ${isfloat} ${size}
+}
+
+do_check "c" "char" 0 1
+do_check "s" "short" 0 2
+do_check "i" "int" 0 4
+do_check "l" "long" 0 8
+if { ${have_cl_khr_fp16} } {
+  do_check "h" "half" 1 2
+}
+do_check "f" "float" 1 4
+if { ${have_cl_khr_fp64} } {
+  do_check "d" "double" 1 8
+}
+# Delete the OpenCL program source
+remote_file target delete ${clprogram}
diff --git a/gdb/testsuite/gdb.opencl/vec_comps.cl b/gdb/testsuite/gdb.opencl/vec_comps.cl
new file mode 100644
index 00000000000..d58f1ba2739
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/vec_comps.cl
@@ -0,0 +1,59 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2010 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>  */
+
+int opencl_version = __OPENCL_VERSION__;
+
+#ifdef HAVE_cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+int have_cl_khr_fp64 = 1;
+#else
+int have_cl_khr_fp64 = 0;
+#endif
+
+#ifdef HAVE_cl_khr_fp16
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+int have_cl_khr_fp16 = 1;
+#else
+int have_cl_khr_fp16 = 0;
+#endif
+
+#define CREATE_VEC(TYPE, NAME)\
+  TYPE NAME =\
+  (TYPE)  (0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
+
+CREATE_VEC(char16, c16)
+CREATE_VEC(uchar16, uc16)
+CREATE_VEC(short16, s16)
+CREATE_VEC(ushort16, us16)
+CREATE_VEC(int16, i16)
+CREATE_VEC(uint16, ui16)
+CREATE_VEC(long16, l16)
+CREATE_VEC(ulong16, ul16)
+#ifdef cl_khr_fp16
+CREATE_VEC(half16, h16)
+#endif
+CREATE_VEC(float16, f16)
+#ifdef cl_khr_fp64
+CREATE_VEC(double16, d16)
+#endif
+
+__kernel void testkernel (__global int *data)
+{
+  data[get_global_id(0)] = 1;
+}
diff --git a/gdb/testsuite/gdb.opencl/vec_comps.exp b/gdb/testsuite/gdb.opencl/vec_comps.exp
new file mode 100644
index 00000000000..e044e96951f
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/vec_comps.exp
@@ -0,0 +1,390 @@
+# Copyright 2010 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.  */
+#
+# Contributed by Ken Werner <ken.werner@de.ibm.com>.
+#
+# Tests component access of OpenCL vectors.
+
+if $tracelevel {
+    strace $tracelevel
+}
+
+load_lib opencl.exp
+
+if { [skip_opencl_tests] } {
+    return 0
+}
+
+set testfile "vec_comps"
+set clprogram [remote_download target ${srcdir}/${subdir}/${testfile}.cl]
+
+# Compile the generic OpenCL host app
+if { [gdb_compile_opencl_hostapp "${clprogram}" "${testfile}" "" ] != "" } {
+    untested ${testfile}.exp
+    return -1
+}
+
+# Load the OpenCL app
+clean_restart ${testfile}
+
+# Set breakpoint at the OpenCL kernel
+gdb_test_multiple "break testkernel" "set pending breakpoint" {
+     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
+            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
+     }
+}
+
+gdb_run_cmd
+gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
+
+# Check if the language was switched to opencl
+gdb_test "show language" "The current source language is \"auto; currently opencl\"\."
+
+# Retrieve some information about the OpenCL version and the availability of extensions
+set opencl_version [get_integer_valueof "opencl_version" 0]
+set have_cl_khr_fp64 [get_integer_valueof "have_cl_khr_fp64" 0]
+set have_cl_khr_fp16 [get_integer_valueof "have_cl_khr_fp16" 0]
+
+# Sanity checks
+proc check_basic { name type size } {
+  gdb_test "ptype ${name}" "type = ${type} \\\[16\\\]"
+  gdb_test "p sizeof(${name})" " = [expr ${size} * 16]"
+  gdb_test "print/d ${name}" " = \\{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15\\}"
+}
+
+proc check_type { name type alttype } {
+  gdb_test "whatis ${name}.lo" "type = ${type}8"
+  gdb_test "whatis ${name}.hi" "type = ${type}8"
+  gdb_test "whatis ${name}.even" "type = ${type}8"
+  gdb_test "whatis ${name}.odd" "type = ${type}8"
+  gdb_test "whatis ${name}.low" "Invalid OpenCL vector component accessor low"
+  gdb_test "whatis ${name}.high" "Invalid OpenCL vector component accessor high"
+
+  gdb_test "whatis ${name}.hi.even" "type = ${type}4"
+  gdb_test "whatis ${name}.odd.odd.lo" "type = ${type}2"
+  gdb_test "whatis ${name}.even.hi.lo.odd" "type = ${alttype}|${type}"
+
+  gdb_test "whatis ${name}.x" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.y" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.z" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.w" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.v" "Invalid OpenCL vector component accessor v"
+
+  gdb_test "whatis ${name}.xy" "type = ${type}2"
+  gdb_test "whatis ${name}.xx" "type = ${type}2"
+  gdb_test "whatis ${name}.wy" "type = ${type}2"
+  gdb_test "whatis ${name}.zv" "Invalid OpenCL vector component accessor zv"
+
+  gdb_test "whatis ${name}.xyz" "type = ${type}3"
+  gdb_test "whatis ${name}.yxy" "type = ${type}3"
+  gdb_test "whatis ${name}.yzx" "type = ${type}3"
+  gdb_test "whatis ${name}.yzv" "Invalid OpenCL vector component accessor yzv"
+
+  gdb_test "whatis ${name}.xywz" "type = ${type}4"
+  gdb_test "whatis ${name}.zzyy" "type = ${type}4"
+  gdb_test "whatis ${name}.wwww" "type = ${type}4"
+  gdb_test "whatis ${name}.yxwv" "Invalid OpenCL vector component accessor yxwv"
+  gdb_test "whatis ${name}.zyxwv" "Invalid OpenCL vector component accessor zyxwv"
+
+  gdb_test "whatis ${name}.xy.x" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.wzyx.yy" "type = ${type}2"
+  gdb_test "whatis ${name}.wzyx.yx.x" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.xyzw.w" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.xy.z" "Invalid OpenCL vector component accessor z"
+
+  gdb_test "whatis ${name}.s0" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.s9" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.sa" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.sf" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.sF" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.sg" "Invalid OpenCL vector component accessor sg"
+  gdb_test "whatis ${name}.sG" "Invalid OpenCL vector component accessor sG"
+  gdb_test "whatis ${name}.Sg" "Invalid OpenCL vector component accessor Sg"
+  gdb_test "whatis ${name}.SG" "Invalid OpenCL vector component accessor SG"
+
+  gdb_test "whatis ${name}.s01" "type = ${type}2"
+  gdb_test "whatis ${name}.s00" "type = ${type}2"
+  gdb_test "whatis ${name}.sF0" "type = ${type}2"
+  gdb_test "whatis ${name}.S42" "type = ${type}2"
+
+  gdb_test "whatis ${name}.s567" "type = ${type}3"
+  gdb_test "whatis ${name}.S333" "type = ${type}3"
+  gdb_test "whatis ${name}.Sf0A" "type = ${type}3"
+  gdb_test "whatis ${name}.SB1D" "type = ${type}3"
+  gdb_test "whatis ${name}.s01g" "Invalid OpenCL vector component accessor s01g"
+
+  gdb_test "whatis ${name}.s9876" "type = ${type}4"
+  gdb_test "whatis ${name}.sFFFF" "type = ${type}4"
+  gdb_test "whatis ${name}.sCafe" "type = ${type}4"
+  gdb_test "whatis ${name}.Sf001" "type = ${type}4"
+  gdb_test "whatis ${name}.s1fg2" "Invalid OpenCL vector component accessor s1fg2"
+  gdb_test "whatis ${name}.s012345" "Invalid OpenCL vector component accessor s012345"
+
+  gdb_test "whatis ${name}.s00000000" "type = ${type}8"
+  gdb_test "whatis ${name}.s00224466" "type = ${type}8"
+  gdb_test "whatis ${name}.sDEADBEEF" "type = ${type}8"
+  gdb_test "whatis ${name}.Sa628c193" "type = ${type}8"
+
+  gdb_test "whatis ${name}.s876543210" "Invalid OpenCL vector component accessor s876543210"
+  gdb_test "whatis ${name}.s0123456789abcde" "Invalid OpenCL vector component accessor s0123456789abcde"
+
+  gdb_test "whatis ${name}.s0123456789aBcDeF" "type = ${type}16"
+  gdb_test "whatis ${name}.s0022446688AACCFF" "type = ${type}16"
+  gdb_test "whatis ${name}.S0123456776543210" "type = ${type}16"
+  gdb_test "whatis ${name}.sFEDCBA9876543210" "type = ${type}16"
+
+  gdb_test "whatis ${name}.sfedcba98.S0246" "type = ${type}4"
+  gdb_test "whatis ${name}.sfedcba98.S0246.s13" "type = ${type}2"
+  gdb_test "whatis ${name}.sfedcba98.S0246.s13.s0" "type = ${alttype}|${type}"
+  gdb_test "whatis ${name}.s0123456789abcdef.s22" "type = ${type}2"
+
+  gdb_test "whatis ${name}.hi.s7654.wx" "type = ${type}2"
+  gdb_test "whatis ${name}.s0123456789abcdef.even.lo" "type = ${type}4"
+  gdb_test "whatis ${name}.odd.xyzw.s23" "type = ${type}2"
+  gdb_test "whatis ${name}.xyzw.hi.odd" "type = ${alttype}|${type}"
+
+  gdb_test "ptype ${name}.lo" "type = ${type} \\\[8\\\]"
+  gdb_test "ptype ${name}.hi" "type = ${type} \\\[8\\\]"
+  gdb_test "ptype ${name}.even" "type = ${type} \\\[8\\\]"
+  gdb_test "ptype ${name}.odd" "type = ${type} \\\[8\\\]"
+
+  gdb_test "ptype ${name}.hi.even" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.odd.odd.lo" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.even.hi.lo.odd" "type = ${alttype}|${type}"
+
+  gdb_test "ptype ${name}.x" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.y" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.z" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.w" "type = ${alttype}|${type}"
+
+  gdb_test "ptype ${name}.xy" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.xx" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.wy" "type = ${type} \\\[2\\\]"
+
+  gdb_test "ptype ${name}.xyz" "type = ${type} \\\[3\\\]"
+  gdb_test "ptype ${name}.yxy" "type = ${type} \\\[3\\\]"
+  gdb_test "ptype ${name}.yzx" "type = ${type} \\\[3\\\]"
+
+  gdb_test "ptype ${name}.xywz" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.zzyy" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.wwww" "type = ${type} \\\[4\\\]"
+
+  gdb_test "ptype ${name}.xy.x" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.wzyx.yy" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.wzyx.yx.x" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.xyzw.w" "type = ${alttype}|${type}"
+
+  gdb_test "ptype ${name}.s0" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.s9" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.sa" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.sf" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.sF" "type = ${alttype}|${type}"
+
+  gdb_test "ptype ${name}.s01" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.s00" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.sF0" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.S42" "type = ${type} \\\[2\\\]"
+
+  gdb_test "ptype ${name}.s567" "type = ${type} \\\[3\\\]"
+  gdb_test "ptype ${name}.S333" "type = ${type} \\\[3\\\]"
+  gdb_test "ptype ${name}.Sf0A" "type = ${type} \\\[3\\\]"
+  gdb_test "ptype ${name}.SB1D" "type = ${type} \\\[3\\\]"
+
+  gdb_test "ptype ${name}.s9876" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.sFFFF" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.sCafe" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.Sf001" "type = ${type} \\\[4\\\]"
+
+  gdb_test "ptype ${name}.s00000000" "type = ${type} \\\[8\\\]"
+  gdb_test "ptype ${name}.s00224466" "type = ${type} \\\[8\\\]"
+  gdb_test "ptype ${name}.sDEADBEEF" "type = ${type} \\\[8\\\]"
+  gdb_test "ptype ${name}.Sa628c193" "type = ${type} \\\[8\\\]"
+
+  gdb_test "ptype ${name}.s0123456789aBcDeF" "type = ${type} \\\[16\\\]"
+  gdb_test "ptype ${name}.s0022446688AACCFF" "type = ${type} \\\[16\\\]"
+  gdb_test "ptype ${name}.S0123456776543210" "type = ${type} \\\[16\\\]"
+  gdb_test "ptype ${name}.sFEDCBA9876543210" "type = ${type} \\\[16\\\]"
+
+  gdb_test "ptype ${name}.sfedcba98.S0246" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.sfedcba98.S0246.s13" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.sfedcba98.S0246.s13.s0" "type = ${alttype}|${type}"
+  gdb_test "ptype ${name}.s0123456789abcdef.s22" "type = ${type} \\\[2\\\]"
+
+  gdb_test "ptype ${name}.hi.s7654.wx" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.s0123456789abcdef.even.lo" "type = ${type} \\\[4\\\]"
+  gdb_test "ptype ${name}.odd.xyzw.s23" "type = ${type} \\\[2\\\]"
+  gdb_test "ptype ${name}.xyzw.hi.odd" "type = ${alttype}|${type}"
+}
+
+proc check_sizeof { name size } {
+  gdb_test "print sizeof (${name}.lo)" " = [expr $size * 8]"
+  gdb_test "print sizeof (${name}.hi)" " = [expr $size * 8]"
+  gdb_test "print sizeof (${name}.even)" " = [expr $size * 8]"
+  gdb_test "print sizeof (${name}.odd)" " = [expr $size * 8]"
+
+  gdb_test "print sizeof (${name}.hi.even)" " = [expr $size * 4]"
+  gdb_test "print sizeof (${name}.odd.odd.lo)" " = [expr $size * 2]"
+  gdb_test "print sizeof (${name}.even.hi.lo.odd)" " = $size"
+
+  gdb_test "print sizeof (${name}.x)" " = $size"
+  gdb_test "print sizeof (${name}.xy)" " = [expr $size * 2]"
+  gdb_test "print sizeof (${name}.xyz)" " = [expr $size * 4]"
+  gdb_test "print sizeof (${name}.xyzw)" " = [expr $size * 4]"
+
+  gdb_test "print sizeof (${name}.xy.x)" " = $size"
+  gdb_test "print sizeof (${name}.wzyx.yy)" " = [expr $size * 2]"
+  gdb_test "print sizeof (${name}.wzyx.yx.x)" " = $size"
+  gdb_test "print sizeof (${name}.xyzw.w)" " = $size"
+
+  gdb_test "print sizeof (${name}.s0)" " = $size"
+  gdb_test "print sizeof (${name}.s01)" " = [expr $size * 2]"
+  gdb_test "print sizeof (${name}.s012)" " = [expr $size * 4]"
+  gdb_test "print sizeof (${name}.s0123)" " = [expr $size * 4]"
+  gdb_test "print sizeof (${name}.s01234567)" " = [expr $size * 8]"
+  gdb_test "print sizeof (${name}.s0123456789abcdef)" " = [expr $size * 16]"
+
+  gdb_test "print sizeof (${name}.sfedcba98.S0246)" " = [expr $size * 4]"
+  gdb_test "print sizeof (${name}.sfedcba98.S0246.s13)" " = [expr $size * 2]"
+  gdb_test "print sizeof (${name}.sfedcba98.S0246.s13.s0)" " = $size"
+  gdb_test "print sizeof (${name}.s0123456789abcdef.s22)" " = [expr $size * 2]"
+
+  gdb_test "print sizeof (${name}.hi.s7654.wx)" " = [expr $size * 2]"
+  gdb_test "print sizeof (${name}.s0123456789abcdef.even.lo)" " = [expr $size * 4]"
+  gdb_test "print sizeof (${name}.odd.xyzw.s23)" " = [expr $size * 2]"
+  gdb_test "print sizeof (${name}.xyzw.hi.odd)" " = $size"
+}
+
+# OpenCL vector component access
+proc check_access { name type } {
+  gdb_test "print/d ${name}.lo" " = \\{0, 1, 2, 3, 4, 5, 6, 7\\}"
+  gdb_test "print/d ${name}.hi" " = \\{8, 9, 10, 11, 12, 13, 14, 15\\}"
+  gdb_test "print/d ${name}.even" " = \\{0, 2, 4, 6, 8, 10, 12, 14\\}"
+  gdb_test "print/d ${name}.odd" " = \\{1, 3, 5, 7, 9, 11, 13, 15\\}"
+
+  gdb_test "print/d ${name}.hi.even" " = \\{8, 10, 12, 14\\}"
+  gdb_test "print/d ${name}.odd.odd.lo" " = \\{3, 7\\}"
+  gdb_test "print/d ${name}.even.hi.lo.odd" " = 10"
+
+  gdb_test "print/d ${name}.x" " = 0"
+  gdb_test "print/d ${name}.y" " = 1"
+  gdb_test "print/d ${name}.z" " = 2"
+  gdb_test "print/d ${name}.w" " = 3"
+
+  gdb_test "print/d ${name}.xy" " = \\{0, 1\\}"
+  gdb_test "print/d ${name}.xx" " = \\{0, 0\\}"
+  gdb_test "print/d ${name}.wy" " = \\{3, 1\\}"
+
+  gdb_test "print/d ${name}.xyz" " = \\{0, 1, 2\\}"
+  gdb_test "print/d ${name}.yxy" " = \\{1, 0, 1\\}"
+  gdb_test "print/d ${name}.yzx" " = \\{1, 2, 0\\}"
+
+  gdb_test "print/d ${name}.xywz" " = \\{0, 1, 3, 2\\}"
+  gdb_test "print/d ${name}.zzyy" " = \\{2, 2, 1, 1\\}"
+  gdb_test "print/d ${name}.wwww" " = \\{3, 3, 3, 3\\}"
+
+  gdb_test "print/d ${name}.xy.x" " = 0"
+  gdb_test "print/d ${name}.wzyx.yy" " = \\{2, 2\\}"
+  gdb_test "print/d ${name}.wzyx.yx.x" " = 2"
+  gdb_test "print/d ${name}.xyzw.w" " = 3"
+
+  for {set i 0} {$i < 16} {incr i} {
+    gdb_test "print/d ${name}.s[format "%x" $i]" " = $i"
+    gdb_test "print/d ${name}.S[format "%x" $i]" " = $i"
+    if {$i > 9} {
+      gdb_test "print/d ${name}.s[format "%X" $i]" " = $i"
+      gdb_test "print/d ${name}.S[format "%X" $i]" " = $i"
+    }
+  }
+
+  gdb_test "print/d ${name}.s01" " = \\{0, 1\\}"
+  gdb_test "print/d ${name}.s00" " = \\{0, 0\\}"
+  gdb_test "print/d ${name}.sF0" " = \\{15, 0\\}"
+  gdb_test "print/d ${name}.S42" " = \\{4, 2\\}"
+
+  gdb_test "print/d ${name}.s567" " = \\{5, 6, 7\\}"
+  gdb_test "print/d ${name}.S333" " = \\{3, 3, 3\\}"
+  gdb_test "print/d ${name}.Sf0A" " = \\{15, 0, 10\\}"
+  gdb_test "print/d ${name}.SB1D" " = \\{11, 1, 13\\}"
+
+  gdb_test "print/d ${name}.s9876" " = \\{9, 8, 7, 6\\}"
+  gdb_test "print/d ${name}.sFFFF" " = \\{15, 15, 15, 15\\}"
+  gdb_test "print/d ${name}.sCafe" " = \\{12, 10, 15, 14\\}"
+  gdb_test "print/d ${name}.Sf001" " = \\{15, 0, 0, 1\\}"
+
+  gdb_test "print/d ${name}.s00000000" " = \\{0, 0, 0, 0, 0, 0, 0, 0\\}"
+  gdb_test "print/d ${name}.s00224466" " = \\{0, 0, 2, 2, 4, 4, 6, 6\\}"
+  gdb_test "print/d ${name}.sDEADBEEF" " = \\{13, 14, 10, 13, 11, 14, 14, 15\\}"
+  gdb_test "print/d ${name}.Sa628c193" " = \\{10, 6, 2, 8, 12, 1, 9, 3\\}"
+
+  gdb_test "print/d ${name}.s0123456789aBcDeF" " = \\{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15\\}"
+  gdb_test "print/d ${name}.s0022446688AACCEE" " = \\{0, 0, 2, 2, 4, 4, 6, 6, 8, 8, 10, 10, 12, 12, 14, 14\\}"
+  gdb_test "print/d ${name}.S0123456776543210" " = \\{0, 1, 2, 3, 4, 5, 6, 7, 7, 6, 5, 4, 3, 2, 1, 0\\}"
+  gdb_test "print/d ${name}.sFEDCBA9876543210" " = \\{15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0\\}"
+
+  gdb_test "print/d ${name}.sfedcba98.S0246" " = \\{15, 13, 11, 9\\}"
+  gdb_test "print/d ${name}.sfedcba98.S0246.s13" " = \\{13, 9\\}"
+  gdb_test "print/d ${name}.sfedcba98.S0246.s13.s0" " = 13"
+  gdb_test "print/d ${name}.s0123456789abcdef.s22" " = \\{2, 2\\}"
+
+  gdb_test "print/d ${name}.hi.s7654.wx" " = \\{12, 15\\}"
+  gdb_test "print/d ${name}.s0123456789abcdef.even.lo" " = \\{0, 2, 4, 6\\}"
+  gdb_test "print/d ${name}.odd.xyzw.s23" " = \\{5, 7\\}"
+  gdb_test "print/d ${name}.xyzw.hi.odd" " = 3"
+
+  # lvalue tests
+  for {set i 0} {$i < 16} {incr i} {
+    gdb_test_no_output "set variable ${name}.s[format "%x" $i] = [expr 15 - $i]"
+    gdb_test "print/d ${name}.s[format "%x" $i]" " = [expr 15 - $i]"
+  }
+  gdb_test "print/d ${name}" " = \\{15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0\\}"
+
+  gdb_test_no_output "set variable ${name}.s02468ace = ${name}.s13579bdf"
+  gdb_test "print/d ${name}" " = \\{14, 14, 12, 12, 10, 10, 8, 8, 6, 6, 4, 4, 2, 2, 0, 0\\}"
+
+  gdb_test_no_output "set variable ${name}.wzyx = ${name}.even.odd"
+  gdb_test "print/d ${name}" " = \\{0, 4, 8, 12, 10, 10, 8, 8, 6, 6, 4, 4, 2, 2, 0, 0\\}"
+
+  gdb_test_no_output "set variable ${name}.odd.lo = ${name}.hi.even"
+  gdb_test "print/d ${name}" " = \\{0, 6, 8, 4, 10, 2, 8, 0, 6, 6, 4, 4, 2, 2, 0, 0\\}"
+
+  gdb_test_no_output "set variable ${name}.hi.hi.hi = ${name}.lo.s1623.lo"
+  gdb_test "print/d ${name}" " = \\{0, 6, 8, 4, 10, 2, 8, 0, 6, 6, 4, 4, 2, 2, 6, 8\\}"
+}
+
+proc do_check { name type alttype size } {
+  check_basic ${name} ${alttype} ${size}
+  check_type  ${name} ${type} ${alttype}
+  check_sizeof ${name} ${size}
+  check_access ${name} ${alttype}
+}
+
+do_check "c16" "char" "char" 1
+do_check "uc16" "uchar" "unsigned char" 1
+do_check "s16" "short" "short" 2
+do_check "us16" "ushort" "unsigned short" 2
+do_check "i16" "int" "int" 4
+do_check "ui16" "uint" "unsigned int" 4
+do_check "l16" "long" "long" 8
+do_check "ul16" "ulong" "unsigned long" 8
+if { ${have_cl_khr_fp16} } {
+  do_check "h16" "half" "half" 2
+}
+do_check "f16" "float" "float" 4
+if { ${have_cl_khr_fp64} } {
+  do_check "d16" "double" "double" 8
+}
+
+# Delete the OpenCL program source
+remote_file target delete ${clprogram}
diff --git a/gdb/testsuite/lib/cl_util.c b/gdb/testsuite/lib/cl_util.c
new file mode 100644
index 00000000000..5b731b2870d
--- /dev/null
+++ b/gdb/testsuite/lib/cl_util.c
@@ -0,0 +1,519 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2010 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>  */
+
+/* Utility macros and functions for OpenCL applications.  */
+
+#include "cl_util.h"
+
+#include <stdlib.h>
+#include <errno.h>
+#include <sys/stat.h>
+#include <string.h>
+
+const char *get_clerror_string (int errcode)
+{
+  switch (errcode)
+    {
+    case CL_SUCCESS:
+      return "CL_SUCCESS";
+    case CL_DEVICE_NOT_FOUND:
+      return "CL_DEVICE_NOT_FOUND";
+    case CL_DEVICE_NOT_AVAILABLE:
+      return "CL_DEVICE_NOT_AVAILABLE";
+    case CL_COMPILER_NOT_AVAILABLE:
+      return "CL_COMPILER_NOT_AVAILABLE";
+    case CL_MEM_OBJECT_ALLOCATION_FAILURE:
+      return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
+    case CL_OUT_OF_RESOURCES:
+      return "CL_OUT_OF_RESOURCES";
+    case CL_OUT_OF_HOST_MEMORY:
+      return "CL_OUT_OF_HOST_MEMORY";
+    case CL_PROFILING_INFO_NOT_AVAILABLE:
+      return "CL_PROFILING_INFO_NOT_AVAILABLE";
+    case CL_MEM_COPY_OVERLAP:
+      return "CL_MEM_COPY_OVERLAP";
+    case CL_IMAGE_FORMAT_MISMATCH:
+      return "CL_IMAGE_FORMAT_MISMATCH";
+    case CL_IMAGE_FORMAT_NOT_SUPPORTED:
+      return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
+    case CL_BUILD_PROGRAM_FAILURE:
+      return "CL_BUILD_PROGRAM_FAILURE";
+    case CL_MAP_FAILURE:
+      return "CL_MAP_FAILURE";
+    case CL_INVALID_VALUE:
+      return "CL_INVALID_VALUE";
+    case CL_INVALID_DEVICE_TYPE:
+      return "CL_INVALID_DEVICE_TYPE";
+    case CL_INVALID_PLATFORM:
+      return "CL_INVALID_PLATFORM";
+    case CL_INVALID_DEVICE:
+      return "CL_INVALID_DEVICE";
+    case CL_INVALID_CONTEXT:
+      return "CL_INVALID_CONTEXT";
+    case CL_INVALID_QUEUE_PROPERTIES:
+      return "CL_INVALID_QUEUE_PROPERTIES";
+    case CL_INVALID_COMMAND_QUEUE:
+      return "CL_INVALID_COMMAND_QUEUE";
+    case CL_INVALID_HOST_PTR:
+      return "CL_INVALID_HOST_PTR";
+    case CL_INVALID_MEM_OBJECT:
+      return "CL_INVALID_MEM_OBJECT";
+    case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
+      return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
+    case CL_INVALID_IMAGE_SIZE:
+      return "CL_INVALID_IMAGE_SIZE";
+    case CL_INVALID_SAMPLER:
+      return "CL_INVALID_SAMPLER";
+    case CL_INVALID_BINARY:
+      return "CL_INVALID_BINARY";
+    case CL_INVALID_BUILD_OPTIONS:
+      return "CL_INVALID_BUILD_OPTIONS";
+    case CL_INVALID_PROGRAM:
+      return "CL_INVALID_PROGRAM";
+    case CL_INVALID_PROGRAM_EXECUTABLE:
+      return "CL_INVALID_PROGRAM_EXECUTABLE";
+    case CL_INVALID_KERNEL_NAME:
+      return "CL_INVALID_KERNEL_NAME";
+    case CL_INVALID_KERNEL_DEFINITION:
+      return "CL_INVALID_KERNEL_DEFINITION";
+    case CL_INVALID_KERNEL:
+      return "CL_INVALID_KERNEL";
+    case CL_INVALID_ARG_INDEX:
+      return "CL_INVALID_ARG_INDEX";
+    case CL_INVALID_ARG_VALUE:
+      return "CL_INVALID_ARG_VALUE";
+    case CL_INVALID_ARG_SIZE:
+      return "CL_INVALID_ARG_SIZE";
+    case CL_INVALID_KERNEL_ARGS:
+      return "CL_INVALID_KERNEL_ARGS";
+    case CL_INVALID_WORK_DIMENSION:
+      return "CL_INVALID_WORK_DIMENSION";
+    case CL_INVALID_WORK_GROUP_SIZE:
+      return "CL_INVALID_WORK_GROUP_SIZE";
+    case CL_INVALID_WORK_ITEM_SIZE:
+      return "CL_INVALID_WORK_ITEM_SIZE";
+    case CL_INVALID_GLOBAL_OFFSET:
+      return "CL_INVALID_GLOBAL_OFFSET";
+    case CL_INVALID_EVENT_WAIT_LIST:
+      return "CL_INVALID_EVENT_WAIT_LIST";
+    case CL_INVALID_EVENT:
+      return "CL_INVALID_EVENT";
+    case CL_INVALID_OPERATION:
+      return "CL_INVALID_OPERATION";
+    case CL_INVALID_GL_OBJECT:
+      return "CL_INVALID_GL_OBJECT";
+    case CL_INVALID_BUFFER_SIZE:
+      return "CL_INVALID_BUFFER_SIZE";
+    case CL_INVALID_MIP_LEVEL:
+      return "CL_INVALID_MIP_LEVEL";
+#ifndef CL_PLATFORM_NVIDIA
+    case CL_INVALID_GLOBAL_WORK_SIZE:
+      return "CL_INVALID_GLOBAL_WORK_SIZE";
+#endif
+    default:
+      return "Unknown";
+    };
+}
+
+
+void print_clinfo ()
+{
+  char *s = NULL;
+  size_t len;
+  unsigned i, j;
+  cl_uint platform_count;
+  cl_platform_id *platforms;
+
+  /* Determine number of OpenCL Platforms available.  */
+  clGetPlatformIDs (0, NULL, &platform_count);
+  printf ("number of OpenCL Platforms available:\t%d\n", platform_count);
+  /* Get platforms.  */
+  platforms
+    = (cl_platform_id*) malloc (sizeof (cl_platform_id) * platform_count);
+  if (platforms == NULL)
+    {
+      fprintf (stderr, "malloc failed\n");
+      exit (EXIT_FAILURE);
+    }
+  clGetPlatformIDs (platform_count, platforms, NULL);
+
+  /* Querying platforms.  */
+  for (i = 0; i < platform_count; i++)
+    {
+      cl_device_id *devices;
+      cl_uint device_count;
+      cl_device_id default_dev;
+      printf (" OpenCL Platform:                       %d\n", i);
+
+#define PRINT_PF_INFO(PARM)\
+      clGetPlatformInfo (platforms[i], PARM, 0, NULL, &len); \
+      s = realloc (s, len); \
+      clGetPlatformInfo (platforms[i], PARM, len, s, NULL); \
+      printf ("  %-36s%s\n", #PARM ":", s);
+
+      PRINT_PF_INFO (CL_PLATFORM_PROFILE)
+      PRINT_PF_INFO (CL_PLATFORM_VERSION)
+      PRINT_PF_INFO (CL_PLATFORM_NAME)
+      PRINT_PF_INFO (CL_PLATFORM_VENDOR)
+      PRINT_PF_INFO (CL_PLATFORM_EXTENSIONS)
+#undef PRINT_PF_INFO
+
+      clGetDeviceIDs (platforms[i], CL_DEVICE_TYPE_DEFAULT, 1, &default_dev,
+		      NULL);
+      clGetDeviceInfo (default_dev, CL_DEVICE_NAME, 0, NULL, &len);
+      s = realloc (s, len);
+      clGetDeviceInfo (default_dev, CL_DEVICE_NAME, len, s, NULL);
+      printf ("  CL_DEVICE_TYPE_DEFAULT:             %s\n", s);
+
+      /* Determine number of devices.  */
+      clGetDeviceIDs (platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &device_count);
+      printf ("\n  number of OpenCL Devices available:   %d\n", device_count);
+      /* Get devices.  */
+      devices = (cl_device_id*) malloc (sizeof (cl_device_id) * device_count);
+      if (devices == NULL)
+	{
+	  fprintf (stderr, "malloc failed\n");
+	  exit (EXIT_FAILURE);
+	}
+      clGetDeviceIDs (platforms[i], CL_DEVICE_TYPE_ALL, device_count, devices,
+		      NULL);
+
+      /* Querying devices.  */
+      for (j = 0; j < device_count; j++)
+	{
+	  cl_device_type dtype;
+	  cl_device_mem_cache_type mctype;
+	  cl_device_local_mem_type mtype;
+	  cl_device_fp_config fpcfg;
+	  cl_device_exec_capabilities xcap;
+	  cl_command_queue_properties qprops;
+	  cl_bool clbool;
+	  cl_uint cluint;
+	  cl_ulong clulong;
+	  size_t sizet;
+	  size_t workitem_size[3];
+	  printf ("   OpenCL Device:                       %d\n", j);
+
+#define PRINT_DEV_INFO(PARM)\
+	  clGetDeviceInfo (devices[j], PARM, 0, NULL, &len); \
+	  s = realloc (s, len); \
+	  clGetDeviceInfo (devices[j], PARM, len, s, NULL); \
+	  printf ("    %-41s%s\n", #PARM ":", s);
+
+	  PRINT_DEV_INFO (CL_DEVICE_NAME)
+	  PRINT_DEV_INFO (CL_DRIVER_VERSION)
+	  PRINT_DEV_INFO (CL_DEVICE_VENDOR)
+	  clGetDeviceInfo (devices[j], CL_DEVICE_VENDOR_ID, sizeof (cluint),
+			   &cluint, NULL);
+	  printf ("    CL_DEVICE_VENDOR_ID:                     %d\n", cluint);
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_TYPE, sizeof (dtype), &dtype, NULL);
+	  if (dtype & CL_DEVICE_TYPE_CPU)
+	    printf ("    CL_DEVICE_TYPE:                          CL_DEVICE_TYPE_CPU\n");
+	  if (dtype & CL_DEVICE_TYPE_GPU)
+	    printf ("    CL_DEVICE_TYPE:                          CL_DEVICE_TYPE_GPU\n");
+	  if (dtype & CL_DEVICE_TYPE_ACCELERATOR)
+	    printf ("    CL_DEVICE_TYPE:                          CL_DEVICE_TYPE_ACCELERATOR\n");
+	  if (dtype & CL_DEVICE_TYPE_DEFAULT)
+	    printf ("    CL_DEVICE_TYPE:                          CL_DEVICE_TYPE_DEFAULT\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_MAX_CLOCK_FREQUENCY:           %d\n", cluint);
+
+	  PRINT_DEV_INFO (CL_DEVICE_PROFILE)
+	  PRINT_DEV_INFO (CL_DEVICE_EXTENSIONS)
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_AVAILABLE, sizeof (clbool), &clbool, NULL);
+	  if (clbool == CL_TRUE)
+	    printf ("    CL_DEVICE_AVAILABLE:                     CL_TRUE\n");
+	  else
+	    printf ("    CL_DEVICE_AVAILABLE:                     CL_FALSE\n");
+	  clGetDeviceInfo (devices[j], CL_DEVICE_ENDIAN_LITTLE, sizeof (clbool), &clbool, NULL);
+	  if (clbool == CL_TRUE)
+	    printf ("    CL_DEVICE_ENDIAN_LITTLE:                 CL_TRUE\n");
+	  else
+	    printf ("    CL_DEVICE_ENDIAN_LITTLE:                 CL_FALSE\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_MAX_COMPUTE_UNITS:             %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof (sizet), &sizet, NULL);
+	  printf ("    CL_DEVICE_MAX_WORK_GROUP_SIZE:           %d\n", sizet);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:      %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof (workitem_size), &workitem_size, NULL);
+	  printf ("    CL_DEVICE_MAX_WORK_ITEM_SIZES:           %d / %d / %d\n", workitem_size[0], workitem_size[1], workitem_size[2]);
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_ADDRESS_BITS, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_ADDRESS_BITS:                  %d\n", cluint);
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof (clulong), &clulong, NULL);
+	  printf ("    CL_DEVICE_MAX_MEM_ALLOC_SIZE:            %llu\n", clulong);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_MEM_BASE_ADDR_ALIGN:           %d\n", cluint);
+	  clGetDeviceInfo(devices[j], CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:      %d\n", cluint);
+	  clGetDeviceInfo(devices[j], CL_DEVICE_MAX_PARAMETER_SIZE, sizeof (sizet), &sizet, NULL);
+	  printf ("    CL_DEVICE_MAX_PARAMETER_SIZE:            %d\n", sizet);
+	  clGetDeviceInfo(devices[j], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof (clulong), &clulong, NULL);
+	  printf ("    CL_DEVICE_GLOBAL_MEM_SIZE:               %llu\n", clulong);
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, sizeof (mctype), &mctype, NULL);
+	  if (mctype & CL_NONE)
+	    printf ("    CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:         CL_NONE\n");
+	  if (mctype & CL_READ_ONLY_CACHE)
+	    printf ("    CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:         CL_READ_ONLY_CACHE\n");
+	  if (mctype & CL_READ_WRITE_CACHE)
+	    printf ("    CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:         CL_READ_WRITE_CACHE\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, sizeof (clulong), &clulong, NULL);
+	  printf ("    CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:         %llu\n", clulong);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:     %d\n", cluint);
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_LOCAL_MEM_TYPE, sizeof (mtype), &mtype, NULL);
+	  if (mtype & CL_LOCAL)
+	    printf ("    CL_DEVICE_LOCAL_MEM_TYPE:                CL_LOCAL\n");
+	  if (mtype & CL_GLOBAL)
+	    printf ("    CL_DEVICE_LOCAL_MEM_TYPE:                CL_GLOBAL\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:      %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_MEM_BASE_ADDR_ALIGN:           %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:   %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:  %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:    %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:   %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:  %d\n", cluint);
+	  clGetDeviceInfo (devices[j], CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, sizeof (cluint), &cluint, NULL);
+	  printf ("    CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE: %d\n", cluint);
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_SINGLE_FP_CONFIG, sizeof (fpcfg), &fpcfg, NULL);
+	  if (fpcfg & CL_FP_DENORM)
+	    printf ("    CL_DEVICE_SINGLE_FP_CONFIG:              CL_FP_DENORM\n");
+	  if (fpcfg & CL_FP_INF_NAN)
+	    printf ("    CL_DEVICE_SINGLE_FP_CONFIG:              CL_FP_INF_NAN\n");
+	  if (fpcfg & CL_FP_ROUND_TO_NEAREST)
+	    printf ("    CL_DEVICE_SINGLE_FP_CONFIG:              CL_FP_ROUND_TO_NEAREST\n");
+	  if (fpcfg & CL_FP_ROUND_TO_ZERO)
+	    printf ("    CL_DEVICE_SINGLE_FP_CONFIG:              CL_FP_ROUND_TO_ZERO\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_EXECUTION_CAPABILITIES, sizeof (xcap), &xcap, NULL);
+	  if (xcap & CL_EXEC_KERNEL )
+	    printf ("    CL_DEVICE_EXECUTION_CAPABILITIES:        CL_EXEC_KERNEL\n");
+	  if (xcap & CL_EXEC_NATIVE_KERNEL)
+	    printf ("    CL_DEVICE_EXECUTION_CAPABILITIES:        CL_EXEC_NATIVE_KERNEL\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_QUEUE_PROPERTIES, sizeof (qprops), &qprops, NULL);
+	  if (qprops & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
+	    printf ("    CL_DEVICE_QUEUE_PROPERTIES:              CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE\n");
+	  if (qprops & CL_QUEUE_PROFILING_ENABLE)
+	    printf ("    CL_DEVICE_QUEUE_PROPERTIES:              CL_QUEUE_PROFILING_ENABLE\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_PROFILING_TIMER_RESOLUTION, sizeof (sizet), &sizet, NULL);
+	  printf ("    CL_DEVICE_PROFILING_TIMER_RESOLUTION:    %d\n", sizet);
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_COMPILER_AVAILABLE, sizeof (clbool), &clbool, NULL);
+	  if (clbool == CL_TRUE)
+	    printf ("    CL_DEVICE_COMPILER_AVAILABLE:            CL_TRUE\n");
+	  else
+	    printf ("    CL_DEVICE_COMPILER_AVAILABLE:            CL_FALSE\n");
+	  clGetDeviceInfo (devices[j], CL_DEVICE_ERROR_CORRECTION_SUPPORT, sizeof (clbool), &clbool, NULL);
+	  if (clbool == CL_TRUE)
+	    printf ("    CL_DEVICE_ERROR_CORRECTION_SUPPORT:      CL_TRUE\n");
+	  else
+	    printf ("    CL_DEVICE_ERROR_CORRECTION_SUPPORT:      CL_FALSE\n");
+
+	  clGetDeviceInfo (devices[j], CL_DEVICE_IMAGE_SUPPORT, sizeof (clbool), &clbool, NULL);
+	  if (clbool == CL_FALSE)
+	    {
+	      printf ("    CL_DEVICE_IMAGE_SUPPORT:                 CL_FALSE\n");
+	    }
+	  else
+	    {
+	      printf ("    CL_DEVICE_IMAGE_SUPPORT:                 CL_TRUE\n");
+	      clGetDeviceInfo (devices[j], CL_DEVICE_MAX_SAMPLERS, sizeof (cluint), &cluint, NULL);
+	      printf ("    CL_DEVICE_MAX_SAMPLERS:                  %d\n", cluint);
+	      clGetDeviceInfo (devices[j], CL_DEVICE_MAX_READ_IMAGE_ARGS, sizeof (cluint), &cluint, NULL);
+	      printf ("    CL_DEVICE_MAX_READ_IMAGE_ARGS:           %d\n", cluint);
+	      clGetDeviceInfo (devices[j], CL_DEVICE_MAX_WRITE_IMAGE_ARGS, sizeof (cluint), &cluint, NULL);
+	      printf ("    CL_DEVICE_MAX_WRITE_IMAGE_ARGS:          %d\n", cluint);
+	      clGetDeviceInfo (devices[j], CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof (sizet), &sizet, NULL);
+	      printf ("    CL_DEVICE_IMAGE2D_MAX_WIDTH:             %d\n", sizet);
+	      clGetDeviceInfo (devices[j], CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof (sizet), &sizet, NULL);
+	      printf ("    CL_DEVICE_IMAGE2D_MAX_HEIGHT:            %d\n", sizet);
+	      clGetDeviceInfo (devices[j], CL_DEVICE_IMAGE3D_MAX_WIDTH, sizeof (sizet), &sizet, NULL);
+	      printf ("    CL_DEVICE_IMAGE3D_MAX_WIDTH:             %d\n", sizet);
+	      clGetDeviceInfo (devices[j], CL_DEVICE_IMAGE3D_MAX_HEIGHT, sizeof (sizet), &sizet, NULL);
+	      printf ("    CL_DEVICE_IMAGE3D_MAX_HEIGHT:            %d\n", sizet);
+	      clGetDeviceInfo (devices[j], CL_DEVICE_IMAGE3D_MAX_DEPTH, sizeof (sizet), &sizet, NULL);
+	      printf ("    CL_DEVICE_IMAGE3D_MAX_DEPTH:             %d\n", sizet);
+	    }
+#undef PRINT_DEV_INFO
+	} /* devices */
+      free (devices);
+    } /* platforms */
+  free (s);
+  free (platforms);
+}
+
+
+const char *
+read_file (const char * const filename, size_t *size)
+{
+  char *buf = NULL;
+  FILE *fd;
+  struct stat st;
+  if (stat (filename, &st) == -1)
+    {
+      /* Check if the file exists.  */
+      if (errno == ENOENT)
+	return buf;
+      perror ("stat failed");
+      exit (EXIT_FAILURE);
+    }
+  buf = (char *) malloc (st.st_size);
+  if (buf == NULL)
+    {
+      fprintf (stderr, "malloc failed\n");
+      exit (EXIT_FAILURE);
+    }
+  fd = fopen (filename, "r");
+  if (fd == NULL)
+    {
+      perror ("fopen failed");
+      free (buf);
+      exit (EXIT_FAILURE);
+    }
+  if (fread (buf, st.st_size, 1, fd) != 1)
+    {
+      fprintf (stderr, "fread failed\n");
+      free (buf);
+      fclose (fd);
+      exit (EXIT_FAILURE);
+    }
+  fclose (fd);
+  *size = st.st_size;
+  return buf;
+}
+
+
+void
+save_program_binaries (cl_program program)
+{
+  cl_device_id *devices;
+  cl_uint device_count;
+  size_t *sizes;
+  unsigned char **binaries;
+  unsigned i, j;
+
+  /* Query the amount of devices for the given program.  */
+  CHK (clGetProgramInfo (program, CL_PROGRAM_NUM_DEVICES, sizeof (cl_uint),
+			&device_count, NULL));
+
+  /* Get the sizes of the binaries.  */
+  sizes = (size_t*) malloc (sizeof (size_t) * device_count);
+  if (sizes == NULL)
+    {
+      fprintf (stderr, "malloc failed\n");
+      exit (EXIT_FAILURE);
+    }
+  CHK (clGetProgramInfo (program, CL_PROGRAM_BINARY_SIZES, sizeof (sizes),
+			 sizes, NULL));
+
+  /* Get the binaries.  */
+  binaries
+    = (unsigned char **) malloc (sizeof (unsigned char *) * device_count);
+  if (binaries == NULL)
+    {
+      fprintf (stderr, "malloc failed\n");
+      exit (EXIT_FAILURE);
+    }
+  for (i = 0; i < device_count; i++)
+    {
+      binaries[i] = (unsigned char *) malloc (sizes[i]);
+      if (binaries[i] == NULL)
+	{
+	  fprintf (stderr, "malloc failed\n");
+	  exit (EXIT_FAILURE);
+	}
+    }
+  CHK (clGetProgramInfo (program, CL_PROGRAM_BINARIES, sizeof (binaries),
+			 binaries, NULL));
+
+  /* Get the devices for the given program to extract the file names.  */
+  devices = (cl_device_id*) malloc (sizeof (cl_device_id) * device_count);
+  if (devices == NULL)
+    {
+      fprintf (stderr, "malloc failed\n");
+      exit (EXIT_FAILURE);
+    }
+  CHK (clGetProgramInfo (program, CL_PROGRAM_DEVICES, sizeof (devices),
+			 devices, NULL));
+
+  for (i = 0; i < device_count; i++)
+    {
+      FILE *fd;
+      char *dev_name = NULL;
+      size_t len;
+      CHK (clGetDeviceInfo (devices[i], CL_DEVICE_NAME, 0, NULL, &len));
+      dev_name = malloc (len);
+      if (dev_name == NULL)
+	{
+	  fprintf (stderr, "malloc failed\n");
+	  exit (EXIT_FAILURE);
+	}
+      CHK (clGetDeviceInfo (devices[i], CL_DEVICE_NAME, len, dev_name, NULL));
+      /* Convert spaces to underscores.  */
+      for (j = 0; j < strlen (dev_name); j++)
+	{
+	  if (dev_name[j] == ' ')
+	    dev_name[j] = '_';
+	}
+
+      /*  Save the binaries.  */
+      printf ("saving program binary for device: %s\n", dev_name);
+      /* Save binaries[i].  */
+      fd = fopen (dev_name, "w");
+      if (fd == NULL)
+	{
+	  perror ("fopen failed");
+	  exit (EXIT_FAILURE);
+	}
+      if (fwrite (binaries[i], sizes[i], 1, fd) != 1)
+	{
+	  fprintf (stderr, "fwrite failed\n");
+	  for (j = i; j < device_count; j++)
+	    free (binaries[j]);
+	  fclose (fd);
+	  exit (EXIT_FAILURE);
+	}
+      fclose (fd);
+      free (binaries[i]);
+      free (dev_name);
+      free (sizes);
+    }
+  free (devices);
+  free (binaries);
+}
diff --git a/gdb/testsuite/lib/cl_util.h b/gdb/testsuite/lib/cl_util.h
new file mode 100644
index 00000000000..acdbc5d137f
--- /dev/null
+++ b/gdb/testsuite/lib/cl_util.h
@@ -0,0 +1,88 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2010 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>  */
+
+/* Utility macros and functions for OpenCL applications.  */
+
+#ifndef CL_UTIL_H
+#define CL_UTIL_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+#include <OpenCL/opencl.h>
+#else
+#include <CL/cl.h>
+#endif
+#include <stdio.h>
+
+/* Executes the given OpenCL function and checks its return value.
+   In case of failure (rc != CL_SUCCESS) an error string will be
+   printed to stderr and the program will be terminated.  This Macro
+   is only intended for OpenCL routines which return cl_int.  */
+
+#define CHK(func)\
+{\
+  int rc = (func);\
+  CHK_ERR (#func, rc);\
+}
+
+/* Macro that checks an OpenCL error code.  In case of failure
+   (err != CL_SUCCESS) an error string will be printed to stderr
+   including the prefix and the program will be terminated.  This
+   Macro is only intended to use in conjunction with OpenCL routines
+   which take a pointer to a cl_int as an argument to place their
+   error code.  */
+
+#define CHK_ERR(prefix, err)\
+if (err != CL_SUCCESS)\
+  {\
+    fprintf (stderr, "CHK_ERR (%s, %d)\n", prefix, err);\
+    fprintf (stderr, "%s:%d error: %s\n", __FILE__, __LINE__,\
+	     get_clerror_string (err));\
+    exit (EXIT_FAILURE);\
+  };
+
+/* Return a pointer to a string that describes the error code specified
+   by the errcode argument.  */
+
+extern const char *get_clerror_string (int errcode);
+
+/* Prints OpenCL information to stdout.  */
+
+extern void print_clinfo ();
+
+/* Reads a given file into the memory and returns a pointer to the data or NULL
+   if the file does not exist.  FILENAME specifies the location of the file to
+   be read.  SIZE is an output parameter that returns the size of the  file in
+   bytes.  */
+
+extern const char *read_file (const char * const filename, size_t *size);
+
+/* Saves all program binaries of the given OpenCL PROGRAM.  The file
+   names are extracted from the devices.  */
+
+extern void save_program_binaries (cl_program program);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif /* CL_UTIL_H */
diff --git a/gdb/testsuite/lib/opencl.exp b/gdb/testsuite/lib/opencl.exp
new file mode 100644
index 00000000000..33d36886c0a
--- /dev/null
+++ b/gdb/testsuite/lib/opencl.exp
@@ -0,0 +1,83 @@
+# Copyright 2010 Free Software Foundation, Inc.
+#
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.
+#
+# Contributed by Ken Werner <ken.werner@de.ibm.com>.
+#
+# Support library for testing OpenCL GDB features
+
+# Compile OpenCL programs using a generic host app.
+proc gdb_compile_opencl_hostapp {clsource executable options} {
+    global srcdir objdir subdir
+    set src "${srcdir}/lib/cl_util.c ${srcdir}/lib/opencl_hostapp.c"
+    set binfile ${objdir}/${subdir}/${executable}
+    set compile_flags [concat additional_flags=-I${srcdir}/lib/ additional_flags=-DCL_SOURCE=$clsource]
+    set options_opencl [concat {debug} $compile_flags $options [list libs=-lOpenCL]]
+    return [gdb_compile ${src} ${binfile} "executable" ${options_opencl}]
+}
+
+# Run a test on the target to check if it supports OpenCL. Return 0 if so, 1 if
+# it does not.
+proc skip_opencl_tests {} {
+    global skip_opencl_tests_saved srcdir objdir subdir gdb_prompt
+
+    # Use the cached value, if it exists.  Cache value per "board" to handle
+    # runs with multiple options (e.g. unix/{-m32,-64}) correctly.
+    set me "skip_opencl_tests"
+    set board [target_info name]
+    if [info exists skip_opencl_tests_saved($board)] {
+        verbose "$me:  returning saved $skip_opencl_tests_saved($board)" 2
+        return $skip_opencl_tests_saved($board)
+    }
+
+    # Set up, compile, and execute an OpenCL program.  Include the current
+    # process ID in the file name of the executable to prevent conflicts with
+    # invocations for multiple testsuites.
+    set clprogram [remote_download target ${srcdir}/lib/opencl_kernel.cl]
+    set executable opencltest[pid].x
+
+    verbose "$me:  compiling OpenCL test app" 2
+    set compile_flags {debug nowarnings quiet}
+
+    if { [gdb_compile_opencl_hostapp "${clprogram}" "${executable}" "" ] != "" } {
+        verbose "$me:  compiling OpenCL binary failed, returning 1" 2
+	return [set skip_opencl_tests_saved($board) 1]
+    }
+
+    # Compilation succeeded so now run it via gdb.
+    clean_restart "$executable"
+    gdb_run_cmd
+    gdb_expect 30 {
+        -re ".*Program exited normally.*${gdb_prompt} $" {
+            verbose -log "\n$me: OpenCL support detected"
+            set skip_opencl_tests_saved($board) 0
+        }
+        -re ".*Program exited with code.*${gdb_prompt} $" {
+            verbose -log "\n$me: OpenCL support not detected"
+            set skip_opencl_tests_saved($board) 1
+        }
+        default {
+            verbose -log "\n$me OpenCL support not detected (default case)"
+            set skip_opencl_tests_saved($board) 1
+        }
+    }
+    gdb_exit
+    remote_file build delete $executable
+
+    # Delete the OpenCL program source file.
+    remote_file target delete ${clprogram}
+
+    verbose "$me:  returning $skip_opencl_tests_saved($board)" 2
+    return $skip_opencl_tests_saved($board)
+}
diff --git a/gdb/testsuite/lib/opencl_hostapp.c b/gdb/testsuite/lib/opencl_hostapp.c
new file mode 100644
index 00000000000..4bc965805b0
--- /dev/null
+++ b/gdb/testsuite/lib/opencl_hostapp.c
@@ -0,0 +1,168 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2010 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+   Contributed by Ken Werner <ken.werner@de.ibm.com>  */
+
+/* Simple OpenCL application that executes a kernel on the default device
+   in a data parallel fashion.  The filename of the OpenCL program source
+   should be specified using the CL_SOURCE define.  The name of the kernel
+   routine is expected to be "testkernel".  */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+#include <CL/cl.h>
+#include "cl_util.h"
+
+#ifndef CL_SOURCE
+#error "Please specify the OpenCL source file using the CL_SOURCE define"
+#endif
+
+#define STRINGIFY(S) _STRINGIFY(S)
+#define _STRINGIFY(S) #S
+
+#define SIZE 16
+
+int
+main ()
+{
+  int err, i;
+  cl_platform_id platform;
+  cl_device_id device;
+  cl_context context;
+  cl_context_properties context_props[3];
+  cl_command_queue queue;
+  cl_program program;
+  cl_kernel kernel;
+  cl_mem buffer;
+
+  size_t len;
+  const char *program_source = NULL;
+  char *device_extensions = NULL;
+  char kernel_build_opts[256];
+  size_t size = sizeof (cl_int) * SIZE;
+  const size_t global_work_size[] = {SIZE, 0, 0}; /* size of each dimension */
+  cl_int *data;
+
+  /* In order to see which devices the OpenCL implementation on your platform
+     provides you may issue a call to the print_clinfo () fuction.  */
+
+  /* Initialize the data the OpenCl program operates on.  */
+  data = (cl_int*) calloc (1, size);
+  if (data == NULL)
+    {
+      fprintf (stderr, "calloc failed\n");
+      exit (EXIT_FAILURE);
+    }
+
+  /* Pick the first platform.  */
+  CHK (clGetPlatformIDs (1, &platform, NULL));
+  /* Get the default device and create context.  */
+  CHK (clGetDeviceIDs (platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, NULL));
+  context_props[0] = CL_CONTEXT_PLATFORM;
+  context_props[1] = (cl_context_properties) platform;
+  context_props[2] = 0;
+  context = clCreateContext (context_props, 1, &device, NULL, NULL, &err);
+  CHK_ERR ("clCreateContext", err);
+  queue = clCreateCommandQueue (context, device, 0, &err);
+  CHK_ERR ("clCreateCommandQueue", err);
+
+  /* Query OpenCL extensions of that device.  */
+  CHK (clGetDeviceInfo (device, CL_DEVICE_EXTENSIONS, 0, NULL, &len));
+  device_extensions = (char *) malloc (len);
+  CHK (clGetDeviceInfo (device, CL_DEVICE_EXTENSIONS, len, device_extensions,
+			NULL));
+  strcpy (kernel_build_opts, "-Werror -cl-opt-disable");
+  if (strstr (device_extensions, "cl_khr_fp64") != NULL)
+    strcpy (kernel_build_opts + strlen (kernel_build_opts),
+	    " -D HAVE_cl_khr_fp64");
+  if (strstr (device_extensions, "cl_khr_fp16") != NULL)
+    strcpy (kernel_build_opts + strlen (kernel_build_opts),
+	    " -D HAVE_cl_khr_fp16");
+
+  /* Read the OpenCL kernel source into the main memory.  */
+  program_source = read_file (STRINGIFY (CL_SOURCE), &len);
+  if (program_source == NULL)
+    {
+      fprintf (stderr, "file does not exist: %s\n", STRINGIFY (CL_SOURCE));
+      exit (EXIT_FAILURE);
+    }
+
+  /* Build the OpenCL kernel.  */
+  program = clCreateProgramWithSource (context, 1, &program_source,
+				       &len, &err);
+  free ((void*) program_source);
+  CHK_ERR ("clCreateProgramWithSource", err);
+  err = clBuildProgram (program, 0, NULL, kernel_build_opts, NULL,
+			NULL);
+  if (err != CL_SUCCESS)
+    {
+      size_t len;
+      char *clbuild_log = NULL;
+      CHK (clGetProgramBuildInfo (program, device, CL_PROGRAM_BUILD_LOG, 0,
+				  NULL, &len));
+      clbuild_log = malloc (len);
+      if (clbuild_log)
+	{
+	  CHK (clGetProgramBuildInfo (program, device, CL_PROGRAM_BUILD_LOG,
+				      len, clbuild_log, NULL));
+	  fprintf (stderr, "clBuildProgram failed with:\n%s\n", clbuild_log);
+ 	  free (clbuild_log);
+        }
+      exit (EXIT_FAILURE);
+  }
+
+  /* In some cases it might be handy to save the OpenCL program binaries to do
+     further analysis on them.  In order to do so you may call the following
+     function: save_program_binaries (program);.  */
+
+  kernel = clCreateKernel (program, "testkernel", &err);
+  CHK_ERR ("clCreateKernel", err);
+
+  /* Setup the input data for the kernel.  */
+  buffer = clCreateBuffer (context, CL_MEM_USE_HOST_PTR, size, data, &err);
+  CHK_ERR ("clCreateBuffer", err);
+
+  /* Execute the kernel (data parallel).  */
+  CHK (clSetKernelArg (kernel, 0, sizeof (buffer), &buffer));
+  CHK (clEnqueueNDRangeKernel (queue, kernel, 1, NULL, global_work_size, NULL,
+			       0, NULL, NULL));
+
+  /* Fetch the results (blocking).  */
+  CHK (clEnqueueReadBuffer (queue, buffer, CL_TRUE, 0, size, data, 0, NULL,
+			    NULL));
+
+  /* Compare the results.  */
+  for (i = 0; i < SIZE; i++)
+    {
+      if (data[i] != 0x1)
+	{
+	  fprintf (stderr, "error: data[%d]: %d != 0x1\n", i, data[i]);
+	  exit (EXIT_FAILURE);
+	}
+    }
+
+  /* Cleanup.  */
+  CHK (clReleaseMemObject (buffer));
+  CHK (clReleaseKernel (kernel));
+  CHK (clReleaseProgram (program));
+  CHK (clReleaseCommandQueue (queue));
+  CHK (clReleaseContext (context));
+  free (data);
+
+  return 0;
+}
diff --git a/gdb/testsuite/lib/opencl_kernel.cl b/gdb/testsuite/lib/opencl_kernel.cl
new file mode 100644
index 00000000000..32cba64a0ff
--- /dev/null
+++ b/gdb/testsuite/lib/opencl_kernel.cl
@@ -0,0 +1,5 @@
+/* OpenCL kernel for testing purposes.  */
+__kernel void testkernel (__global int *data)
+{
+  data[get_global_id(0)] = 0x1;
+}