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.

10 files changed, 796 insertions, 12 deletions

diff --git a/gdb/ChangeLog b/gdb/ChangeLog
index 0197e8b..87f31f3 100644
--- a/gdb/ChangeLog
+++ b/gdb/ChangeLog
@@ -1,5 +1,23 @@
 2011-02-08  Ulrich Weigand  <uweigand@de.ibm.com>
 
+	* dwarf2read.c (read_subroutine_type): Set special calling
+	convention flag for functions compiled by IBM XL C for OpenCL.
+	* ppc-sysv-tdep.c: Include "dwarf2.h"
+	(ppc_sysv_abi_push_dummy_call): Implement IBM OpenCL vector types
+	calling convention.
+	(do_ppc_sysv_return_value): Add FUNC_TYPE argument.  Implement
+	IBM OpenCL vector types calling convention.
+	(ppc_sysv_abi_return_value): Pass through FUNC_TYPE.
+	(ppc_sysv_abi_broken_return_value): Likewise.
+	(ppc64_sysv_abi_push_dummy_call): Implement IBM OpenCL vector
+	types calling convention.
+	(ppc64_sysv_abi_return_value): Likewise.
+	* spu-tdep.c: Include "dwarf2.h"
+	(spu_return_value): Implement IBM OpenCL vector types calling
+	convention.
+
+2011-02-08  Ulrich Weigand  <uweigand@de.ibm.com>
+
 	* ppc-sysv-tdep.c (ppc64_sysv_abi_push_dummy_call): Implement
 	correct ABI for AltiVec vector arguments.
 
diff --git a/gdb/dwarf2read.c b/gdb/dwarf2read.c
index f269dee..837fa3b 100644
--- a/gdb/dwarf2read.c
+++ b/gdb/dwarf2read.c
@@ -7955,7 +7955,12 @@ read_subroutine_type (struct die_info *die, struct dwarf2_cu *cu)
      the subroutine die.  Otherwise set the calling convention to
      the default value DW_CC_normal.  */
   attr = dwarf2_attr (die, DW_AT_calling_convention, cu);
-  TYPE_CALLING_CONVENTION (ftype) = attr ? DW_UNSND (attr) : DW_CC_normal;
+  if (attr)
+    TYPE_CALLING_CONVENTION (ftype) = DW_UNSND (attr);
+  else if (cu->producer && strstr (cu->producer, "IBM XL C for OpenCL"))
+    TYPE_CALLING_CONVENTION (ftype) = DW_CC_GDB_IBM_OpenCL;
+  else
+    TYPE_CALLING_CONVENTION (ftype) = DW_CC_normal;
 
   /* We need to add the subroutine type to the die immediately so
      we don't infinitely recurse when dealing with parameters
diff --git a/gdb/ppc-sysv-tdep.c b/gdb/ppc-sysv-tdep.c
index 300dcac..872117d 100644
--- a/gdb/ppc-sysv-tdep.c
+++ b/gdb/ppc-sysv-tdep.c
@@ -30,6 +30,7 @@
 #include "target.h"
 #include "objfiles.h"
 #include "infcall.h"
+#include "dwarf2.h"
 
 /* Pass the arguments in either registers, or in the stack.  Using the
    ppc sysv ABI, the first eight words of the argument list (that might
@@ -50,6 +51,8 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function,
 {
   struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
   enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
+  struct type *ftype;
+  int opencl_abi = 0;
   ULONGEST saved_sp;
   int argspace = 0;		/* 0 is an initial wrong guess.  */
   int write_pass;
@@ -59,6 +62,13 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function,
   regcache_cooked_read_unsigned (regcache, gdbarch_sp_regnum (gdbarch),
 				 &saved_sp);
 
+  ftype = check_typedef (value_type (function));
+  if (TYPE_CODE (ftype) == TYPE_CODE_PTR)
+    ftype = check_typedef (TYPE_TARGET_TYPE (ftype));
+  if (TYPE_CODE (ftype) == TYPE_CODE_FUNC
+      && TYPE_CALLING_CONVENTION (ftype) == DW_CC_GDB_IBM_OpenCL)
+    opencl_abi = 1;
+
   /* Go through the argument list twice.
 
      Pass 1: Figure out how much new stack space is required for
@@ -327,6 +337,126 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function,
 		 Hence we increase freg even when writing to memory.  */
 	      freg += 2;
 	    }
