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.
This commit is contained in:
Ulrich Weigand
parent
d6dafb7c8e
commit
54fcddd0ac
10 changed files with 796 additions and 12 deletions
|
|
@ -1,3 +1,21 @@
|
|||
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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -30,6 +30,7 @@
|
|||
#include "target.h"
|
||||
#include "objfiles.h"
|
||||
#include "infcall.h"
|
||||
#include "dwarf2.h"
|
||||
|
||||
/* Pass the arguments in either registers, or in the stack. Using the
|
||||
ppc sysv ABI, the first eight words of the argument list (that might
|
||||
|
|
@ -50,6 +51,8 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function,
|
|||
{
|
||||
struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
|
||||
enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
|
||||
struct type *ftype;
|
||||
int opencl_abi = 0;
|
||||
ULONGEST saved_sp;
|
||||
int argspace = 0; /* 0 is an initial wrong guess. */
|
||||
int write_pass;
|
||||
|
|
@ -59,6 +62,13 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function,
|
|||
regcache_cooked_read_unsigned (regcache, gdbarch_sp_regnum (gdbarch),
|
||||
&saved_sp);
|
||||
|
||||
ftype = check_typedef (value_type (function));
|
||||
if (TYPE_CODE (ftype) == TYPE_CODE_PTR)
|
||||
ftype = check_typedef (TYPE_TARGET_TYPE (ftype));
|
||||
if (TYPE_CODE (ftype) == TYPE_CODE_FUNC
|
||||
&& TYPE_CALLING_CONVENTION (ftype) == DW_CC_GDB_IBM_OpenCL)
|
||||
opencl_abi = 1;
|
||||
|
||||
/* Go through the argument list twice.
|
||||
|
||||
Pass 1: Figure out how much new stack space is required for
|
||||
|
|
@ -327,6 +337,126 @@ ppc_sysv_abi_push_dummy_call (struct gdbarch *gdbarch, struct value *function,
|
|||
Hence we increase freg even when writing to memory. */
|
||||
freg += 2;
|
||||
}
|
||||
else if (len < 16
|
||||
&& TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
&& opencl_abi)
|
||||
{
|
||||
/* OpenCL vectors shorter than 16 bytes are passed as if
|
||||
a series of independent scalars. */
|
||||
struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type));
|
||||
int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype);
|
||||
|
||||
for (i = 0; i < nelt; i++)
|
||||
{
|
||||
const gdb_byte *elval = val + i * TYPE_LENGTH (eltype);
|
||||
|
||||
if (TYPE_CODE (eltype) == TYPE_CODE_FLT && !tdep->soft_float)
|
||||
{
|
||||
if (freg <= 8)
|
||||
{
|
||||
if (write_pass)
|
||||
{
|
||||
int regnum = tdep->ppc_fp0_regnum + freg;
|
||||
gdb_byte regval[MAX_REGISTER_SIZE];
|
||||
struct type *regtype
|
||||
= register_type (gdbarch, regnum);
|
||||
convert_typed_floating (elval, eltype,
|
||||
regval, regtype);
|
||||
regcache_cooked_write (regcache, regnum, regval);
|
||||
}
|
||||
freg++;
|
||||
}
|
||||
else
|
||||
{
|
||||
argoffset = align_up (argoffset, len);
|
||||
if (write_pass)
|
||||
write_memory (sp + argoffset, val, len);
|
||||
argoffset += len;
|
||||
}
|
||||
}
|
||||
else if (TYPE_LENGTH (eltype) == 8)
|
||||
{
|
||||
if (greg > 9)
|
||||
{
|
||||
/* Just in case GREG was 10. */
|
||||
greg = 11;
|
||||
argoffset = align_up (argoffset, 8);
|
||||
if (write_pass)
|
||||
write_memory (sp + argoffset, elval,
|
||||
TYPE_LENGTH (eltype));
|
||||
argoffset += 8;
|
||||
}
|
||||
else
|
||||
{
|
||||
/* Must start on an odd register - r3/r4 etc. */
|
||||
if ((greg & 1) == 0)
|
||||
greg++;
|
||||
if (write_pass)
|
||||
{
|
||||
int regnum = tdep->ppc_gp0_regnum + greg;
|
||||
regcache_cooked_write (regcache,
|
||||
regnum + 0, elval + 0);
|
||||
regcache_cooked_write (regcache,
|
||||
regnum + 1, elval + 4);
|
||||
}
|
||||
greg += 2;
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
gdb_byte word[MAX_REGISTER_SIZE];
|
||||
store_unsigned_integer (word, tdep->wordsize, byte_order,
|
||||
unpack_long (eltype, elval));
|
||||
|
||||
if (greg <= 10)
|
||||
{
|
||||
if (write_pass)
|
||||
regcache_cooked_write (regcache,
|
||||
tdep->ppc_gp0_regnum + greg,
|
||||
word);
|
||||
greg++;
|
||||
}
|
||||
else
|
||||
{
|
||||
argoffset = align_up (argoffset, tdep->wordsize);
|
||||
if (write_pass)
|
||||
write_memory (sp + argoffset, word, tdep->wordsize);
|
||||
argoffset += tdep->wordsize;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
else if (len >= 16
|
||||
&& TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
&& opencl_abi)
|
||||
{
|
||||
/* OpenCL vectors 16 bytes or longer are passed as if
|
||||
a series of AltiVec vectors. */
|
||||
int i;
|
||||
|
||||
for (i = 0; i < len / 16; i++)
|
||||
{
|
||||
const gdb_byte *elval = val + i * 16;
|
||||
|
||||
if (vreg <= 13)
|
||||
{
|
||||
if (write_pass)
|
||||
regcache_cooked_write (regcache,
|
||||
tdep->ppc_vr0_regnum + vreg,
|
||||
elval);
|
||||
vreg++;
|
||||
}
|
||||
else
|
||||
{
|
||||
argoffset = align_up (argoffset, 16);
|
||||
if (write_pass)
|
||||
write_memory (sp + argoffset, elval, 16);
|
||||
argoffset += 16;
|
||||
}
|
||||
}
|
||||
}
|
||||
else if (len == 16
|
||||
&& TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
|
|
@ -552,13 +682,21 @@ get_decimal_float_return_value (struct gdbarch *gdbarch, struct type *valtype,
|
|||
when returned in general-purpose registers. */
|
||||
|
||||
static enum return_value_convention
|
||||
do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *type,
|
||||
struct regcache *regcache, gdb_byte *readbuf,
|
||||
const gdb_byte *writebuf, int broken_gcc)
|
||||
do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *func_type,
|
||||
struct type *type, struct regcache *regcache,
|
||||
gdb_byte *readbuf, const gdb_byte *writebuf,
|
||||
int broken_gcc)
|
||||
{
|
||||
struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
|
||||
enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
|
||||
int opencl_abi = 0;
|
||||
|
||||
if (func_type
|
||||
&& TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL)
|
||||
opencl_abi = 1;
|
||||
|
||||
gdb_assert (tdep->wordsize == 4);
|
||||
|
||||
if (TYPE_CODE (type) == TYPE_CODE_FLT
|
||||
&& TYPE_LENGTH (type) <= 8
|
||||
&& !tdep->soft_float)
|
||||
|
|
@ -691,6 +829,83 @@ do_ppc_sysv_return_value (struct gdbarch *gdbarch, struct type *type,
|
|||
}
|
||||
return RETURN_VALUE_REGISTER_CONVENTION;
|
||||
}
|
||||
/* OpenCL vectors < 16 bytes are returned as distinct
|
||||
scalars in f1..f2 or r3..r10. */
|
||||
if (TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
&& TYPE_LENGTH (type) < 16
|
||||
&& opencl_abi)
|
||||
{
|
||||
struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type));
|
||||
int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype);
|
||||
|
||||
for (i = 0; i < nelt; i++)
|
||||
{
|
||||
int offset = i * TYPE_LENGTH (eltype);
|
||||
|
||||
if (TYPE_CODE (eltype) == TYPE_CODE_FLT)
|
||||
{
|
||||
int regnum = tdep->ppc_fp0_regnum + 1 + i;
|
||||
gdb_byte regval[MAX_REGISTER_SIZE];
|
||||
struct type *regtype = register_type (gdbarch, regnum);
|
||||
|
||||
if (writebuf != NULL)
|
||||
{
|
||||
convert_typed_floating (writebuf + offset, eltype,
|
||||
regval, regtype);
|
||||
regcache_cooked_write (regcache, regnum, regval);
|
||||
}
|
||||
if (readbuf != NULL)
|
||||
{
|
||||
regcache_cooked_read (regcache, regnum, regval);
|
||||
convert_typed_floating (regval, regtype,
|
||||
readbuf + offset, eltype);
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
int regnum = tdep->ppc_gp0_regnum + 3 + i;
|
||||
ULONGEST regval;
|
||||
|
||||
if (writebuf != NULL)
|
||||
{
|
||||
regval = unpack_long (eltype, writebuf + offset);
|
||||
regcache_cooked_write_unsigned (regcache, regnum, regval);
|
||||
}
|
||||
if (readbuf != NULL)
|
||||
{
|
||||
regcache_cooked_read_unsigned (regcache, regnum, ®val);
|
||||
store_unsigned_integer (readbuf + offset,
|
||||
TYPE_LENGTH (eltype), byte_order,
|
||||
regval);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return RETURN_VALUE_REGISTER_CONVENTION;
|
||||
}
|
||||
/* OpenCL vectors >= 16 bytes are returned in v2..v9. */
|
||||
if (TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
&& TYPE_LENGTH (type) >= 16
|
||||
&& opencl_abi)
|
||||
{
|
||||
int n_regs = TYPE_LENGTH (type) / 16;
|
||||
int i;
|
||||
|
||||
for (i = 0; i < n_regs; i++)
|
||||
{
|
||||
int offset = i * 16;
|
||||
int regnum = tdep->ppc_vr0_regnum + 2 + i;
|
||||
|
||||
if (writebuf != NULL)
|
||||
regcache_cooked_write (regcache, regnum, writebuf + offset);
|
||||
if (readbuf != NULL)
|
||||
regcache_cooked_read (regcache, regnum, readbuf + offset);
|
||||
}
|
||||
|
||||
return RETURN_VALUE_REGISTER_CONVENTION;
|
||||
}
|
||||
if (TYPE_LENGTH (type) == 16
|
||||
&& TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
|
|
@ -826,8 +1041,8 @@ ppc_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type,
|
|||
struct type *valtype, struct regcache *regcache,
|
||||
gdb_byte *readbuf, const gdb_byte *writebuf)
|
||||
{
|
||||
return do_ppc_sysv_return_value (gdbarch, valtype, regcache, readbuf,
|
||||
writebuf, 0);
|
||||
return do_ppc_sysv_return_value (gdbarch, func_type, valtype, regcache,
|
||||
readbuf, writebuf, 0);
|
||||
}
|
||||
|
||||
enum return_value_convention
|
||||
|
|
@ -837,8 +1052,8 @@ ppc_sysv_abi_broken_return_value (struct gdbarch *gdbarch,
|
|||
struct regcache *regcache,
|
||||
gdb_byte *readbuf, const gdb_byte *writebuf)
|
||||
{
|
||||
return do_ppc_sysv_return_value (gdbarch, valtype, regcache, readbuf,
|
||||
writebuf, 1);
|
||||
return do_ppc_sysv_return_value (gdbarch, func_type, valtype, regcache,
|
||||
readbuf, writebuf, 1);
|
||||
}
|
||||
|
||||
/* The helper function for 64-bit SYSV push_dummy_call. Converts the
|
||||
|
|
@ -899,6 +1114,8 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch,
|
|||
CORE_ADDR func_addr = find_function_addr (function, NULL);
|
||||
struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
|
||||
enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
|
||||
struct type *ftype;
|
||||
int opencl_abi = 0;
|
||||
ULONGEST back_chain;
|
||||
/* See for-loop comment below. */
|
||||
int write_pass;
|
||||
|
|
@ -925,6 +1142,13 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch,
|
|||
regcache_cooked_read_unsigned (regcache, gdbarch_sp_regnum (gdbarch),
|
||||
&back_chain);
|
||||
|
||||
ftype = check_typedef (value_type (function));
|
||||
if (TYPE_CODE (ftype) == TYPE_CODE_PTR)
|
||||
ftype = check_typedef (TYPE_TARGET_TYPE (ftype));
|
||||
if (TYPE_CODE (ftype) == TYPE_CODE_FUNC
|
||||
&& TYPE_CALLING_CONVENTION (ftype) == DW_CC_GDB_IBM_OpenCL)
|
||||
opencl_abi = 1;
|
||||
|
||||
/* Go through the argument list twice.
|
||||
|
||||
Pass 1: Compute the function call's stack space and register
|
||||
|
|
@ -1133,6 +1357,109 @@ ppc64_sysv_abi_push_dummy_call (struct gdbarch *gdbarch,
|
|||
greg += 2;
|
||||
gparam = align_up (gparam + TYPE_LENGTH (type), tdep->wordsize);
|
||||
}
|
||||
else if (TYPE_LENGTH (type) < 16
|
||||
&& TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
&& opencl_abi)
|
||||
{
|
||||
/* OpenCL vectors shorter than 16 bytes are passed as if
|
||||
a series of independent scalars. */
|
||||
struct type *eltype = check_typedef (TYPE_TARGET_TYPE (type));
|
||||
int i, nelt = TYPE_LENGTH (type) / TYPE_LENGTH (eltype);
|
||||
|
||||
for (i = 0; i < nelt; i++)
|
||||
{
|
||||
const gdb_byte *elval = val + i * TYPE_LENGTH (eltype);
|
||||
|
||||
if (TYPE_CODE (eltype) == TYPE_CODE_FLT)
|
||||
{
|
||||
if (write_pass)
|
||||
{
|
||||
gdb_byte regval[MAX_REGISTER_SIZE];
|
||||
const gdb_byte *p;
|
||||
|
||||
if (TYPE_LENGTH (eltype) == 4)
|
||||
{
|
||||
memcpy (regval, elval, 4);
|
||||
memcpy (regval + 4, elval, 4);
|
||||
p = regval;
|
||||
}
|
||||
else
|
||||
p = elval;
|
||||
|
||||
write_memory (gparam, p, 8);
|
||||
|
||||
if (freg <= 13)
|
||||
{
|
||||
int regnum = tdep->ppc_fp0_regnum + freg;
|
||||
struct type *regtype
|
||||
= register_type (gdbarch, regnum);
|
||||
|
||||
convert_typed_floating (elval, eltype,
|
||||
regval, regtype);
|
||||
regcache_cooked_write (regcache, regnum, regval);
|
||||
}
|
||||
|
||||
if (greg <= 10)
|
||||
regcache_cooked_write (regcache,
|
||||
tdep->ppc_gp0_regnum + greg,
|
||||
regval);
|
||||
}
|
||||
|
||||
freg++;
|
||||
greg++;
|
||||
gparam = align_up (gparam + 8, tdep->wordsize);
|
||||
}
|
||||
else
|
||||
{
|
||||
if (write_pass)
|
||||
{
|
||||
ULONGEST word = unpack_long (eltype, elval);
|
||||
if (greg <= 10)
|
||||
regcache_cooked_write_unsigned
|
||||
(regcache, tdep->ppc_gp0_regnum + greg, word);
|
||||
|
||||
write_memory_unsigned_integer
|
||||
(gparam, tdep->wordsize, byte_order, word);
|
||||
}
|
||||
|
||||
greg++;
|
||||
gparam = align_up (gparam + TYPE_LENGTH (eltype),
|
||||
tdep->wordsize);
|
||||
}
|
||||
}
|
||||
}
|
||||
else if (TYPE_LENGTH (type) >= 16
|
||||
&& TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (type)
|
||||
&& opencl_abi)
|
||||
{
|
||||
/* OpenCL vectors 16 bytes or longer are passed as if
|
||||
a series of AltiVec vectors. */
|
||||
int i;
|
||||
|
||||
for (i = 0; i < TYPE_LENGTH (type) / 16; i++)
|
||||
{
|
||||
const gdb_byte *elval = val + i * 16;
|
||||
|
||||
gparam = align_up (gparam, 16);
|
||||
greg += greg & 1;
|
||||
|
||||
if (write_pass)
|
||||
{
|
||||
if (vreg <= 13)
|
||||
regcache_cooked_write (regcache,
|
||||
tdep->ppc_vr0_regnum + vreg,
|
||||
elval);
|
||||
|
||||
write_memory (gparam, elval, 16);
|
||||
}
|
||||
|
||||
greg += 2;
|
||||
vreg++;
|
||||
gparam += 16;
|
||||
}
|
||||
}
|
||||
else if (TYPE_LENGTH (type) == 16 && TYPE_VECTOR (type)
|
||||
&& TYPE_CODE (type) == TYPE_CODE_ARRAY
|
||||
&& tdep->ppc_vr0_regnum >= 0)
|
||||
|
|
@ -1358,6 +1685,11 @@ ppc64_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type,
|
|||
{
|
||||
struct gdbarch_tdep *tdep = gdbarch_tdep (gdbarch);
|
||||
enum bfd_endian byte_order = gdbarch_byte_order (gdbarch);
|
||||
int opencl_abi = 0;
|
||||
|
||||
if (func_type
|
||||
&& TYPE_CALLING_CONVENTION (func_type) == DW_CC_GDB_IBM_OpenCL)
|
||||
opencl_abi = 1;
|
||||
|
||||
/* This function exists to support a calling convention that
|
||||
requires floating-point registers. It shouldn't be used on
|
||||
|
|
@ -1420,6 +1752,83 @@ ppc64_sysv_abi_return_value (struct gdbarch *gdbarch, struct type *func_type,
|
|||
regcache_cooked_read (regcache, tdep->ppc_gp0_regnum + 3, readbuf);
|
||||
return RETURN_VALUE_REGISTER_CONVENTION;
|
||||
}
|
||||
/* OpenCL vectors < 16 bytes are returned as distinct
|
||||
scalars in f1..f2 or r3..r10. */
|
||||
if (TYPE_CODE (valtype) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (valtype)
|
||||
&& TYPE_LENGTH (valtype) < 16
|
||||
&& opencl_abi)
|
||||
{
|
||||
struct type *eltype = check_typedef (TYPE_TARGET_TYPE (valtype));
|
||||
int i, nelt = TYPE_LENGTH (valtype) / TYPE_LENGTH (eltype);
|
||||
|
||||
for (i = 0; i < nelt; i++)
|
||||
{
|
||||
int offset = i * TYPE_LENGTH (eltype);
|
||||
|
||||
if (TYPE_CODE (eltype) == TYPE_CODE_FLT)
|
||||
{
|
||||
int regnum = tdep->ppc_fp0_regnum + 1 + i;
|
||||
gdb_byte regval[MAX_REGISTER_SIZE];
|
||||
struct type *regtype = register_type (gdbarch, regnum);
|
||||
|
||||
if (writebuf != NULL)
|
||||
{
|
||||
convert_typed_floating (writebuf + offset, eltype,
|
||||
regval, regtype);
|
||||
regcache_cooked_write (regcache, regnum, regval);
|
||||
}
|
||||
if (readbuf != NULL)
|
||||
{
|
||||
regcache_cooked_read (regcache, regnum, regval);
|
||||
convert_typed_floating (regval, regtype,
|
||||
readbuf + offset, eltype);
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
int regnum = tdep->ppc_gp0_regnum + 3 + i;
|
||||
ULONGEST regval;
|
||||
|
||||
if (writebuf != NULL)
|
||||
{
|
||||
regval = unpack_long (eltype, writebuf + offset);
|
||||
regcache_cooked_write_unsigned (regcache, regnum, regval);
|
||||
}
|
||||
if (readbuf != NULL)
|
||||
{
|
||||
regcache_cooked_read_unsigned (regcache, regnum, ®val);
|
||||
store_unsigned_integer (readbuf + offset,
|
||||
TYPE_LENGTH (eltype), byte_order,
|
||||
regval);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return RETURN_VALUE_REGISTER_CONVENTION;
|
||||
}
|
||||
/* OpenCL vectors >= 16 bytes are returned in v2..v9. */
|
||||
if (TYPE_CODE (valtype) == TYPE_CODE_ARRAY
|
||||
&& TYPE_VECTOR (valtype)
|
||||
&& TYPE_LENGTH (valtype) >= 16
|
||||
&& opencl_abi)
|
||||
{
|
||||
int n_regs = TYPE_LENGTH (valtype) / 16;
|
||||
int i;
|
||||
|
||||
for (i = 0; i < n_regs; i++)
|
||||
{
|
||||
int offset = i * 16;
|
||||
int regnum = tdep->ppc_vr0_regnum + 2 + i;
|
||||
|
||||
if (writebuf != NULL)
|
||||
regcache_cooked_write (regcache, regnum, writebuf + offset);
|
||||
if (readbuf != NULL)
|
||||
regcache_cooked_read (regcache, regnum, readbuf + offset);
|
||||
}
|
||||
|
||||
return RETURN_VALUE_REGISTER_CONVENTION;
|
||||
}
|
||||
/* Array type has more than one use. */
|
||||
if (TYPE_CODE (valtype) == TYPE_CODE_ARRAY)
|
||||
{
|
||||
|
|
|
|||
|
|
@ -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:
|
||||
|
|
|
|||
|
|
@ -1,3 +1,9 @@
|
|||
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
|
||||
|
|
|
|||
|
|
@ -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 $@..."
|
||||
|
|
|
|||
218
gdb/testsuite/gdb.opencl/callfuncs.cl
Normal file
218
gdb/testsuite/gdb.opencl/callfuncs.cl
Normal file
|
|
@ -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;
|
||||
}
|
||||
102
gdb/testsuite/gdb.opencl/callfuncs.exp
Normal file
102
gdb/testsuite/gdb.opencl/callfuncs.exp
Normal file
|
|
@ -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}
|
||||
|
|
@ -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.
|
||||
|
|
|
|||
|
|
@ -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. */
|
||||
|
|
|
|||
Loading…
Reference in a new issue