From 54fcddd0ace6123d9dadc3040c39acd1b42ffd09 Mon Sep 17 00:00:00 2001 From: Ulrich Weigand Date: Tue, 8 Feb 2011 13:30:10 +0000 Subject: [PATCH] include/ChangeLog: * dwarf2.h (enum dwarf_calling_convention): Add DW_CC_GDB_IBM_OpenCL. gdb/ChangeLog: * dwarf2read.c (read_subroutine_type): Set special calling convention flag for functions compiled by IBM XL C for OpenCL. * ppc-sysv-tdep.c: Include "dwarf2.h" (ppc_sysv_abi_push_dummy_call): Implement IBM OpenCL vector types calling convention. (do_ppc_sysv_return_value): Add FUNC_TYPE argument. Implement IBM OpenCL vector types calling convention. (ppc_sysv_abi_return_value): Pass through FUNC_TYPE. (ppc_sysv_abi_broken_return_value): Likewise. (ppc64_sysv_abi_push_dummy_call): Implement IBM OpenCL vector types calling convention. (ppc64_sysv_abi_return_value): Likewise. * spu-tdep.c: Include "dwarf2.h" (spu_return_value): Implement IBM OpenCL vector types calling convention. gdb/testsuite/ChangeLog: * gdb.opencl/callfuncs.cl: New file. * gdb.opencl/callfuncs.exp: New test. * gdb.opencl/Makefile.in (EXECUTABLES): Add callfuncs. --- gdb/ChangeLog | 18 ++ gdb/dwarf2read.c | 7 +- gdb/ppc-sysv-tdep.c | 423 ++++++++++++++++++++++++- gdb/spu-tdep.c | 18 +- gdb/testsuite/ChangeLog | 6 + gdb/testsuite/gdb.opencl/Makefile.in | 2 +- gdb/testsuite/gdb.opencl/callfuncs.cl | 218 +++++++++++++ gdb/testsuite/gdb.opencl/callfuncs.exp | 102 ++++++ include/ChangeLog | 4 + include/dwarf2.h | 10 +- 10 files changed, 796 insertions(+), 12 deletions(-) create mode 100644 gdb/testsuite/gdb.opencl/callfuncs.cl create mode 100644 gdb/testsuite/gdb.opencl/callfuncs.exp diff --git a/gdb/ChangeLog b/gdb/ChangeLog index 0197e8bce1..87f31f341b 100644 --- a/gdb/ChangeLog +++ b/gdb/ChangeLog @@ -1,3 +1,21 @@ +2011-02-08 Ulrich Weigand + + * dwarf2read.c (read_subroutine_type): Set special calling + convention flag for functions compiled by IBM XL C for OpenCL. + * ppc-sysv-tdep.c: Include "dwarf2.h" + (ppc_sysv_abi_push_dummy_call): Implement IBM OpenCL vector types + calling convention. + (do_ppc_sysv_return_value): Add FUNC_TYPE argument. Implement + IBM OpenCL vector types calling convention. + (ppc_sysv_abi_return_value): Pass through FUNC_TYPE. + (ppc_sysv_abi_broken_return_value): Likewise. + (ppc64_sysv_abi_push_dummy_call): Implement IBM OpenCL vector + types calling convention. + (ppc64_sysv_abi_return_value): Likewise. + * spu-tdep.c: Include "dwarf2.h" + (spu_return_value): Implement IBM OpenCL vector types calling + convention. + 2011-02-08 Ulrich Weigand * ppc-sysv-tdep.c (ppc64_sysv_abi_push_dummy_call): Implement diff --git a/gdb/dwarf2read.c b/gdb/dwarf2read.c index f269deebc4..837fa3bc6f 100644 --- a/gdb/dwarf2read.c +++ b/gdb/dwarf2read.c @@ -7955,7 +7955,12 @@ read_subroutine_type (struct die_info *die, struct dwarf2_cu *cu) the subroutine die. Otherwise set the calling convention to the default value DW_CC_normal. */ attr = dwarf2_attr (die, DW_AT_calling_convention, cu); - TYPE_CALLING_CONVENTION (ftype) = attr ? DW_UNSND (attr) : DW_CC_normal; + if (attr) + TYPE_CALLING_CONVENTION (ftype) = DW_UNSND (attr); + else if (cu->producer && strstr (cu->producer, "IBM XL C for OpenCL")) + TYPE_CALLING_CONVENTION (ftype) = DW_CC_GDB_IBM_OpenCL; + else + TYPE_CALLING_CONVENTION (ftype) = DW_CC_normal; /* We need to add the subroutine type to the die immediately so we don't infinitely recurse when dealing with parameters diff --git a/gdb/ppc-sysv-tdep.c b/gdb/ppc-sysv-tdep.c index 300dcac6f1..872117df31 100644 --- a/gdb/ppc-sysv-tdep.c +++ b/gdb/ppc-sysv-tdep.c @@ -30,6 +30,7 @@ #include "target.h" #include "objfiles.h" #include "infcall.h" +#include "dwarf2.h" /* Pass the arguments in either registers, or in the stack. Using the ppc sysv ABI, the first eight words of the argument list (that might @@ -50,6 +51,8 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function, { struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch); enum bfd_endian byte_order = gdbarch_byte_order (gdbarch); + struct type *ftype; + int opencl_abi = 0; ULONGEST saved_sp; int argspace = 0; /* 0 is an initial wrong guess. */ int write_pass; @@ -59,6 +62,13 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function, regcache_cooked_read_unsigned (regcache, gdbarch_sp_regnum (gdbarch), &saved_sp); + ftype = check_typedef (value_type (function)); + if (TYPE_CODE (ftype) == TYPE_CODE_PTR) + ftype = check_typedef (TYPE_TARGET_TYPE (ftype)); + if (TYPE_CODE (ftype) == TYPE_CODE_FUNC + && TYPE_CALLING_CONVENTION (ftype) == DW_CC_GDB_IBM_OpenCL) + opencl_abi = 1; + /* Go through the argument list twice. Pass 1: Figure out how much new stack space is required for @@ -327,6 +337,126 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function, Hence we increase freg even when writing to memory. */ freg += 2; } + else if (len < 16 + && TYPE_CODE (type) == TYPE_CODE_ARRAY + && TYPE_VECTOR (type) + && opencl_abi) + { + /* OpenCL vectors shorter than 16 bytes are passed as if + a series of independent scalars. */ + struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type)); + int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype); + + for (i = 0; i < nelt; i++) + { + const gdb_byte *elval = val + i * TYPE_LENGTH (eltype); + + if (TYPE_CODE (eltype) == TYPE_CODE_FLT && !tdep->soft_float) + { + if (freg <= 8) + { + if (write_pass) + { + int regnum = tdep->ppc_fp0_regnum + freg; + gdb_byte regval[MAX_REGISTER_SIZE]; + struct type *regtype + = register_type (gdbarch, regnum); + convert_typed_floating (elval, eltype, + regval, regtype); + regcache_cooked_write (regcache, regnum, regval); + } + freg++; + } + else + { + argoffset = align_up (argoffset, len); + if (write_pass) + write_memory (sp + argoffset, val, len); + argoffset += len; + } + } + else if (TYPE_LENGTH (eltype) == 8) + { + if (greg > 9) + { + /* Just in case GREG was 10. */ + greg = 11; + argoffset = align_up (argoffset, 8); + if (write_pass) + write_memory (sp + argoffset, elval, + TYPE_LENGTH (eltype)); + argoffset += 8; + } + else + { + /* Must start on an odd register - r3/r4 etc. */ + if ((greg & 1) == 0) + greg++; + if (write_pass) + { + int regnum = tdep->ppc_gp0_regnum + greg; + regcache_cooked_write (regcache, + regnum + 0, elval + 0); + regcache_cooked_write (regcache, + regnum + 1, elval + 4); + } + greg += 2; + } + } + else + { + gdb_byte word[MAX_REGISTER_SIZE]; + store_unsigned_integer (word, tdep->wordsize, byte_order, + unpack_long (eltype, elval)); + + if (greg <= 10) + { + if (write_pass) + regcache_cooked_write (regcache, + tdep->ppc_gp0_regnum + greg, + word); + greg++; + } + else + { + argoffset = align_up (argoffset, tdep->wordsize); + if (write_pass) + write_memory (sp + argoffset, word, tdep->wordsize); + argoffset += tdep->wordsize; + } + } + } + } + else if (len >= 16 + && TYPE_CODE (type) == TYPE_CODE_ARRAY + && TYPE_VECTOR (type) + && opencl_abi) + { + /* OpenCL vectors 16 bytes or longer are passed as if + a series of AltiVec vectors. */ + int i; + + for (i = 0; i < len / 16; i++) + { + const gdb_byte *elval = val + i * 16; + + if (vreg <= 13) + { + if (write_pass) + regcache_cooked_write (regcache, + tdep->ppc_vr0_regnum + vreg, + elval); + vreg++; + } + else + { + argoffset = align_up (argoffset, 16); + if (write_pass) + write_memory (sp + argoffset, elval, 16); + argoffset += 16; + } + } + } else if (len == 16 && TYPE_CODE (type) == TYPE_CODE_ARRAY && TYPE_VECTOR (type) @@ -552,13 +682,21 @@ get_decimal_float_return_value (struct gdbarch *gdbarch, struct type *valtype, when returned in general-purpose registers. */ static enum return_value_convention -do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *type, - struct regcache *regcache, gdb_byte *readbuf, - const gdb_byte *writebuf, int broken_gcc) +do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *func_type, + struct type *type, struct regcache *regcache, + gdb_byte *readbuf, const gdb_byte *writebuf, + int broken_gcc) { struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch); enum bfd_endian byte_order = gdbarch_byte_order (gdbarch); + int opencl_abi = 0; + + if (func_type + && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL) + opencl_abi = 1; + gdb_assert (tdep->wordsize == 4); + if (TYPE_CODE (type) == TYPE_CODE_FLT && TYPE_LENGTH (type) <= 8 && !tdep->soft_float) @@ -691,6 +829,83 @@ do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *type, } return RETURN_VALUE_REGISTER_CONVENTION; } + /* OpenCL vectors < 16 bytes are returned as distinct + scalars in f1..f2 or r3..r10. */ + if (TYPE_CODE (type) == TYPE_CODE_ARRAY + && TYPE_VECTOR (type) + && TYPE_LENGTH (type) < 16 + && opencl_abi) + { + struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type)); + int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype); + + for (i = 0; i < nelt; i++) + { + int offset = i * TYPE_LENGTH (eltype); + + if (TYPE_CODE (eltype) == TYPE_CODE_FLT) + { + int regnum = tdep->ppc_fp0_regnum + 1 + i; + gdb_byte regval[MAX_REGISTER_SIZE]; + struct type *regtype = register_type (gdbarch, regnum); + + if (writebuf != NULL) + { + convert_typed_floating (writebuf + offset, eltype, + regval, regtype); + regcache_cooked_write (regcache, regnum, regval); + } + if (readbuf != NULL) + { + regcache_cooked_read (regcache, regnum, regval); + convert_typed_floating (regval, regtype, + readbuf + offset, eltype); + } + } + else + { + int regnum = tdep->ppc_gp0_regnum + 3 + i; + ULONGEST regval; + + if (writebuf != NULL) + { + regval = unpack_long (eltype, writebuf + offset); + regcache_cooked_write_unsigned (regcache, regnum, regval); + } + if (readbuf != NULL) + { + regcache_cooked_read_unsigned (regcache, regnum, ®val); + store_unsigned_integer (readbuf + offset, + TYPE_LENGTH (eltype), byte_order, + regval); + } + } + } + + return RETURN_VALUE_REGISTER_CONVENTION; + } + /* OpenCL vectors >= 16 bytes are returned in v2..v9. */ + if (TYPE_CODE (type) == TYPE_CODE_ARRAY + && TYPE_VECTOR (type) + && TYPE_LENGTH (type) >= 16 + && opencl_abi) + { + int n_regs = TYPE_LENGTH (type) / 16; + int i; + + for (i = 0; i < n_regs; i++) + { + int offset = i * 16; + int regnum = tdep->ppc_vr0_regnum + 2 + i; + + if (writebuf != NULL) + regcache_cooked_write (regcache, regnum, writebuf + offset); + if (readbuf != NULL) + regcache_cooked_read (regcache, regnum, readbuf + offset); + } + + return RETURN_VALUE_REGISTER_CONVENTION; + } if (TYPE_LENGTH (type) == 16 && TYPE_CODE (type) == TYPE_CODE_ARRAY && TYPE_VECTOR (type) @@ -826,8 +1041,8 @@ ppc_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type, struct type *valtype, struct regcache *regcache, gdb_byte *readbuf, const gdb_byte *writebuf) { - return do_ppc_sysv_return_value (gdbarch, valtype, regcache, readbuf, - writebuf, 0); + return do_ppc_sysv_return_value (gdbarch, func_type, valtype, regcache, + readbuf, writebuf, 0); } enum return_value_convention @@ -837,8 +1052,8 @@ ppc_sysv_abi_broken_return_value (struct gdbarch *gdbarch, struct regcache *regcache, gdb_byte *readbuf, const gdb_byte *writebuf) { - return do_ppc_sysv_return_value (gdbarch, valtype, regcache, readbuf, - writebuf, 1); + return do_ppc_sysv_return_value (gdbarch, func_type, valtype, regcache, + readbuf, writebuf, 1); } /* The helper function for 64-bit SYSV push_dummy_call. Converts the @@ -899,6 +1114,8 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, CORE_ADDR func_addr = find_function_addr (function, NULL); struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch); enum bfd_endian byte_order = gdbarch_byte_order (gdbarch); + struct type *ftype; + int opencl_abi = 0; ULONGEST back_chain; /* See for-loop comment below. */ int write_pass; @@ -925,6 +1142,13 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, regcache_cooked_read_unsigned (regcache, gdbarch_sp_regnum (gdbarch), &back_chain); + ftype = check_typedef (value_type (function)); + if (TYPE_CODE (ftype) == TYPE_CODE_PTR) + ftype = check_typedef (TYPE_TARGET_TYPE (ftype)); + if (TYPE_CODE (ftype) == TYPE_CODE_FUNC + && TYPE_CALLING_CONVENTION (ftype) == DW_CC_GDB_IBM_OpenCL) + opencl_abi = 1; + /* Go through the argument list twice. Pass 1: Compute the function call's stack space and register @@ -1133,6 +1357,109 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, greg += 2; gparam = align_up (gparam + TYPE_LENGTH (type), tdep->wordsize); } + else if (TYPE_LENGTH (type) < 16 + && TYPE_CODE (type) == TYPE_CODE_ARRAY + && TYPE_VECTOR (type) + && opencl_abi) + { + /* OpenCL vectors shorter than 16 bytes are passed as if + a series of independent scalars. */ + struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type)); + int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype); + + for (i = 0; i < nelt; i++) + { + const gdb_byte *elval = val + i * TYPE_LENGTH (eltype); + + if (TYPE_CODE (eltype) == TYPE_CODE_FLT) + { + if (write_pass) + { + gdb_byte regval[MAX_REGISTER_SIZE]; + const gdb_byte *p; + + if (TYPE_LENGTH (eltype) == 4) + { + memcpy (regval, elval, 4); + memcpy (regval + 4, elval, 4); + p = regval; + } + else + p = elval; + + write_memory (gparam, p, 8); + + if (freg <= 13) + { + int regnum = tdep->ppc_fp0_regnum + freg; + struct type *regtype + = register_type (gdbarch, regnum); + + convert_typed_floating (elval, eltype, + regval, regtype); + regcache_cooked_write (regcache, regnum, regval); + } + + if (greg <= 10) + regcache_cooked_write (regcache, + tdep->ppc_gp0_regnum + greg, + regval); + } + + freg++; + greg++; + gparam = align_up (gparam + 8, tdep->wordsize); + } + else + { + if (write_pass) + { + ULONGEST word = unpack_long (eltype, elval); + if (greg <= 10) + regcache_cooked_write_unsigned + (regcache, tdep->ppc_gp0_regnum + greg, word); + + write_memory_unsigned_integer + (gparam, tdep->wordsize, byte_order, word); + } + + greg++; + gparam = align_up (gparam + TYPE_LENGTH (eltype), + tdep->wordsize); + } + } + } + else if (TYPE_LENGTH (type) >= 16 + && TYPE_CODE (type) == TYPE_CODE_ARRAY + && TYPE_VECTOR (type) + && opencl_abi) + { + /* OpenCL vectors 16 bytes or longer are passed as if + a series of AltiVec vectors. */ + int i; + + for (i = 0; i < TYPE_LENGTH (type) / 16; i++) + { + const gdb_byte *elval = val + i * 16; + + gparam = align_up (gparam, 16); + greg += greg & 1; + + if (write_pass) + { + if (vreg <= 13) + regcache_cooked_write (regcache, + tdep->ppc_vr0_regnum + vreg, + elval); + + write_memory (gparam, elval, 16); + } + + greg += 2; + vreg++; + gparam += 16; + } + } else if (TYPE_LENGTH (type) == 16 && TYPE_VECTOR (type) && TYPE_CODE (type) == TYPE_CODE_ARRAY && tdep->ppc_vr0_regnum >= 0) @@ -1358,6 +1685,11 @@ ppc64_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type, { struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch); enum bfd_endian byte_order = gdbarch_byte_order (gdbarch); + int opencl_abi = 0; + + if (func_type + && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL) + opencl_abi = 1; /* This function exists to support a calling convention that requires floating-point registers. It shouldn't be used on @@ -1420,6 +1752,83 @@ ppc64_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type, regcache_cooked_read (regcache, tdep->ppc_gp0_regnum + 3, readbuf); return RETURN_VALUE_REGISTER_CONVENTION; } + /* OpenCL vectors < 16 bytes are returned as distinct + scalars in f1..f2 or r3..r10. */ + if (TYPE_CODE (valtype) == TYPE_CODE_ARRAY + && TYPE_VECTOR (valtype) + && TYPE_LENGTH (valtype) < 16 + && opencl_abi) + { + struct type *eltype = check_typedef (TYPE_TARGET_TYPE (valtype)); + int i, nelt = TYPE_LENGTH (valtype) / TYPE_LENGTH (eltype); + + for (i = 0; i < nelt; i++) + { + int offset = i * TYPE_LENGTH (eltype); + + if (TYPE_CODE (eltype) == TYPE_CODE_FLT) + { + int regnum = tdep->ppc_fp0_regnum + 1 + i; + gdb_byte regval[MAX_REGISTER_SIZE]; + struct type *regtype = register_type (gdbarch, regnum); + + if (writebuf != NULL) + { + convert_typed_floating (writebuf + offset, eltype, + regval, regtype); + regcache_cooked_write (regcache, regnum, regval); + } + if (readbuf != NULL) + { + regcache_cooked_read (regcache, regnum, regval); + convert_typed_floating (regval, regtype, + readbuf + offset, eltype); + } + } + else + { + int regnum = tdep->ppc_gp0_regnum + 3 + i; + ULONGEST regval; + + if (writebuf != NULL) + { + regval = unpack_long (eltype, writebuf + offset); + regcache_cooked_write_unsigned (regcache, regnum, regval); + } + if (readbuf != NULL) + { + regcache_cooked_read_unsigned (regcache, regnum, ®val); + store_unsigned_integer (readbuf + offset, + TYPE_LENGTH (eltype), byte_order, + regval); + } + } + } + + return RETURN_VALUE_REGISTER_CONVENTION; + } + /* OpenCL vectors >= 16 bytes are returned in v2..v9. */ + if (TYPE_CODE (valtype) == TYPE_CODE_ARRAY + && TYPE_VECTOR (valtype) + && TYPE_LENGTH (valtype) >= 16 + && opencl_abi) + { + int n_regs = TYPE_LENGTH (valtype) / 16; + int i; + + for (i = 0; i < n_regs; i++) + { + int offset = i * 16; + int regnum = tdep->ppc_vr0_regnum + 2 + i; + + if (writebuf != NULL) + regcache_cooked_write (regcache, regnum, writebuf + offset); + if (readbuf != NULL) + regcache_cooked_read (regcache, regnum, readbuf + offset); + } + + return RETURN_VALUE_REGISTER_CONVENTION; + } /* Array type has more than one use. */ if (TYPE_CODE (valtype) == TYPE_CODE_ARRAY) { diff --git a/gdb/spu-tdep.c b/gdb/spu-tdep.c index d424659868..0b0ea4ee60 100644 --- a/gdb/spu-tdep.c +++ b/gdb/spu-tdep.c @@ -44,6 +44,7 @@ #include "block.h" #include "observer.h" #include "infcall.h" +#include "dwarf2.h" #include "spu-tdep.h" @@ -1448,6 +1449,13 @@ spu_return_value (struct gdbarch *gdbarch, struct type *func_type, gdb_byte *out, const gdb_byte *in) { enum return_value_convention rvc; + int opencl_vector = 0; + + if (func_type + && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL + && TYPE_CODE (type) == TYPE_CODE_ARRAY + && TYPE_VECTOR (type)) + opencl_vector = 1; if (TYPE_LENGTH (type) <= (SPU_ARGN_REGNUM - SPU_ARG1_REGNUM + 1) * 16) rvc = RETURN_VALUE_REGISTER_CONVENTION; @@ -1459,7 +1467,10 @@ spu_return_value (struct gdbarch *gdbarch, struct type *func_type, switch (rvc) { case RETURN_VALUE_REGISTER_CONVENTION: - spu_value_to_regcache (regcache, SPU_ARG1_REGNUM, type, in); + if (opencl_vector && TYPE_LENGTH (type) == 2) + regcache_cooked_write_part (regcache, SPU_ARG1_REGNUM, 2, 2, in); + else + spu_value_to_regcache (regcache, SPU_ARG1_REGNUM, type, in); break; case RETURN_VALUE_STRUCT_CONVENTION: @@ -1472,7 +1483,10 @@ spu_return_value (struct gdbarch *gdbarch, struct type *func_type, switch (rvc) { case RETURN_VALUE_REGISTER_CONVENTION: - spu_regcache_to_value (regcache, SPU_ARG1_REGNUM, type, out); + if (opencl_vector && TYPE_LENGTH (type) == 2) + regcache_cooked_read_part (regcache, SPU_ARG1_REGNUM, 2, 2, out); + else + spu_regcache_to_value (regcache, SPU_ARG1_REGNUM, type, out); break; case RETURN_VALUE_STRUCT_CONVENTION: diff --git a/gdb/testsuite/ChangeLog b/gdb/testsuite/ChangeLog index 3cb8d41f5e..19aec4c09e 100644 --- a/gdb/testsuite/ChangeLog +++ b/gdb/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2011-02-08 Ulrich Weigand + + * gdb.opencl/callfuncs.cl: New file. + * gdb.opencl/callfuncs.exp: New test. + * gdb.opencl/Makefile.in (EXECUTABLES): Add callfuncs. + 2011-02-08 Ulrich Weigand * gdb.arch/altivec-abi.c (vec_func): Make use of intv_on_stack_f diff --git a/gdb/testsuite/gdb.opencl/Makefile.in b/gdb/testsuite/gdb.opencl/Makefile.in index c12aef3e81..7dec34c2df 100644 --- a/gdb/testsuite/gdb.opencl/Makefile.in +++ b/gdb/testsuite/gdb.opencl/Makefile.in @@ -1,7 +1,7 @@ VPATH = @srcdir@ srcdir = @srcdir@ -EXECUTABLES = datatypes vec_comps convs_casts operators +EXECUTABLES = callfuncs datatypes vec_comps convs_casts operators all info install-info dvi install uninstall installcheck check: @echo "Nothing to be done for $@..." diff --git a/gdb/testsuite/gdb.opencl/callfuncs.cl b/gdb/testsuite/gdb.opencl/callfuncs.cl new file mode 100644 index 0000000000..6d53ee00d6 --- /dev/null +++ b/gdb/testsuite/gdb.opencl/callfuncs.cl @@ -0,0 +1,218 @@ +/* This testcase is part of GDB, the GNU debugger. + + Copyright 2011 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 . + + Contributed by Ulrich Weigand */ + +__constant int opencl_version = __OPENCL_VERSION__; + +#ifdef HAVE_cl_khr_fp64 +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +__constant int have_cl_khr_fp64 = 1; +#else +__constant int have_cl_khr_fp64 = 0; +#endif + +#ifdef HAVE_cl_khr_fp16 +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +__constant int have_cl_khr_fp16 = 1; +#else +__constant int have_cl_khr_fp16 = 0; +#endif + +#define def_call_func(type) \ + type call_##type (type a, type b) { return a + b; } + +#ifdef CL_VERSION_1_1 +#define def_call_family(type) \ + def_call_func(type) \ + def_call_func(type##2) \ + def_call_func(type##3) \ + def_call_func(type##4) \ + def_call_func(type##8) \ + def_call_func(type##16) +#else +#define def_call_family(type) \ + def_call_func(type) \ + def_call_func(type##2) \ + def_call_func(type##4) \ + def_call_func(type##8) \ + def_call_func(type##16) +#endif + +def_call_family(char) +def_call_family(uchar) +def_call_family(short) +def_call_family(ushort) +def_call_family(int) +def_call_family(uint) +def_call_family(long) +def_call_family(ulong) +#ifdef cl_khr_fp16 +def_call_family(half) +#endif +def_call_family(float) +#ifdef cl_khr_fp64 +def_call_family(double) +#endif + +#define call_func(type, var) \ + var = call_##type (var, var); + +#ifdef CL_VERSION_1_1 +#define call_family(type, var) \ + call_func(type, var) \ + call_func(type##2, var##2) \ + call_func(type##3, var##3) \ + call_func(type##4, var##4) \ + call_func(type##8, var##8) \ + call_func(type##16, var##16) +#else +#define call_family(type, var) \ + call_func(type, var) \ + call_func(type##2, var##2) \ + call_func(type##4, var##4) \ + call_func(type##8, var##8) \ + call_func(type##16, var##16) +#endif + +__kernel void testkernel (__global int *data) +{ + 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); + +#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 + + /* marker! */ + + call_family (char, c); + call_family (uchar, uc); + call_family (short, s); + call_family (ushort, us); + call_family (int, i); + call_family (uint, ui); + call_family (long, l); + call_family (ulong, ul); +#ifdef cl_khr_fp16 + call_family (half, h); +#endif + call_family (float, f); +#ifdef cl_khr_fp64 + call_family (double, d); +#endif + + data[get_global_id(0)] = 1; +} diff --git a/gdb/testsuite/gdb.opencl/callfuncs.exp b/gdb/testsuite/gdb.opencl/callfuncs.exp new file mode 100644 index 0000000000..f43558924d --- /dev/null +++ b/gdb/testsuite/gdb.opencl/callfuncs.exp @@ -0,0 +1,102 @@ +# Copyright 2011 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 . */ +# +# Contributed by Ulrich Weigand . +# +# Tests OpenCL function calling conventions. + +if $tracelevel { + strace $tracelevel +} + +load_lib opencl.exp + +if { [skip_opencl_tests] } { + return 0 +} + +set testfile "callfuncs" +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 + +# Load the OpenCL app +gdb_reinitialize_dir $srcdir/$subdir +gdb_load ${objdir}/${subdir}/${testfile} + +# Set breakpoint at the OpenCL kernel +gdb_test "tbreak testkernel" \ + "" \ + "Set pending breakpoint" \ + ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" \ + "y" + +gdb_run_cmd +gdb_test "" ".*reakpoint.*1.*testkernel.*" "run" + +# Continue to the marker +gdb_breakpoint [gdb_get_line_number "marker" "${clprogram}"] +gdb_continue_to_breakpoint "marker" + +# Check if the language was switched to opencl +gdb_test "show language" "The current source language is \"auto; currently opencl\"\." + +# Prevent multi-threaded execution during inferior calls +gdb_test_no_output "set scheduler-locking on" + +# 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 function call / return sequence +proc call_test { type var } { + global opencl_version + + gdb_test "print/d call_${type} (${var}, ${var})" " = 2" + gdb_test "print/d call_${type}2 (${var}2, ${var}2)" " = \\{2, 4\\}" + if { ${opencl_version} >= 110 } { + gdb_test "print/d call_${type}3 (${var}3, ${var}3)" " = \\{2, 4, 6\\}" + } + gdb_test "print/d call_${type}4 (${var}4, ${var}4)" " = \\{2, 4, 6, 8\\}" + gdb_test "print/d call_${type}8 (${var}8, ${var}8)" " = \\{2, 4, 6, 8, 10, 12, 14, 16\\}" + gdb_test "print/d call_${type}16 (${var}16, ${var}16)" " = \\{2, 4, 6, 8, 10, 12, 14, 16, 18, 20, 22, 24, 26, 28, 30, 32\\}" +} + +call_test "char" "c" +call_test "uchar" "uc" +call_test "short" "s" +call_test "ushort" "us" +call_test "int" "i" +call_test "uint" "ui" +call_test "long" "l" +call_test "ulong" "ul" +if { ${have_cl_khr_fp16} } { + call_test "half" "h" +} +call_test "float" "f" +if { ${have_cl_khr_fp64} } { + call_test "double" "d" +} + +# Delete the OpenCL program source +remote_file target delete ${clprogram} diff --git a/include/ChangeLog b/include/ChangeLog index c96d35855b..2a19c6133f 100644 --- a/include/ChangeLog +++ b/include/ChangeLog @@ -1,3 +1,7 @@ +2011-02-08 Ulrich Weigand + + * dwarf2.h (enum dwarf_calling_convention): Add DW_CC_GDB_IBM_OpenCL. + 2011-01-12 Iain Sandoe * dwarf2.h: Update value for DW_AT_hi_user. diff --git a/include/dwarf2.h b/include/dwarf2.h index ad00aac09b..7729ad86dd 100644 --- a/include/dwarf2.h +++ b/include/dwarf2.h @@ -754,7 +754,15 @@ enum dwarf_calling_convention DW_CC_hi_user = 0xff, DW_CC_GNU_renesas_sh = 0x40, - DW_CC_GNU_borland_fastcall_i386 = 0x41 + DW_CC_GNU_borland_fastcall_i386 = 0x41, + + /* This DW_CC_ value is not currently generated by any toolchain. It is + used internally to GDB to indicate OpenCL C functions that have been + compiled with the IBM XL C for OpenCL compiler and use a non-platform + calling convention for passing OpenCL C vector types. This value may + be changed freely as long as it does not conflict with any other DW_CC_ + value defined here. */ + DW_CC_GDB_IBM_OpenCL = 0xff }; /* Inline attribute. */