+	  else if (len < 16
+		   && TYPE_CODE (type) == TYPE_CODE_ARRAY
+		   && TYPE_VECTOR (type)
+		   && opencl_abi)
+	    {
+	      /* OpenCL vectors shorter than 16 bytes are passed as if
+		 a series of independent scalars.  */
+	      struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type));
+	      int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype);
+
+	      for (i = 0; i < nelt; i++)
+		{
+		  const gdb_byte *elval = val + i * TYPE_LENGTH (eltype);
+
+		  if (TYPE_CODE (eltype) == TYPE_CODE_FLT && !tdep->soft_float)
+		    {
+		      if (freg <= 8)
+			{
+			  if (write_pass)
+			    {
+			      int regnum = tdep->ppc_fp0_regnum + freg;
+			      gdb_byte regval[MAX_REGISTER_SIZE];
+			      struct type *regtype
+				= register_type (gdbarch, regnum);
+			      convert_typed_floating (elval, eltype,
+						      regval, regtype);
+			      regcache_cooked_write (regcache, regnum, regval);
+			    }
+			  freg++;
+			}
+		      else
+			{
+			  argoffset = align_up (argoffset, len);
+			  if (write_pass)
+			    write_memory (sp + argoffset, val, len);
+			  argoffset += len;
+			}
+		    }
+		  else if (TYPE_LENGTH (eltype) == 8)
+		    {
+		      if (greg > 9)
+			{
+			  /* Just in case GREG was 10.  */
+			  greg = 11;
+			  argoffset = align_up (argoffset, 8);
+			  if (write_pass)
+			    write_memory (sp + argoffset, elval,
+					  TYPE_LENGTH (eltype));
+			  argoffset += 8;
+			}
+		      else
+			{
+			  /* Must start on an odd register - r3/r4 etc.  */
+			  if ((greg & 1) == 0)
+			    greg++;
+			  if (write_pass)
+			    {
+			      int regnum = tdep->ppc_gp0_regnum + greg;
+			      regcache_cooked_write (regcache,
+						     regnum + 0, elval + 0);
+			      regcache_cooked_write (regcache,
+						     regnum + 1, elval + 4);
+			    }
+			  greg += 2;
+			}
+		    }
+		  else
+		    {
+		      gdb_byte word[MAX_REGISTER_SIZE];
+		      store_unsigned_integer (word, tdep->wordsize, byte_order,
+					      unpack_long (eltype, elval));
+
+		      if (greg <= 10)
+			{
+			  if (write_pass)
+			    regcache_cooked_write (regcache,
+						   tdep->ppc_gp0_regnum + greg,
+						   word);
+			  greg++;
+			}
+		      else
+			{
+			  argoffset = align_up (argoffset, tdep->wordsize);
+			  if (write_pass)
+			    write_memory (sp + argoffset, word, tdep->wordsize);
+			  argoffset += tdep->wordsize;
+			}
+		    }
+		}
+	    }
+	  else if (len >= 16
+		   && TYPE_CODE (type) == TYPE_CODE_ARRAY
+		   && TYPE_VECTOR (type)
+		   && opencl_abi)
+	    {
+	      /* OpenCL vectors 16 bytes or longer are passed as if
+		 a series of AltiVec vectors.  */
+	      int i;
+
+	      for (i = 0; i < len / 16; i++)
+		{
+		  const gdb_byte *elval = val + i * 16;
+
+		  if (vreg <= 13)
+		    {
+		      if (write_pass)
+			regcache_cooked_write (regcache,
+					       tdep->ppc_vr0_regnum + vreg,
+					       elval);
+		      vreg++;
+		    }
+		  else
+		    {
+		      argoffset = align_up (argoffset, 16);
+		      if (write_pass)
+			write_memory (sp + argoffset, elval, 16);
+		      argoffset += 16;
+		    }
+		}
+	    }
 	  else if (len == 16
 		   && TYPE_CODE (type) == TYPE_CODE_ARRAY
 		   && TYPE_VECTOR (type)
@@ -552,13 +682,21 @@ get_decimal_float_return_value (struct gdbarch *gdbarch, struct type *valtype,
    when returned in general-purpose registers.  */
 
 static enum return_value_convention
-do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *type,
-			  struct regcache *regcache, gdb_byte *readbuf,
-			  const gdb_byte *writebuf, int broken_gcc)
+do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *func_type,
+			  struct type *type, struct regcache *regcache,
+			  gdb_byte *readbuf, const gdb_byte *writebuf,
+			  int broken_gcc)
 {
   struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
   enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
+  int opencl_abi = 0;
+
+  if (func_type
+      && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL)
+    opencl_abi = 1;
+
   gdb_assert (tdep->wordsize == 4);
+
   if (TYPE_CODE (type) == TYPE_CODE_FLT
       && TYPE_LENGTH (type) <= 8
       && !tdep->soft_float)
@@ -691,6 +829,83 @@ do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *type,
 	}
       return RETURN_VALUE_REGISTER_CONVENTION;
     }
