[rfc] Implement support for IBM XL C for OpenCL vector ABI

Ulrich Weigand
2011-02-02T17:48:28Z

Hello,

code generated for OpenCL C kernels does not necessarily need to
adhere to a platform-defined ABI, since OpenCL does not allow to
link binary components together. However, since GDB allows for
inferior function calls to routines defined as part of an OpenCL
C kernel, it needs to understand the de-facto ABI used on any
given implementation.

With the IBM XL C for OpenCL compiler, we mostly use the existing
platform ABI for the PowerPC and SPU architectures. However, the
OpenCL C language defines a large set of vector types that do not
correspond to any of the pre-existing data types. For those, the
compiler chose to implement an ABI. GDB's current implementation
does not always match this ABI.

The following patch implements support for the actually implemented
ABI in OpenCL C for PowerPC and SPU. To do so, we need to actually
know whether any given function uses the OpenCL C ABI (as opposed
to the regular platform ABI). Ideally, we'd want to know if the
inferior function to be called originates in an OpenCL C source
file compiled with the IBM XL compiler, but this information is
no longer directly available in the push_dummy_call etc. callbacks.

What *is* available is the TYPE_CALLING_CONVENTION attribute. However,
this is determined from DWARF DW_AT_calling_convention attributes,
which the OpenCL compiler does not actually set. To work around this,
the patch below hard-codes a special flag to be used as value of
TYPE_CALLING_CONVENTION, which is set depending on the compiler
that built the source file (i.e. DWARF "producer").

This extra flag is defined by GDB itself, and has a value outside
the defined range of DW_AT_calling_convention attribute values,
so there should be no potential conflict.

Does this look reasonable? If anyone sees a better way to implement
this, I'd appreciate any suggestions ...

Using this value, the patch below then implements the OpenCL ABI
for both PowerPC (32-bit and 64-bit) and SPU, both for function
calls and function return.

Tested on powerpc64-linux and Cell/B.E. using the IBM XL C for
OpenCL compiler and OpenCL runtime.

Note that this patch assumes the PowerPC AltiVec ABI fix here:
http://sourceware.org/ml/gdb-patches/2011-02/msg00021.html
is already applied.

Any comments welcome! I'm planning on committing this in a
week or so.

Bye,
Ulrich

ChangeLog:

* gdbtypes.h (DW_CC_GDB_IBM_OpenCL): Define.
* dwarf2read.c (read_subroutine_type): Set special calling
convention flag for functions compiled by IBM XL C for OpenCL.
* ppc-sysv-tdep.c (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 (spu_return_value): Likewise.

testsuite/ChangeLog:

* gdb.opencl/callfuncs.cl: New file.
* gdb.opencl/callfuncs.exp: New test.
* gdb.opencl/Makefile.in (EXECUTABLES): Add callfuncs.

diff -urNp gdb-orig/gdb/dwarf2read.c gdb-head/gdb/dwarf2read.c
--- gdb-orig/gdb/dwarf2read.c 2011-02-01 15:46:45.000000000 +0100
+++ gdb-head/gdb/dwarf2read.c 2011-02-01 19:48:02.000000000 +0100
@@ -7907,7 +7907,12 @@ read_subroutine_type (struct die_info *d
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 -urNp gdb-orig/gdb/gdbtypes.h gdb-head/gdb/gdbtypes.h
--- gdb-orig/gdb/gdbtypes.h 2011-02-01 15:46:45.000000000 +0100
+++ gdb-head/gdb/gdbtypes.h 2011-02-01 19:48:02.000000000 +0100
@@ -597,6 +597,12 @@ struct main_type
supporting multiple ABIs. Right now this is only fetched from
the Dwarf-2 DW_AT_calling_convention attribute. */
unsigned calling_convention;
+ /* GDB uses the following values, in addition to the DW_CC_... values
+ defined in include/dwarf2.h, to describe implicitly derived extra
+ calling conventions. These must not overlap the range defined by
+ DWARF for DW_CC_... values. */
+#define DW_CC_GDB_IBM_OpenCL 0x1000
+
} type_specific;
};

