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

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

[Ulrich Weigand]

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

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

[Tom Tromey]

>>>>> "Ulrich" == Ulrich Weigand <uweigand@de.ibm.com> writes:

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

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

I think it would be slightly more future-proof to choose a value in the
lo_user-hi_user range and put the new name, along with a detailed
comment, into include/dwarf2.h.

That way, if somebody finds a conflicting use, presumably they will be
editing the same place, see the comment, and change the value to
something that does not conflict.

Tom

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

[Ulrich Weigand]

Tom Tromey wrote:
> >>>>> "Ulrich" == Ulrich Weigand <uweigand@de.ibm.com> writes:
> 
> Ulrich> This extra flag is defined by GDB itself, and has a value outside
> Ulrich> the defined range of DW_AT_calling_convention attribute values,
> Ulrich> so there should be no potential conflict.
> 
> Ulrich> Does this look reasonable?  If anyone sees a better way to implement
> Ulrich> this, I'd appreciate any suggestions ...
> 
> I think it would be slightly more future-proof to choose a value in the
> lo_user-hi_user range and put the new name, along with a detailed
> comment, into include/dwarf2.h.
> 
> That way, if somebody finds a conflicting use, presumably they will be
> editing the same place, see the comment, and change the value to
> something that does not conflict.

Hmm, I had deliberately attempted to avoid touching dwarf2.h, since this
value is used only by GDB.  But I can see where your argument makes sense ...

The version below implements your suggestion.  Does this look good to you?

Thanks,
Ulrich


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.


diff -urNp gdb-orig/gdb/dwarf2read.c gdb-head/gdb/dwarf2read.c
--- gdb-orig/gdb/dwarf2read.c	2011-02-07 19:10:15.000000000 +0100
+++ gdb-head/gdb/dwarf2read.c	2011-02-07 19:14:17.000000000 +0100
@@ -7955,7 +7955,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/ppc-sysv-tdep.c gdb-head/gdb/ppc-sysv-tdep.c
--- gdb-orig/gdb/ppc-sysv-tdep.c	2011-02-07 19:11:14.000000000 +0100
+++ gdb-head/gdb/ppc-sysv-tdep.c	2011-02-07 19:29:09.000000000 +0100
@@ -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 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 +62,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 +337,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 +682,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 +829,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 +1041,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 +1052,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 +1114,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 +1142,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 +1357,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 +1685,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 +1752,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-02 14:36:50.000000000 +0100
+++ gdb-head/gdb/spu-tdep.c	2011-02-07 19:29:34.000000000 +0100
@@ -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 *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 +1467,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 +1483,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-07 19:14:17.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-07 19:14:17.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-02 14:36:50.000000000 +0100
+++ gdb-head/gdb/testsuite/gdb.opencl/Makefile.in	2011-02-07 19:14:17.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 $@..."
diff -urNp gdb-orig/include/dwarf2.h gdb-head/include/dwarf2.h
--- gdb-orig/include/dwarf2.h	2011-01-17 20:31:04.000000000 +0100
+++ gdb-head/include/dwarf2.h	2011-02-07 19:23:31.000000000 +0100
@@ -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.  */


-- 
  Dr. Ulrich Weigand
  GNU Toolchain for Linux on System z and Cell BE
  Ulrich.Weigand@de.ibm.com

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

[Tom Tromey]

>>>>> "Ulrich" == Ulrich Weigand <uweigand@de.ibm.com> writes:

Ulrich> The version below implements your suggestion.  Does this look
Ulrich> good to you?

Yes, thanks.

Tom

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

[Ulrich Weigand]

Tom Tromey:
> >>>>> "Ulrich" == Ulrich Weigand <uweigand@de.ibm.com> writes:
> 
> Ulrich> The version below implements your suggestion.  Does this look
> Ulrich> good to you?
> 
> Yes, thanks.

OK, I've checked this in now.

Thanks,
Ulrich

-- 
  Dr. Ulrich Weigand
  GNU Toolchain for Linux on System z and Cell BE
  Ulrich.Weigand@de.ibm.com

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

[Luis Machado]

On 02/08/2011 11:30 AM, Ulrich Weigand wrote:
> Tom Tromey:
>>>>>>> "Ulrich" == Ulrich Weigand<uweigand@de.ibm.com>  writes:
>>
>> Ulrich>  The version below implements your suggestion.  Does this look
>> Ulrich>  good to you?
>>
>> Yes, thanks.
>
> OK, I've checked this in now.
>
> Thanks,
> Ulrich
>

Hi Ulrich,

Did you forget the "gdbtypes.h" hunk? HEAD seems to be broken due to the 
missing DW_CC_GDB_IBM_OpenCL #define.

../../HEAD-sandbox/gdb/dwarf2read.c: In function ‘read_subroutine_type’:
../../HEAD-sandbox/gdb/dwarf2read.c:7961: error: ‘DW_CC_GDB_IBM_OpenCL’ 
undeclared (first use in this function)

Regards,
Luis

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

[Luis Machado]