+  /* OpenCL vectors < 16 bytes are returned as distinct
+     scalars in f1..f2 or r3..r10.  */
+  if (TYPE_CODE (type) == TYPE_CODE_ARRAY
+      && TYPE_VECTOR (type)
+      && TYPE_LENGTH (type) < 16
+      && opencl_abi)
+    {
+      struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type));
+      int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype);
+
+      for (i = 0; i < nelt; i++)
+	{
+	  int offset = i * TYPE_LENGTH (eltype);
+
+	  if (TYPE_CODE (eltype) == TYPE_CODE_FLT)
+	    {
+	      int regnum = tdep->ppc_fp0_regnum + 1 + i;
+	      gdb_byte regval[MAX_REGISTER_SIZE];
+	      struct type *regtype = register_type (gdbarch, regnum);
+
+	      if (writebuf != NULL)
+		{
+		  convert_typed_floating (writebuf + offset, eltype,
+					  regval, regtype);
+		  regcache_cooked_write (regcache, regnum, regval);
+		}
+	      if (readbuf != NULL)
+		{
+		  regcache_cooked_read (regcache, regnum, regval);
+		  convert_typed_floating (regval, regtype,
+					  readbuf + offset, eltype);
+		}
+	    }
+	  else
+	    {
+	      int regnum = tdep->ppc_gp0_regnum + 3 + i;
+	      ULONGEST regval;
+
+	      if (writebuf != NULL)
+		{
+		  regval = unpack_long (eltype, writebuf + offset);
+		  regcache_cooked_write_unsigned (regcache, regnum, regval);
+		}
+	      if (readbuf != NULL)
+		{
+		  regcache_cooked_read_unsigned (regcache, regnum, &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 gdbarch *gdbarch, struct type *func_type,
 			   struct type *valtype, struct regcache *regcache,
 			   gdb_byte *readbuf, const gdb_byte *writebuf)
 {
-  return do_ppc_sysv_return_value (gdbarch, valtype, regcache, readbuf,
-				   writebuf, 0);
+  return do_ppc_sysv_return_value (gdbarch, func_type, valtype, regcache,
+				   readbuf, writebuf, 0);
 }
 
 enum return_value_convention
@@ -837,8 +1052,8 @@ ppc_sysv_abi_broken_return_value (struct gdbarch *gdbarch,
 				  struct regcache *regcache,
 				  gdb_byte *readbuf, const gdb_byte *writebuf)
 {
-  return do_ppc_sysv_return_value (gdbarch, valtype, regcache, readbuf,
-				   writebuf, 1);
+  return do_ppc_sysv_return_value (gdbarch, func_type, valtype, regcache,
+				   readbuf, writebuf, 1);
 }
 
 /* The helper function for 64-bit SYSV push_dummy_call.  Converts the
@@ -899,6 +1114,8 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch,
   CORE_ADDR func_addr = find_function_addr (function, NULL);
   struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
   enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
+  struct type *ftype;
+  int opencl_abi = 0;
   ULONGEST back_chain;
   /* See for-loop comment below.  */
   int write_pass;
@@ -925,6 +1142,13 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch,
   regcache_cooked_read_unsigned (regcache, gdbarch_sp_regnum (gdbarch),
 				 &back_chain);
 
+  ftype = check_typedef (value_type (function));
+  if (TYPE_CODE (ftype) == TYPE_CODE_PTR)
+    ftype = check_typedef (TYPE_TARGET_TYPE (ftype));
+  if (TYPE_CODE (ftype) == TYPE_CODE_FUNC
+      && TYPE_CALLING_CONVENTION (ftype) == DW_CC_GDB_IBM_OpenCL)
+    opencl_abi = 1;
+
   /* Go through the argument list twice.
 
      Pass 1: Compute the function call's stack space and register
@@ -1133,6 +1357,109 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch,
 	      greg += 2;
 	      gparam = align_up (gparam + TYPE_LENGTH (type), tdep->wordsize);
 	    }
+	  else if (TYPE_LENGTH (type) < 16
+		   && TYPE_CODE (type) == TYPE_CODE_ARRAY
+		   && TYPE_VECTOR (type)
+		   && opencl_abi)
+	    {
+	      /* OpenCL vectors shorter than 16 bytes are passed as if
+		 a series of independent scalars.  */
+	      struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type));
+	      int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype);
+
+	      for (i = 0; i < nelt; i++)
+		{
+		  const gdb_byte *elval = val + i * TYPE_LENGTH (eltype);
+
+		  if (TYPE_CODE (eltype) == TYPE_CODE_FLT)
+		    {
+		      if (write_pass)
+			{
+			  gdb_byte regval[MAX_REGISTER_SIZE];
+			  const gdb_byte *p;
+
+			  if (TYPE_LENGTH (eltype) == 4)
+			    {
+			      memcpy (regval, elval, 4);
+			      memcpy (regval + 4, elval, 4);
+			      p = regval;
+			    }
+			  else
+			    p = elval;
+
+			  write_memory (gparam, p, 8);
+
+			  if (freg <= 13)
+			    {
+			      int regnum = tdep->ppc_fp0_regnum + freg;
+			      struct type *regtype
+				= register_type (gdbarch, regnum);
+
+			      convert_typed_floating (elval, eltype,
+						      regval, regtype);
+			      regcache_cooked_write (regcache, regnum, regval);
+			    }
+
+			  if (greg <= 10)
+			    regcache_cooked_write (regcache,
+						   tdep->ppc_gp0_regnum + greg,
+						   regval);
+			}
+
+		      freg++;
+		      greg++;
+		      gparam = align_up (gparam + 8, tdep->wordsize);
+		    }
+		  else
+		    {
+		      if (write_pass)
+			{
+			  ULONGEST word = unpack_long (eltype, elval);
+			  if (greg <= 10)
+			    regcache_cooked_write_unsigned
+			      (regcache, tdep->ppc_gp0_regnum + greg, word);
+
+			  write_memory_unsigned_integer
+			    (gparam, tdep->wordsize, byte_order, word);
+			}
+
+		      greg++;
+		      gparam = align_up (gparam + TYPE_LENGTH (eltype),
+					 tdep->wordsize);
+		    }
+		}
+	    }
+	  else if (TYPE_LENGTH (type) >= 16
+		   && TYPE_CODE (type) == TYPE_CODE_ARRAY
+		   && TYPE_VECTOR (type)
+		   && opencl_abi)
+	    {
+	      /* OpenCL vectors 16 bytes or longer are passed as if
+		 a series of AltiVec vectors.  */
+	      int i;
+
+	      for (i = 0; i < TYPE_LENGTH (type) / 16; i++)
+		{
+		  const gdb_byte *elval = val + i * 16;
+
+		  gparam = align_up (gparam, 16);
+		  greg += greg & 1;
+
+		  if (write_pass)
+		    {
+		      if (vreg <= 13)
+			regcache_cooked_write (regcache,
+					       tdep->ppc_vr0_regnum + vreg,
+					       elval);
+
+		      write_memory (gparam, elval, 16);
+		    }
+
+		  greg += 2;
+		  vreg++;
+		  gparam += 16;
+		}
+	    }
 	  else if (TYPE_LENGTH (type) == 16 && TYPE_VECTOR (type)
 		   && TYPE_CODE (type) == TYPE_CODE_ARRAY
 		   && tdep->ppc_vr0_regnum >= 0)