diff -urNp gdb-orig/gdb/ppc-sysv-tdep.c gdb-head/gdb/ppc-sysv-tdep.c
--- gdb-orig/gdb/ppc-sysv-tdep.c 2011-02-01 18:27:46.000000000 +0100
+++ gdb-head/gdb/ppc-sysv-tdep.c 2011-02-01 19:49:42.000000000 +0100
@@ -50,6 +50,8 @@ ppc_sysv_abi_push_dummy_call (struct gdb
{
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 +61,13 @@ ppc_sysv_abi_push_dummy_call (struct gdb
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 +336,126 @@ ppc_sysv_abi_push_dummy_call (struct gdb
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 +681,21 @@ get_decimal_float_return_value (struct g
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 +828,83 @@ do_ppc_sysv_return_value (struct gdbarch
}
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, &regval);
+ 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 +1040,8 @@ ppc_sysv_abi_return_value (struct gdbarc
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 +1051,8 @@ ppc_sysv_abi_broken_return_value (struct
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 +1113,8 @@ ppc64_sysv_abi_push_dummy_call (struct g
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 +1141,13 @@ ppc64_sysv_abi_push_dummy_call (struct g
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 +1356,109 @@ ppc64_sysv_abi_push_dummy_call (struct g
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 +1684,11 @@ ppc64_sysv_abi_return_value (struct gdba
{
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 +1751,83 @@ ppc64_sysv_abi_return_value (struct gdba
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, &regval);
+ 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 -urNp gdb-orig/gdb/spu-tdep.c gdb-head/gdb/spu-tdep.c
--- gdb-orig/gdb/spu-tdep.c 2011-02-01 15:46:45.000000000 +0100
+++ gdb-head/gdb/spu-tdep.c 2011-02-01 19:48:02.000000000 +0100
@@ -1448,6 +1448,13 @@ spu_return_value (struct gdbarch *gdbarc
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 +1466,10 @@ spu_return_value (struct gdbarch *gdbarc
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 +1482,10 @@ spu_return_value (struct gdbarch *gdbarc
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 -urNp gdb-orig/gdb/testsuite/gdb.opencl/callfuncs.cl gdb-head/gdb/testsuite/gdb.opencl/callfuncs.cl
--- gdb-orig/gdb/testsuite/gdb.opencl/callfuncs.cl 1970-01-01 01:00:00.000000000 +0100
+++ gdb-head/gdb/testsuite/gdb.opencl/callfuncs.cl 2011-02-01 19:48:02.000000000 +0100
@@ -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 <http://www.gnu.org/licenses/>.
+
+ Contributed by Ulrich Weigand <ulrich.weigand.ibm.com> */
+
+__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 -urNp gdb-orig/gdb/testsuite/gdb.opencl/callfuncs.exp gdb-head/gdb/testsuite/gdb.opencl/callfuncs.exp
--- gdb-orig/gdb/testsuite/gdb.opencl/callfuncs.exp 1970-01-01 01:00:00.000000000 +0100
+++ gdb-head/gdb/testsuite/gdb.opencl/callfuncs.exp 2011-02-01 19:48:02.000000000 +0100
@@ -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 <http://www.gnu.org/licenses/>. */
+#
+# Contributed by Ulrich Weigand <ulrich.weigand@de.ibm.com>.
+#
+# 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 -urNp gdb-orig/gdb/testsuite/gdb.opencl/Makefile.in gdb-head/gdb/testsuite/gdb.opencl/Makefile.in
--- gdb-orig/gdb/testsuite/gdb.opencl/Makefile.in 2011-02-01 15:46:45.000000000 +0100
+++ gdb-head/gdb/testsuite/gdb.opencl/Makefile.in 2011-02-01 19:48:02.000000000 +0100
@@ -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 $@..."
--
Dr. Ulrich Weigand
GNU Toolchain for Linux on System z and Cell BE
Ulrich.Weigand@de.ibm.com