On 02/08/2011 11:30 AM, Ulrich Weigand wrote:
> Tom Tromey:
>>>>>>> "Ulrich" == Ulrich Weigand<uweigand@de.ibm.com>  writes:
>>
>> Ulrich>  The version below implements your suggestion.  Does this look
>> Ulrich>  good to you?
>>
>> Yes, thanks.
>
> OK, I've checked this in now.
>
> Thanks,
> Ulrich
>

Hi Ulrich,

Did you forget the "gdbtypes.h" hunk? HEAD seems to be broken due to the 
missing DW_CC_GDB_IBM_OpenCL #define.

../../HEAD-sandbox/gdb/dwarf2read.c: In function ‘read_subroutine_type’:
../../HEAD-sandbox/gdb/dwarf2read.c:7961: error: ‘DW_CC_GDB_IBM_OpenCL’ 
undeclared (first use in this function)

Regards,
Luis

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

[Yao Qi]

On 02/14/2011 08:06 PM, Luis Machado wrote:
> Did you forget the "gdbtypes.h" hunk? HEAD seems to be broken due to the
> missing DW_CC_GDB_IBM_OpenCL #define.

No.  DW_CC_GDB_IBM_OpenCL is in include/dwarf2.h, by Ulrich's 2nd
version of this patch.

include/ChangeLog:

	* dwarf2.h (enum dwarf_calling_convention): Add DW_CC_GDB_IBM_OpenCL.

http://sourceware.org/ml/gdb-patches/2011-02/msg00167.html

-- 
Yao (齐尧)

[commit/powerpc] crash trying to allocate memory in inferior

[Joel Brobecker]

Our testsuite noticed a crash when trying to call a function which
requires GDB to allocate memory in the inferior. Typically, this
happens when one of the parameters is a string.  For instance, our
testcase tries:

    (gdb) call debug.trace (me, "You")
    [1]    32737 segmentation fault /path/to/gdb

What happens is that GDB sees the string, and thus tries to allocate
memory for it in the inferior:

> /* Allocate NBYTES of space in the inferior using the inferior's
>    malloc and return a value that is a pointer to the allocated
>    space.  */
>
> struct value *
> value_allocate_space_in_inferior (int len)
> {
>   struct objfile *objf;
>   struct value *val = find_function_in_inferior ("malloc", &objf);

And find_function_in_inferior first searches the symtab in case
we have debug info.  But, in our case (bareboard powerpc), we don't,
so it gets "malloc"'s address from the minimal symbols, and builds
a value whose type is a TYPE_CODE_PTR, not a TYPE_CODE_FUNC.

As a result, when we later try to make the call to malloc, we end up
inside the powerpc tdep code that has:

> do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *func_type,
[...]
>   if (func_type
>       && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL)

The problem is that func_type is not a TYPE_CODE_FUNC, and thus
the type-specific kind is not TYPE_SPECIFIC_FUNC, and so we do
TYPE_CALLING_CONVENTION is an invalid access.

Interestingly, the other call to TYPE_CALLING_CONVENTION is correctly
preceded by a check of the type's TYPE_CODE (making sure that it is
TYPE_CODE_FUNC).

gdb/ChangeLog:

        * ppc-sysv-tdep.c (do_ppc_sysv_return_value): Do not check
        FUNC_TYPE's calling convention if FUNC_TYPE is not a function.

tested on powerpc-elf. Checked in.

---
 gdb/ChangeLog       |    5 +++++
 gdb/ppc-sysv-tdep.c |    1 +
 2 files changed, 6 insertions(+), 0 deletions(-)

diff --git a/gdb/ChangeLog b/gdb/ChangeLog
index 30cf144..f5cdd45 100644
--- a/gdb/ChangeLog
+++ b/gdb/ChangeLog
@@ -1,3 +1,8 @@
+2011-10-24  Joel Brobecker  <brobecker@adacore.com>
+
+	* ppc-sysv-tdep.c (do_ppc_sysv_return_value): Do not check
+	FUNC_TYPE's calling convention if FUNC_TYPE is not a function.
+
 2011-10-24  Pedro Alves  <pedro@codesourcery.com>
 
 	* linux-nat.c (linux_handle_extended_wait): When handling a clone
diff --git a/gdb/ppc-sysv-tdep.c b/gdb/ppc-sysv-tdep.c
index e431363..bda4544 100644
--- a/gdb/ppc-sysv-tdep.c
+++ b/gdb/ppc-sysv-tdep.c
@@ -692,6 +692,7 @@ do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *func_type,
   int opencl_abi = 0;
 
   if (func_type
+      && TYPE_CODE (func_type) == TYPE_CODE_FUNC
       && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL)
     opencl_abi = 1;
 
-- 
1.7.1

Re: [commit/powerpc] crash trying to allocate memory in inferior

[Ulrich Weigand]

Joel Brobeker wrote:

>         * ppc-sysv-tdep.c (do_ppc_sysv_return_value): Do not check
>         FUNC_TYPE's calling convention if FUNC_TYPE is not a function.

Huh, I had just been testing a different patch for the same problem ...

