This is the mail archive of the gdb-patches@sourceware.org mailing list for the GDB project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

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


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


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]