diff options
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/ChangeLog | 11 | ||||
| -rw-r--r-- | gcc/builtins.c | 71 | ||||
| -rw-r--r-- | gcc/builtins.def | 4 | ||||
| -rw-r--r-- | gcc/doc/extend.texi | 10 | ||||
| -rw-r--r-- | gcc/fortran/ChangeLog | 5 | ||||
| -rw-r--r-- | gcc/fortran/f95-lang.c | 4 | ||||
| -rw-r--r-- | gcc/omp-builtins.def | 5 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size-2.c | 37 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size.c | 79 |
10 files changed, 232 insertions, 0 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 1c2a828..806f74c 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,14 @@ +2018-05-02 Tom de Vries <tom@codesourcery.com> + + PR libgomp/82428 + * builtins.def (DEF_GOACC_BUILTIN_ONLY): Define. + * omp-builtins.def (BUILT_IN_GOACC_PARLEVEL_ID) + (BUILT_IN_GOACC_PARLEVEL_SIZE): New builtin. + * builtins.c (expand_builtin_goacc_parlevel_id_size): New function. + (expand_builtin): Call expand_builtin_goacc_parlevel_id_size. + * doc/extend.texi (Other Builtins): Add __builtin_goacc_parlevel_id and + __builtin_goacc_parlevel_size. + 2018-05-02 Richard Biener <rguenther@suse.de> PR tree-optimization/85597 diff --git a/gcc/builtins.c b/gcc/builtins.c index a71555e..300e13c 100644 --- a/gcc/builtins.c +++ b/gcc/builtins.c @@ -71,6 +71,8 @@ along with GCC; see the file COPYING3. If not see #include "gimple-fold.h" #include "intl.h" #include "file-prefix-map.h" /* remap_macro_filename() */ +#include "gomp-constants.h" +#include "omp-general.h" struct target_builtins default_target_builtins; #if SWITCHABLE_TARGET @@ -6628,6 +6630,71 @@ expand_stack_save (void) return ret; } +/* Emit code to get the openacc gang, worker or vector id or size. */ + +static rtx +expand_builtin_goacc_parlevel_id_size (tree exp, rtx target, int ignore) +{ + const char *name; + rtx fallback_retval; + rtx_insn *(*gen_fn) (rtx, rtx); + switch (DECL_FUNCTION_CODE (get_callee_fndecl (exp))) + { + case BUILT_IN_GOACC_PARLEVEL_ID: + name = "__builtin_goacc_parlevel_id"; + fallback_retval = const0_rtx; + gen_fn = targetm.gen_oacc_dim_pos; + break; + case BUILT_IN_GOACC_PARLEVEL_SIZE: + name = "__builtin_goacc_parlevel_size"; + fallback_retval = const1_rtx; + gen_fn = targetm.gen_oacc_dim_size; + break; + default: + gcc_unreachable (); + } + + if (oacc_get_fn_attrib (current_function_decl) == NULL_TREE) + { + error ("%qs only supported in OpenACC code", name); + return const0_rtx; + } + + tree arg = CALL_EXPR_ARG (exp, 0); + if (TREE_CODE (arg) != INTEGER_CST) + { + error ("non-constant argument 0 to %qs", name); + return const0_rtx; + } + + int dim = TREE_INT_CST_LOW (arg); + switch (dim) + { + case GOMP_DIM_GANG: + case GOMP_DIM_WORKER: + case GOMP_DIM_VECTOR: + break; + default: + error ("illegal argument 0 to %qs", name); + return const0_rtx; + } + + if (ignore) + return target; + + if (!targetm.have_oacc_dim_size ()) + { + emit_move_insn (target, fallback_retval); + return target; + } + + rtx reg = MEM_P (target) ? gen_reg_rtx (GET_MODE (target)) : target; + emit_insn (gen_fn (reg, GEN_INT (dim))); + if (reg != target) + emit_move_insn (target, reg); + + return target; +} /* Expand an expression EXP that calls a built-in function, with result going to TARGET if that's convenient @@ -7758,6 +7825,10 @@ expand_builtin (tree exp, rtx target, rtx subtarget, machine_mode mode, folding. */ break; + case BUILT_IN_GOACC_PARLEVEL_ID: + case BUILT_IN_GOACC_PARLEVEL_SIZE: + return expand_builtin_goacc_parlevel_id_size (exp, target, ignore); + default: /* just do library call, if unknown builtin */ break; } diff --git a/gcc/builtins.def b/gcc/builtins.def index 17f825d..449d08d 100644 --- a/gcc/builtins.def +++ b/gcc/builtins.def @@ -214,6 +214,10 @@ along with GCC; see the file COPYING3. If not see #define DEF_GOACC_BUILTIN_COMPILER(ENUM, NAME, TYPE, ATTRS) \ DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \ flag_openacc, true, true, ATTRS, false, true) +#undef DEF_GOACC_BUILTIN_ONLY +#define DEF_GOACC_BUILTIN_ONLY(ENUM, NAME, TYPE, ATTRS) \ + DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, BT_LAST, \ + false, false, true, ATTRS, false, flag_openacc) #undef DEF_GOMP_BUILTIN #define DEF_GOMP_BUILTIN(ENUM, NAME, TYPE, ATTRS) \ DEF_BUILTIN (ENUM, "__builtin_" NAME, BUILT_IN_NORMAL, TYPE, TYPE, \ diff --git a/gcc/doc/extend.texi b/gcc/doc/extend.texi index 5571d05..f751b08 100644 --- a/gcc/doc/extend.texi +++ b/gcc/doc/extend.texi @@ -12437,6 +12437,16 @@ Aarch64. This function is mainly useful when writing inline assembly code. @end deftypefn +@deftypefn {Built-in Function} int __builtin_goacc_parlevel_id (int x) +Returns the openacc gang, worker or vector id depending on whether @var{x} is +0, 1 or 2. +@end deftypefn + +@deftypefn {Built-in Function} int __builtin_goacc_parlevel_size (int x) +Returns the openacc gang, worker or vector size depending on whether @var{x} is +0, 1 or 2. +@end deftypefn + @node Target Builtins @section Built-in Functions Specific to Particular Target Machines diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog index 46c234f..787f0ae 100644 --- a/gcc/fortran/ChangeLog +++ b/gcc/fortran/ChangeLog @@ -1,3 +1,8 @@ +2018-05-02 Tom de Vries <tom@codesourcery.com> + + PR libgomp/82428 + * f95-lang.c (DEF_GOACC_BUILTIN_ONLY): Define. + 2018-04-24 Steven G. Kargl <kargl@gcc.gnu.org> PR fortran/85520 diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c index 5fe34b2..0f39f0c 100644 --- a/gcc/fortran/f95-lang.c +++ b/gcc/fortran/f95-lang.c @@ -1202,6 +1202,10 @@ gfc_init_builtin_functions (void) #undef DEF_GOACC_BUILTIN_COMPILER #define DEF_GOACC_BUILTIN_COMPILER(code, name, type, attr) \ gfc_define_builtin (name, builtin_types[type], code, name, attr); +#undef DEF_GOACC_BUILTIN_ONLY +#define DEF_GOACC_BUILTIN_ONLY(code, name, type, attr) \ + gfc_define_builtin ("__builtin_" name, builtin_types[type], code, NULL, \ + attr); #undef DEF_GOMP_BUILTIN #define DEF_GOMP_BUILTIN(code, name, type, attr) /* ignore */ #include "../omp-builtins.def" diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index 8b5d663..55e5633 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -51,6 +51,11 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait", DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device", BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST) +DEF_GOACC_BUILTIN_ONLY (BUILT_IN_GOACC_PARLEVEL_ID, "goacc_parlevel_id", + BT_FN_INT_INT, ATTR_NOTHROW_LEAF_LIST) +DEF_GOACC_BUILTIN_ONLY (BUILT_IN_GOACC_PARLEVEL_SIZE, "goacc_parlevel_size", + BT_FN_INT_INT, ATTR_NOTHROW_LEAF_LIST) + DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_THREAD_NUM, "omp_get_thread_num", BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_THREADS, "omp_get_num_threads", diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 63750af..9c290eb 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2018-05-02 Tom de Vries <tom@codesourcery.com> + + PR libgomp/82428 + * c-c++-common/goacc/builtin-goacc-parlevel-id-size-2.c: New test. + * c-c++-common/goacc/builtin-goacc-parlevel-id-size.c: New test. + 2018-05-02 David Pagan <dave.pagan@oracle.com> PR c/30552 diff --git a/gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size-2.c b/gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size-2.c new file mode 100644 index 0000000..16c7b34 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size-2.c @@ -0,0 +1,37 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-O2" } */ + +#include "../../../../include/gomp-constants.h" + +void +foo (void) +{ + __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + /* { dg-error "'__builtin_goacc_parlevel_id' only supported in OpenACC code" "" { target *-*-* } .-1 } */ + + __builtin_goacc_parlevel_size (GOMP_DIM_GANG); + /* { dg-error "'__builtin_goacc_parlevel_size' only supported in OpenACC code" "" { target *-*-* } .-1 } */ +} + +#pragma acc routine +void +foo2 (int arg) +{ + __builtin_goacc_parlevel_id (arg); + /* { dg-error "non-constant argument 0 to '__builtin_goacc_parlevel_id'" "" { target *-*-* } .-1 } */ + + __builtin_goacc_parlevel_size (arg); + /* { dg-error "non-constant argument 0 to '__builtin_goacc_parlevel_size'" "" { target *-*-* } .-1 } */ + + __builtin_goacc_parlevel_id (-1); + /* { dg-error "illegal argument 0 to '__builtin_goacc_parlevel_id'" "" { target *-*-* } .-1 } */ + + __builtin_goacc_parlevel_id (-1); + /* { dg-error "illegal argument 0 to '__builtin_goacc_parlevel_id'" "" { target *-*-* } .-1 } */ + + __builtin_goacc_parlevel_size (-1); + /* { dg-error "illegal argument 0 to '__builtin_goacc_parlevel_size'" "" { target *-*-* } .-1 } */ + + __builtin_goacc_parlevel_size (3); + /* { dg-error "illegal argument 0 to '__builtin_goacc_parlevel_size'" "" { target *-*-* } .-1 } */ +} diff --git a/gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size.c b/gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size.c new file mode 100644 index 0000000..5cda818 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/builtin-goacc-parlevel-id-size.c @@ -0,0 +1,79 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-O2" } */ + +#include "../../../../include/gomp-constants.h" + +#pragma acc routine +int +foo (void) +{ + int res; + + __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + + __builtin_goacc_parlevel_size (GOMP_DIM_GANG); + __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); + __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); + + res += __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + res += __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + res += __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + + res += __builtin_goacc_parlevel_size (GOMP_DIM_GANG); + res += __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); + res += __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); + + return res; +} + +void +foo2 (void) +{ + int res; + +#pragma acc parallel + { + __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + + __builtin_goacc_parlevel_size (GOMP_DIM_GANG); + __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); + __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); + + res += __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + res += __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + res += __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + + res += __builtin_goacc_parlevel_size (GOMP_DIM_GANG); + res += __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); + res += __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); + } +} + +void +foo3 (void) +{ + int res; + +#pragma acc kernels + { + __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + + __builtin_goacc_parlevel_size (GOMP_DIM_GANG); + __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); + __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); + + res += __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + res += __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + res += __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + + res += __builtin_goacc_parlevel_size (GOMP_DIM_GANG); + res += __builtin_goacc_parlevel_size (GOMP_DIM_WORKER); + res += __builtin_goacc_parlevel_size (GOMP_DIM_VECTOR); + } +} |