diff options
31 files changed, 4550 insertions, 43 deletions
diff --git a/gdb/ChangeLog b/gdb/ChangeLog index 9ea16e9..793c062 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 568fdb5..550badf 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 \ @@ -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 57e09b3..2e6c371 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 40c4172..015ba16 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 a04fbb2..dc571a4 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); @@ -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 a8df072..63f97ec 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 069dce4..993b0fb 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 a91f14a..404faf8 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; @@ -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 0000000..088d49a --- /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 d68e850..1a2dde5 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 8d8d704..d02689b 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 7b1248a..b523d1b 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 c8668e5..2748108 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 7afa865..b6ecdcb 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 0000000..c12aef3 --- /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 0000000..a024c51 --- /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 0000000..34ea635 --- /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 0000000..c0d2a1e --- /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 0000000..45c9e52 --- /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 0000000..0974c04 --- /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 0000000..b60c65c --- /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 0000000..d58f1ba --- /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 0000000..e044e96 --- /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 0000000..5b731b2 --- /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 0000000..acdbc5d --- /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 0000000..33d3688 --- /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 0000000..4bc9658 --- /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 0000000..32cba64 --- /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; +} |