Since this patch is a bit more general (it also fixes the 64-bit case,
and actually handles pointers to OpenCL functions), I've checked it
in anyway.

Tested on powerpc64-linux.

Bye,
Ulrich


ChangeLog:

	* ppc-sysv-tdep.c (ppc_sysv_use_opencl_abi): New function.
	(ppc_sysv_abi_push_dummy_call): Use it.
	(do_ppc_sysv_return_value): Likewise.
	(ppc64_sysv_abi_push_dummy_call): Likewise.
	(ppc64_sysv_abi_return_value): Likewise.


Index: gdb/ppc-sysv-tdep.c
===================================================================
RCS file: /cvs/src/src/gdb/ppc-sysv-tdep.c,v
retrieving revision 1.64
diff -u -p -r1.64 ppc-sysv-tdep.c
--- gdb/ppc-sysv-tdep.c	24 Oct 2011 16:51:36 -0000	1.64
+++ gdb/ppc-sysv-tdep.c	26 Oct 2011 16:12:56 -0000
@@ -32,6 +32,22 @@
 #include "infcall.h"
 #include "dwarf2.h"
 
+
+/* Check whether FTPYE is a (pointer to) function type that should use
+   the OpenCL vector ABI.  */
+
+static int
+ppc_sysv_use_opencl_abi (struct type *ftype)
+{
+  ftype = check_typedef (ftype);
+
+  if (TYPE_CODE (ftype) == TYPE_CODE_PTR)
+    ftype = check_typedef (TYPE_TARGET_TYPE (ftype));
+
+  return (TYPE_CODE (ftype) == TYPE_CODE_FUNC
+	  && TYPE_CALLING_CONVENTION (ftype) == DW_CC_GDB_IBM_OpenCL);
+}
+
 /* 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
    be less than eight parameters if some parameters occupy more than one
@@ -51,8 +67,7 @@ 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;
+  int opencl_abi = ppc_sysv_use_opencl_abi (value_type (function));
   ULONGEST saved_sp;
   int argspace = 0;		/* 0 is an initial wrong guess.  */
   int write_pass;
@@ -62,13 +77,6 @@ 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
@@ -689,12 +697,7 @@ do_ppc_sysv_return_value (struct gdbarch
 {
   struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
   enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
-  int opencl_abi = 0;
-
-  if (func_type
-      && TYPE_CODE (func_type) == TYPE_CODE_FUNC
-      && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL)
-    opencl_abi = 1;
+  int opencl_abi = func_type? ppc_sysv_use_opencl_abi (func_type) : 0;
 
   gdb_assert (tdep->wordsize == 4);
 
@@ -1115,8 +1118,7 @@ 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;
+  int opencl_abi = ppc_sysv_use_opencl_abi (value_type (function));
   ULONGEST back_chain;
   /* See for-loop comment below.  */
   int write_pass;
@@ -1146,13 +1148,6 @@ 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
@@ -1721,11 +1716,7 @@ 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;
+  int opencl_abi = func_type? ppc_sysv_use_opencl_abi (func_type) : 0;
 
   /* This function exists to support a calling convention that
      requires floating-point registers.  It shouldn't be used on


-- 
  Dr. Ulrich Weigand
  GNU Toolchain for Linux on System z and Cell BE
  Ulrich.Weigand@de.ibm.com

Re: [commit/powerpc] crash trying to allocate memory in inferior

[Joel Brobecker]

> Huh, I had just been testing a different patch for the same problem ...
> 
> Since this patch is a bit more general (it also fixes the 64-bit case,
> and actually handles pointers to OpenCL functions), I've checked it
> in anyway.

Nice! Thank you.

-- 
Joel

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

[Mark Kettenis]

> Date: Wed, 2 Feb 2011 18:48:00 +0100 (CET)
> From: "Ulrich Weigand" <uweigand@de.ibm.com>
> 
> 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.

I didn't look too closely at the diff yet, but given that
push_dummy_call() functions tend to be fairly complex already, would
it be possible to move the OpenCL calling convention into a seperate
function?

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

[Ulrich Weigand]

Mark Kettenis wrote:

> I didn't look too closely at the diff yet, but given that
> push_dummy_call() functions tend to be fairly complex already, would
> it be possible to move the OpenCL calling convention into a seperate
> function?

Having a completely separate function at the push_dummy_call level
would lead to significant code duplication, since most of the "regular"
data types are also available in OpenCL C.  The new code just adds
support for the special vector types.

I certainly agree that those functions are quite complex already.  One
reason for that is that pushing arguments is inherently stateful, and
that state is currently spread out across a significant number of local
variables in the push_dummy_call functions.  This makes using subroutines
and helper functions difficult.

It seems one way towards a refactoring of those routines would be to
extract that state into an explicit data structure (along the lines of
CUMULATIVE_ARGS in GCC), and then split handling of different types
of arguments into functions of their own, passing that state around.

If this looks useful, I can work on patch along those lines ...

Bye,
Ulrich

-- 
  Dr. Ulrich Weigand
  GNU Toolchain for Linux on System z and Cell BE
  Ulrich.Weigand@de.ibm.com