@@ -1358,6 +1685,11 @@ ppc64_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type,
 {
   struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
   enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
+  int opencl_abi = 0;
+
+  if (func_type
+      && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL)
+    opencl_abi = 1;
 
   /* This function exists to support a calling convention that
      requires floating-point registers.  It shouldn't be used on
@@ -1420,6 +1752,83 @@ ppc64_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type,
 	regcache_cooked_read (regcache, tdep->ppc_gp0_regnum + 3, readbuf);
       return RETURN_VALUE_REGISTER_CONVENTION;
     }
+  /* OpenCL vectors < 16 bytes are returned as distinct
+     scalars in f1..f2 or r3..r10.  */
+  if (TYPE_CODE (valtype) == TYPE_CODE_ARRAY
+      && TYPE_VECTOR (valtype)
+      && TYPE_LENGTH (valtype) < 16
+      && opencl_abi)
+    {
+      struct type *eltype = check_typedef (TYPE_TARGET_TYPE (valtype));
+      int i, nelt = TYPE_LENGTH (valtype) / TYPE_LENGTH (eltype);
+
+      for (i = 0; i < nelt; i++)
+	{
+	  int offset = i * TYPE_LENGTH (eltype);
+
+	  if (TYPE_CODE (eltype) == TYPE_CODE_FLT)
+	    {
+	      int regnum = tdep->ppc_fp0_regnum + 1 + i;
+	      gdb_byte regval[MAX_REGISTER_SIZE];
+	      struct type *regtype = register_type (gdbarch, regnum);
+
+	      if (writebuf != NULL)
+		{
+		  convert_typed_floating (writebuf + offset, eltype,
+					  regval, regtype);
+		  regcache_cooked_write (regcache, regnum, regval);
+		}
+	      if (readbuf != NULL)
+		{
+		  regcache_cooked_read (regcache, regnum, regval);
+		  convert_typed_floating (regval, regtype,
+					  readbuf + offset, eltype);
+		}
+	    }
+	  else
+	    {
+	      int regnum = tdep->ppc_gp0_regnum + 3 + i;
+	      ULONGEST regval;
+
+	      if (writebuf != NULL)
+		{
+		  regval = unpack_long (eltype, writebuf + offset);
+		  regcache_cooked_write_unsigned (regcache, regnum, regval);
+		}
+	      if (readbuf != NULL)
+		{
+		  regcache_cooked_read_unsigned (regcache, regnum, &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 --git a/gdb/spu-tdep.c b/gdb/spu-tdep.c
index d424659..0b0ea4e 100644
--- a/gdb/spu-tdep.c
+++ b/gdb/spu-tdep.c
@@ -44,6 +44,7 @@
 #include "block.h"
 #include "observer.h"
 #include "infcall.h"
+#include "dwarf2.h"
 
 #include "spu-tdep.h"
 
@@ -1448,6 +1449,13 @@ spu_return_value (struct gdbarch *gdbarch, struct type *func_type,
 		  gdb_byte *out, const gdb_byte *in)
 {
   enum return_value_convention rvc;
+  int opencl_vector = 0;
+
+  if (func_type
+      && TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL
+      && TYPE_CODE (type) == TYPE_CODE_ARRAY
+      && TYPE_VECTOR (type))
+    opencl_vector = 1;
 
   if (TYPE_LENGTH (type) <= (SPU_ARGN_REGNUM - SPU_ARG1_REGNUM + 1) * 16)
     rvc = RETURN_VALUE_REGISTER_CONVENTION;
@@ -1459,7 +1467,10 @@ spu_return_value (struct gdbarch *gdbarch, struct type *func_type,
       switch (rvc)
 	{
 	case RETURN_VALUE_REGISTER_CONVENTION:
-	  spu_value_to_regcache (regcache, SPU_ARG1_REGNUM, type, in);
+	  if (opencl_vector && TYPE_LENGTH (type) == 2)
+	    regcache_cooked_write_part (regcache, SPU_ARG1_REGNUM, 2, 2, in);
+	  else
+	    spu_value_to_regcache (regcache, SPU_ARG1_REGNUM, type, in);
 	  break;
 
 	case RETURN_VALUE_STRUCT_CONVENTION:
@@ -1472,7 +1483,10 @@ spu_return_value (struct gdbarch *gdbarch, struct type *func_type,
       switch (rvc)
 	{
 	case RETURN_VALUE_REGISTER_CONVENTION:
-	  spu_regcache_to_value (regcache, SPU_ARG1_REGNUM, type, out);
+	  if (opencl_vector && TYPE_LENGTH (type) == 2)
+	    regcache_cooked_read_part (regcache, SPU_ARG1_REGNUM, 2, 2, out);
+	  else
+	    spu_regcache_to_value (regcache, SPU_ARG1_REGNUM, type, out);
 	  break;
 
 	case RETURN_VALUE_STRUCT_CONVENTION:
diff --git a/gdb/testsuite/ChangeLog b/gdb/testsuite/ChangeLog
index 3cb8d41..19aec4c 100644
--- a/gdb/testsuite/ChangeLog
+++ b/gdb/testsuite/ChangeLog
@@ -1,5 +1,11 @@
 2011-02-08  Ulrich Weigand  <uweigand@de.ibm.com>
 
+	* gdb.opencl/callfuncs.cl: New file.
+	* gdb.opencl/callfuncs.exp: New test.
+	* gdb.opencl/Makefile.in (EXECUTABLES): Add callfuncs.
+
+2011-02-08  Ulrich Weigand  <uweigand@de.ibm.com>
+
 	* gdb.arch/altivec-abi.c (vec_func): Make use of intv_on_stack_f
 	when computing result.
 	* gdb.arch/altivec-abi.exp: Update expected results.
diff --git a/gdb/testsuite/gdb.opencl/Makefile.in b/gdb/testsuite/gdb.opencl/Makefile.in
index c12aef3..7dec34c 100644
--- a/gdb/testsuite/gdb.opencl/Makefile.in
+++ b/gdb/testsuite/gdb.opencl/Makefile.in
@@ -1,7 +1,7 @@
 VPATH = @srcdir@
 srcdir = @srcdir@
 
-EXECUTABLES = datatypes vec_comps convs_casts operators
+EXECUTABLES = callfuncs datatypes vec_comps convs_casts operators
 
 all info install-info dvi install uninstall installcheck check:
 	@echo "Nothing to be done for $@..."
diff --git a/gdb/testsuite/gdb.opencl/callfuncs.cl b/gdb/testsuite/gdb.opencl/callfuncs.cl
new file mode 100644
index 0000000..6d53ee0
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/callfuncs.cl
@@ -0,0 +1,218 @@
+/* This testcase is part of GDB, the GNU debugger.
+
+   Copyright 2011 Free Software Foundation, Inc.
+
+   This program is free software; you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3 of the License, or
+   (at your option) any later version.
+
+   This program is distributed in the hope that it will be useful,
+   but WITHOUT ANY WARRANTY; without even the implied warranty of
+   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+   GNU General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with this program.  If not, see <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 --git a/gdb/testsuite/gdb.opencl/callfuncs.exp b/gdb/testsuite/gdb.opencl/callfuncs.exp
new file mode 100644
index 0000000..f435589
--- /dev/null
+++ b/gdb/testsuite/gdb.opencl/callfuncs.exp
@@ -0,0 +1,102 @@
+# Copyright 2011 Free Software Foundation, Inc.
+
+# This program is free software; you can redistribute it and/or modify
+# it under the terms of the GNU General Public License as published by
+# the Free Software Foundation; either version 3 of the License, or
+# (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program.  If not, see <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 --git a/include/ChangeLog b/include/ChangeLog
index c96d358..2a19c61 100644
--- a/include/ChangeLog
+++ b/include/ChangeLog
@@ -1,3 +1,7 @@
+2011-02-08  Ulrich Weigand  <uweigand@de.ibm.com>
+
+	* dwarf2.h (enum dwarf_calling_convention): Add DW_CC_GDB_IBM_OpenCL.
+
 2011-01-12  Iain Sandoe  <iains@gcc.gnu.org>
 
 	* dwarf2.h: Update value for DW_AT_hi_user.
diff --git a/include/dwarf2.h b/include/dwarf2.h
index ad00aac..7729ad8 100644
--- a/include/dwarf2.h
+++ b/include/dwarf2.h
@@ -754,7 +754,15 @@ enum dwarf_calling_convention
     DW_CC_hi_user = 0xff,
 
     DW_CC_GNU_renesas_sh = 0x40,
-    DW_CC_GNU_borland_fastcall_i386 = 0x41
+    DW_CC_GNU_borland_fastcall_i386 = 0x41,
+
+    /* This DW_CC_ value is not currently generated by any toolchain.  It is
+       used internally to GDB to indicate OpenCL C functions that have been
+       compiled with the IBM XL C for OpenCL compiler and use a non-platform
+       calling convention for passing OpenCL C vector types.  This value may
+       be changed freely as long as it does not conflict with any other DW_CC_
+       value defined here.  */
+    DW_CC_GDB_IBM_OpenCL = 0xff
   };
 
 /* Inline attribute.  */