This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC 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] | |
Hi!
In r220529, I have committed a merge from trunk r219682 (2015-01-15) into
gomp-4_0-branch. This is the trunk ÂMerge current set of OpenACC changes
From gomp-4_0-branch commit, which -- obviously -- mostly has been
present on gomp-4_0-branch already; here's the additional cleanup that I
merged in:
contrib/ChangeLog | 5 +
gcc/ChangeLog | 180 +++++++++++
gcc/ada/ChangeLog | 5 +
gcc/builtin-types.def | 5 +-
gcc/c-family/ChangeLog | 38 +++
gcc/c-family/c-omp.c | 1 -
gcc/c/ChangeLog | 51 +++
gcc/c/c-parser.c | 9 +-
gcc/cp/ChangeLog | 45 +++
gcc/cp/parser.c | 9 +-
gcc/doc/invoke.texi | 4 +
gcc/fortran/ChangeLog | 196 ++++++++++++
gcc/fortran/gfortran.texi | 75 +++--
gcc/fortran/intrinsic.texi | 56 ++--
gcc/fortran/invoke.texi | 4 +
gcc/fortran/types.def | 5 +-
gcc/lto/ChangeLog | 7 +
gcc/testsuite/ChangeLog | 130 ++++++++
gcc/testsuite/c-c++-common/goacc/asyncwait-1.c | 97 +-----
gcc/testsuite/c-c++-common/goacc/clauses-fail.c | 8 +-
gcc/testsuite/c-c++-common/goacc/data-2.c | 2 +-
include/ChangeLog | 6 +
libgomp/ChangeLog | 351 +++++++++++++++++++++
libgomp/Makefile.am | 6 +-
libgomp/Makefile.in | 12 +-
libgomp/env.c | 14 +-
libgomp/error.c | 2 +-
libgomp/libgomp-plugin.c | 9 +-
libgomp/libgomp-plugin.h | 10 +-
libgomp/libgomp.h | 67 ++--
libgomp/libgomp.map | 8 +-
libgomp/libgomp.texi | 5 +
libgomp/oacc-cuda.c | 2 +-
libgomp/oacc-host.c | 23 +-
libgomp/oacc-init.c | 18 +-
libgomp/oacc-int.h | 14 +-
libgomp/oacc-mem.c | 34 +-
libgomp/oacc-parallel.c | 1 -
libgomp/oacc-plugin.c | 2 +-
libgomp/oacc-plugin.h | 6 +-
libgomp/openacc.h | 131 ++++----
libgomp/plugin/plugin-host.c | 66 ++--
libgomp/plugin/plugin-nvptx.c | 16 +-
libgomp/splay-tree.c | 9 +-
libgomp/splay-tree.h | 21 --
libgomp/target.c | 32 +-
libgomp/testsuite/lib/libgomp.exp | 9 +-
.../{abort.c => abort-1.c} | 0
.../testsuite/libgomp.oacc-c-c++-common/data-3.c | 2 +-
libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c | 71 ++++-
.../testsuite/libgomp.oacc-c-c++-common/nested-2.c | 142 +++++++--
libgomp/testsuite/libgomp.oacc-fortran/lib-11.f90 | 82 -----
libgomp/testsuite/libgomp.oacc-fortran/lib-9.f90 | 83 -----
liboffloadmic/ChangeLog | 5 +
54 files changed, 1565 insertions(+), 626 deletions(-)
diff --git contrib/ChangeLog contrib/ChangeLog
index 1f02d95..f062ea9 100644
--- contrib/ChangeLog
+++ contrib/ChangeLog
@@ -1,3 +1,8 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gcc_update (files_and_dependencies): Update rules for new
+ libgomp/plugin/Makefrag.am and libgomp/plugin/configfrag.ac files.
+
2015-01-12 Yury Gribov <y.gribov@samsung.com>
* check_GNU_style.sh: Support patches coming from stdin.
diff --git gcc/ChangeLog gcc/ChangeLog
index d40a3f0..3a27df9 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,183 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ Bernd Schmidt <bernds@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Tom de Vries <tom@codesourcery.com>
+ Ilmir Usmanov <i.usmanov@samsung.com>
+ Dmitry Bocharnikov <dmitry.b@samsung.com>
+ Evgeny Gavrin <e.gavrin@samsung.com>
+ Jakub Jelinek <jakub@redhat.com>
+
+ * builtin-types.def (BT_FN_VOID_INT_INT_VAR)
+ (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR)
+ (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR):
+ New function types.
+ * builtins.c: Include "gomp-constants.h".
+ (expand_builtin_acc_on_device): New function.
+ (expand_builtin, is_inexpensive_builtin): Handle
+ BUILT_IN_ACC_ON_DEVICE.
+ * builtins.def (DEF_GOACC_BUILTIN, DEF_GOACC_BUILTIN_COMPILER):
+ New macros.
+ * cgraph.c (cgraph_node::create): Consider flag_openacc next to
+ flag_openmp.
+ * config.gcc <nvptx-*> (tm_file): Add nvptx/offload.h.
+ <*-intelmic-* | *-intelmicemul-*> (tm_file): Add
+ i386/intelmic-offload.h.
+ * gcc.c (LINK_COMMAND_SPEC, GOMP_SELF_SPECS): For -fopenacc, link
+ to libgomp and its dependencies.
+ * config/arc/arc.h (LINK_COMMAND_SPEC): Likewise.
+ * config/darwin.h (LINK_COMMAND_SPEC_A): Likewise.
+ * config/i386/mingw32.h (GOMP_SELF_SPECS): Likewise.
+ * config/ia64/hpux.h (LIB_SPEC): Likewise.
+ * config/pa/pa-hpux11.h (LIB_SPEC): Likewise.
+ * config/pa/pa64-hpux.h (LIB_SPEC): Likewise.
+ * doc/generic.texi: Update for OpenACC changes.
+ * doc/gimple.texi: Likewise.
+ * doc/invoke.texi: Likewise.
+ * doc/sourcebuild.texi: Likewise.
+ * gimple-pretty-print.c (dump_gimple_omp_for): Handle
+ GF_OMP_FOR_KIND_OACC_LOOP.
+ (dump_gimple_omp_target): Handle GF_OMP_TARGET_KIND_OACC_KERNELS,
+ GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_DATA,
+ GF_OMP_TARGET_KIND_OACC_UPDATE,
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
+ Dump more data.
+ * gimple.c: Update comments for OpenACC changes.
+ * gimple.def: Likewise.
+ * gimple.h: Likewise.
+ (enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP,
+ GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_KERNELS,
+ GF_OMP_TARGET_KIND_OACC_DATA, GF_OMP_TARGET_KIND_OACC_UPDATE,
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
+ (gimple_omp_for_cond, gimple_omp_for_set_cond): Sort in the
+ appropriate place.
+ (is_gimple_omp_oacc, is_gimple_omp_offloaded): New functions.
+ * gimplify.c: Include "gomp-constants.h".
+ Update comments for OpenACC changes.
+ (is_gimple_stmt): Handle OACC_PARALLEL, OACC_KERNELS, OACC_DATA,
+ OACC_HOST_DATA, OACC_DECLARE, OACC_UPDATE, OACC_ENTER_DATA,
+ OACC_EXIT_DATA, OACC_CACHE, OACC_LOOP.
+ (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
+ OMP_CLAUSE__CACHE_, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
+ OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
+ OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER,
+ OMP_CLAUSE_VECTOR, OMP_CLAUSE_DEVICE_RESIDENT,
+ OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO,
+ OMP_CLAUSE_SEQ.
+ (gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Use
+ GOMP_MAP_* instead of OMP_CLAUSE_MAP_*. Use
+ OMP_CLAUSE_SET_MAP_KIND.
+ (gimplify_oacc_cache): New function.
+ (gimplify_omp_for): Handle OACC_LOOP.
+ (gimplify_omp_workshare): Handle OACC_KERNELS, OACC_PARALLEL,
+ OACC_DATA.
+ (gimplify_omp_target_update): Handle OACC_ENTER_DATA,
+ OACC_EXIT_DATA, OACC_UPDATE.
+ (gimplify_expr): Handle OACC_LOOP, OACC_CACHE, OACC_HOST_DATA,
+ OACC_DECLARE, OACC_KERNELS, OACC_PARALLEL, OACC_DATA,
+ OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE.
+ (gimplify_body): Consider flag_openacc next to flag_openmp.
+ * lto-streamer-out.c: Include "gomp-constants.h".
+ * omp-builtins.def (BUILT_IN_ACC_GET_DEVICE_TYPE)
+ (BUILT_IN_GOACC_DATA_START, BUILT_IN_GOACC_DATA_END)
+ (BUILT_IN_GOACC_ENTER_EXIT_DATA, BUILT_IN_GOACC_PARALLEL)
+ (BUILT_IN_GOACC_UPDATE, BUILT_IN_GOACC_WAIT)
+ (BUILT_IN_GOACC_GET_THREAD_NUM, BUILT_IN_GOACC_GET_NUM_THREADS)
+ (BUILT_IN_ACC_ON_DEVICE): New builtins.
+ * omp-low.c: Include "gomp-constants.h".
+ Update comments for OpenACC changes.
+ (struct omp_context): Add reduction_map, gwv_below, gwv_this
+ members.
+ (extract_omp_for_data, use_pointer_for_field, install_var_field)
+ (new_omp_context, delete_omp_context, scan_sharing_clauses)
+ (create_omp_child_function, scan_omp_for, scan_omp_target)
+ (check_omp_nesting_restrictions, lower_reduction_clauses)
+ (build_omp_regions_1, diagnose_sb_0, make_gimple_omp_edges):
+ Update for OpenACC changes.
+ (scan_sharing_clauses): Handle OMP_CLAUSE_NUM_GANGS:
+ OMP_CLAUSE_NUM_WORKERS: OMP_CLAUSE_VECTOR_LENGTH,
+ OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT, OMP_CLAUSE_GANG,
+ OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_DEVICE_RESIDENT,
+ OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE__CACHE_, OMP_CLAUSE_INDEPENDENT,
+ OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ. Use GOMP_MAP_* instead of
+ OMP_CLAUSE_MAP_*.
+ (expand_omp_for_static_nochunk, expand_omp_for_static_chunk):
+ Handle GF_OMP_FOR_KIND_OACC_LOOP.
+ (expand_omp_target, lower_omp_target): Handle
+ GF_OMP_TARGET_KIND_OACC_PARALLEL, GF_OMP_TARGET_KIND_OACC_KERNELS,
+ GF_OMP_TARGET_KIND_OACC_UPDATE,
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA,
+ GF_OMP_TARGET_KIND_OACC_DATA.
+ (pass_expand_omp::execute, execute_lower_omp)
+ (pass_diagnose_omp_blocks::gate): Consider flag_openacc next to
+ flag_openmp.
+ (offload_symbol_decl): New variable.
+ (oacc_get_reduction_array_id, oacc_max_threads)
+ (get_offload_symbol_decl, get_base_type, lookup_oacc_reduction)
+ (maybe_lookup_oacc_reduction, enclosing_target_ctx)
+ (oacc_loop_or_target_p, oacc_lower_reduction_var_helper)
+ (oacc_gimple_assign, oacc_initialize_reduction_data)
+ (oacc_finalize_reduction_data, oacc_process_reduction_data): New
+ functions.
+ (is_targetreg_ctx): Remove function.
+ * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__CACHE_,
+ OMP_CLAUSE_DEVICE_RESIDENT, OMP_CLAUSE_USE_DEVICE,
+ OMP_CLAUSE_GANG, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
+ OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ, OMP_CLAUSE_INDEPENDENT,
+ OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_NUM_GANGS,
+ OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH.
+ * tree.c (omp_clause_code_name, walk_tree_1): Update accordingly.
+ * tree.h (OMP_CLAUSE_GANG_EXPR, OMP_CLAUSE_GANG_STATIC_EXPR)
+ (OMP_CLAUSE_ASYNC_EXPR, OMP_CLAUSE_WAIT_EXPR)
+ (OMP_CLAUSE_VECTOR_EXPR, OMP_CLAUSE_WORKER_EXPR)
+ (OMP_CLAUSE_NUM_GANGS_EXPR, OMP_CLAUSE_NUM_WORKERS_EXPR)
+ (OMP_CLAUSE_VECTOR_LENGTH_EXPR): New macros.
+ * tree-core.h: Update comments for OpenACC changes.
+ (enum omp_clause_map_kind): Remove.
+ (struct tree_omp_clause): Change type of map_kind member from enum
+ omp_clause_map_kind to unsigned char.
+ * tree-inline.c: Update comments for OpenACC changes.
+ * tree-nested.c: Likewise. Include "gomp-constants.h".
+ (convert_nonlocal_reference_stmt, convert_local_reference_stmt)
+ (convert_tramp_reference_stmt, convert_gimple_call): Update for
+ OpenACC changes. Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*. Use
+ OMP_CLAUSE_SET_MAP_KIND.
+ * tree-pretty-print.c: Include "gomp-constants.h".
+ (dump_omp_clause): Handle OMP_CLAUSE_DEVICE_RESIDENT,
+ OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE__CACHE_, OMP_CLAUSE_GANG,
+ OMP_CLAUSE_ASYNC, OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ,
+ OMP_CLAUSE_WAIT, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR,
+ OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
+ OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_INDEPENDENT. Use GOMP_MAP_*
+ instead of OMP_CLAUSE_MAP_*.
+ (dump_generic_node): Handle OACC_PARALLEL, OACC_KERNELS,
+ OACC_DATA, OACC_HOST_DATA, OACC_DECLARE, OACC_UPDATE,
+ OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_CACHE, OACC_LOOP.
+ * tree-streamer-in.c: Include "gomp-constants.h".
+ (unpack_ts_omp_clause_value_fields) Use GOMP_MAP_* instead of
+ OMP_CLAUSE_MAP_*. Use OMP_CLAUSE_SET_MAP_KIND.
+ * tree-streamer-out.c: Include "gomp-constants.h".
+ (pack_ts_omp_clause_value_fields): Use GOMP_MAP_* instead of
+ OMP_CLAUSE_MAP_*.
+ * tree.def (OACC_PARALLEL, OACC_KERNELS, OACC_DATA)
+ (OACC_HOST_DATA, OACC_LOOP, OACC_CACHE, OACC_DECLARE)
+ (OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_UPDATE): New tree codes.
+ * tree.c (omp_clause_num_ops): Update accordingly.
+ * tree.h (OMP_BODY, OMP_CLAUSES, OMP_LOOP_CHECK, OMP_CLAUSE_SIZE):
+ Likewise.
+ (OACC_PARALLEL_BODY, OACC_PARALLEL_CLAUSES, OACC_KERNELS_BODY)
+ (OACC_KERNELS_CLAUSES, OACC_DATA_BODY, OACC_DATA_CLAUSES)
+ (OACC_HOST_DATA_BODY, OACC_HOST_DATA_CLAUSES, OACC_CACHE_CLAUSES)
+ (OACC_DECLARE_CLAUSES, OACC_ENTER_DATA_CLAUSES)
+ (OACC_EXIT_DATA_CLAUSES, OACC_UPDATE_CLAUSES)
+ (OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): New macros.
+ * tree.h (OMP_CLAUSE_MAP_KIND): Cast it to enum gomp_map_kind.
+ (OMP_CLAUSE_SET_MAP_KIND): New macro.
+ * varpool.c (varpool_node::get_create): Consider flag_openacc next
+ to flag_openmp.
+ * config/i386/intelmic-offload.h: New file.
+ * config/nvptx/offload.h: Likewise.
+
2015-01-15 Prathamesh Kulkarni <prathamesh.kulkarni@linaro.org>
* explow.h: Remove duplicate contents.
diff --git gcc/ada/ChangeLog gcc/ada/ChangeLog
index 93efb49..c130f7d 100644
--- gcc/ada/ChangeLog
+++ gcc/ada/ChangeLog
@@ -1,3 +1,8 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gcc-interface/utils.c (DEF_FUNCTION_TYPE_VAR_8)
+ (DEF_FUNCTION_TYPE_VAR_12): New macros.
+
2015-01-09 Michael Collison <michael.collison@linaro.org>
* gcc-interface/cuintp.c: Include hash-set.h, machmode.h,
diff --git gcc/builtin-types.def gcc/builtin-types.def
index 8ab4300..3412677 100644
--- gcc/builtin-types.def
+++ gcc/builtin-types.def
@@ -593,8 +593,9 @@ DEF_FUNCTION_TYPE_VAR_8 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
BT_PTR, BT_INT, BT_INT)
DEF_FUNCTION_TYPE_VAR_12 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
- BT_PTR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+ BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT, BT_INT,
+ BT_INT, BT_INT)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_VAR, BT_FN_VOID_VAR)
DEF_FUNCTION_TYPE_3 (BT_FN_PTR_PTR_FN_VOID_VAR_PTR_SIZE,
diff --git gcc/c-family/ChangeLog gcc/c-family/ChangeLog
index 9764045..e0ad215 100644
--- gcc/c-family/ChangeLog
+++ gcc/c-family/ChangeLog
@@ -1,3 +1,41 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ Bernd Schmidt <bernds@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ Ilmir Usmanov <i.usmanov@samsung.com>
+ Jakub Jelinek <jakub@redhat.com>
+
+ * c.opt (fopenacc): New option.
+ * c-cppbuiltin.c (c_cpp_builtins): Conditionally define _OPENACC.
+ * c-common.c (DEF_FUNCTION_TYPE_VAR_8, DEF_FUNCTION_TYPE_VAR_12):
+ New macros.
+ * c-common.h (c_finish_oacc_wait): New prototype.
+ * c-omp.c: Include "omp-low.h" and "gomp-constants.h".
+ (c_finish_oacc_wait): New function.
+ * c-pragma.c (oacc_pragmas): New variable.
+ (c_pp_lookup_pragma, init_pragma): Handle it.
+ * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_CACHE,
+ PRAGMA_OACC_DATA, PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA,
+ PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL,
+ PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
+ (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ASYNC,
+ PRAGMA_OACC_CLAUSE_AUTO, PRAGMA_OACC_CLAUSE_COLLAPSE,
+ PRAGMA_OACC_CLAUSE_COPY, PRAGMA_OACC_CLAUSE_COPYIN,
+ PRAGMA_OACC_CLAUSE_COPYOUT, PRAGMA_OACC_CLAUSE_CREATE,
+ PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICE,
+ PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_FIRSTPRIVATE,
+ PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST,
+ PRAGMA_OACC_CLAUSE_IF, PRAGMA_OACC_CLAUSE_NUM_GANGS,
+ PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_PRIVATE,
+ PRAGMA_OACC_CLAUSE_REDUCTION, PRAGMA_OACC_CLAUSE_SELF,
+ PRAGMA_OACC_CLAUSE_SEQ, PRAGMA_OACC_CLAUSE_VECTOR,
+ PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT,
+ PRAGMA_OACC_CLAUSE_WORKER.
+
2015-01-14 Marcos Diaz <marcos.diaz@tallertechnologies.com>
* c-cppbuiltin.c (c_cpp_builtins): New cpp define __SSP_EXPLICIT__
diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c
index eea1e3a..8715045 100644
--- gcc/c-family/c-omp.c
+++ gcc/c-family/c-omp.c
@@ -714,7 +714,6 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
enum c_omp_clause_split s;
int i;
- gcc_assert (code != OACC_PARALLEL);
for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++)
cclauses[i] = NULL;
/* Add implicit nowait clause on
diff --git gcc/c/ChangeLog gcc/c/ChangeLog
index f42f53b..4652409 100644
--- gcc/c/ChangeLog
+++ gcc/c/ChangeLog
@@ -1,3 +1,54 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ Bernd Schmidt <bernds@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Jakub Jelinek <jakub@redhat.com>
+ Ilmir Usmanov <i.usmanov@samsung.com>
+
+ * c-parser.c: Include "gomp-constants.h".
+ (c_parser_omp_clause_map): Use enum gomp_map_kind instead of enum
+ omp_clause_map_kind. Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
+ Use OMP_CLAUSE_SET_MAP_KIND.
+ (c_parser_pragma): Handle PRAGMA_OACC_ENTER_DATA,
+ PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_UPDATE.
+ (c_parser_omp_construct): Handle PRAGMA_OACC_CACHE,
+ PRAGMA_OACC_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP,
+ PRAGMA_OACC_PARALLEL, PRAGMA_OACC_WAIT.
+ (c_parser_omp_clause_name): Handle "auto", "async", "copy",
+ "copyout", "create", "delete", "deviceptr", "gang", "host",
+ "num_gangs", "num_workers", "present", "present_or_copy", "pcopy",
+ "present_or_copyin", "pcopyin", "present_or_copyout", "pcopyout",
+ "present_or_create", "pcreate", "seq", "self", "vector",
+ "vector_length", "wait", "worker".
+ (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
+ (OACC_ENTER_DATA_CLAUSE_MASK, OACC_EXIT_DATA_CLAUSE_MASK)
+ (OACC_LOOP_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK)
+ (OACC_UPDATE_CLAUSE_MASK, OACC_WAIT_CLAUSE_MASK): New macros.
+ (c_parser_omp_variable_list): Handle OMP_CLAUSE__CACHE_.
+ (c_parser_oacc_wait_list, c_parser_oacc_data_clause)
+ (c_parser_oacc_data_clause_deviceptr)
+ (c_parser_omp_clause_num_gangs, c_parser_omp_clause_num_workers)
+ (c_parser_oacc_clause_async, c_parser_oacc_clause_wait)
+ (c_parser_omp_clause_vector_length, c_parser_oacc_all_clauses)
+ (c_parser_oacc_cache, c_parser_oacc_data, c_parser_oacc_kernels)
+ (c_parser_oacc_enter_exit_data, c_parser_oacc_loop)
+ (c_parser_oacc_parallel, c_parser_oacc_update)
+ (c_parser_oacc_wait): New functions.
+ * c-tree.h (c_finish_oacc_parallel, c_finish_oacc_kernels)
+ (c_finish_oacc_data): New prototypes.
+ * c-typeck.c: Include "gomp-constants.h".
+ (handle_omp_array_sections): Handle GOMP_MAP_FORCE_DEVICEPTR. Use
+ GOMP_MAP_* instead of OMP_CLAUSE_MAP_*. Use
+ OMP_CLAUSE_SET_MAP_KIND.
+ (c_finish_oacc_parallel, c_finish_oacc_kernels)
+ (c_finish_oacc_data): New functions.
+ (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_,
+ OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
+ OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT,
+ OMP_CLAUSE_AUTO, OMP_CLAUSE_SEQ, OMP_CLAUSE_GANG,
+ OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, and OMP_CLAUSE_MAP's
+ GOMP_MAP_FORCE_DEVICEPTR.
+
2015-01-09 Michael Collison <michael.collison@linaro.org>
* c-array-notation.c: Include hash-set.h, machmode.h,
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index fc6661b..665ee42 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1250,11 +1250,11 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
vec<tree, va_gc> **, location_t *,
tree *, vec<location_t> *,
unsigned int * = NULL);
+static void c_parser_oacc_enter_exit_data (c_parser *, bool);
+static void c_parser_oacc_update (c_parser *);
static tree c_parser_oacc_loop (location_t, c_parser *, char *);
static void c_parser_omp_construct (c_parser *);
static void c_parser_omp_threadprivate (c_parser *);
-static void c_parser_oacc_enter_exit_data (c_parser *, bool);
-static void c_parser_oacc_update (c_parser *);
static void c_parser_omp_barrier (c_parser *);
static void c_parser_omp_flush (c_parser *);
static tree c_parser_omp_for_loop (location_t, c_parser *, enum tree_code,
@@ -11699,7 +11699,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
c_name = "wait";
break;
default:
- c_parser_error (parser, "expected clause");
+ c_parser_error (parser, "expected %<#pragma acc%> clause");
goto saw_error;
}
@@ -11928,7 +11928,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
c_name = "simdlen";
break;
default:
- c_parser_error (parser, "expected clause");
+ c_parser_error (parser, "expected %<#pragma omp%> clause");
goto saw_error;
}
@@ -13046,7 +13046,6 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
if (cclauses != NULL
&& cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL] != NULL)
{
- gcc_assert (code != OACC_LOOP);
tree *c;
for (c = &cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL]; *c ; )
if (OMP_CLAUSE_CODE (*c) != OMP_CLAUSE_FIRSTPRIVATE
diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog
index 02b4fac..543f4d9 100644
--- gcc/cp/ChangeLog
+++ gcc/cp/ChangeLog
@@ -1,3 +1,48 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ Ilmir Usmanov <i.usmanov@samsung.com>
+ Jakub Jelinek <jakub@redhat.com>
+
+ * parser.c: Include "gomp-constants.h".
+ (cp_parser_omp_clause_map): Use enum gomp_map_kind instead of enum
+ omp_clause_map_kind. Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
+ Use OMP_CLAUSE_SET_MAP_KIND.
+ (cp_parser_omp_construct, cp_parser_pragma): Handle
+ PRAGMA_OACC_CACHE, PRAGMA_OACC_DATA, PRAGMA_OACC_ENTER_DATA,
+ PRAGMA_OACC_EXIT_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_PARALLEL,
+ PRAGMA_OACC_LOOP, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT.
+ (cp_parser_omp_clause_name): Handle "async", "copy", "copyout",
+ "create", "delete", "deviceptr", "host", "num_gangs",
+ "num_workers", "present", "present_or_copy", "pcopy",
+ "present_or_copyin", "pcopyin", "present_or_copyout", "pcopyout",
+ "present_or_create", "pcreate", "vector_length", "wait".
+ (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
+ (OACC_EXIT_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
+ (OACC_LOOP_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK)
+ (OACC_UPDATE_CLAUSE_MASK, OACC_WAIT_CLAUSE_MASK): New macros.
+ (cp_parser_omp_var_list_no_open): Handle OMP_CLAUSE__CACHE_.
+ (cp_parser_oacc_data_clause, cp_parser_oacc_data_clause_deviceptr)
+ (cp_parser_oacc_clause_vector_length, cp_parser_oacc_wait_list)
+ (cp_parser_oacc_clause_wait, cp_parser_omp_clause_num_gangs)
+ (cp_parser_omp_clause_num_workers, cp_parser_oacc_clause_async)
+ (cp_parser_oacc_all_clauses, cp_parser_oacc_cache)
+ (cp_parser_oacc_data, cp_parser_oacc_enter_exit_data)
+ (cp_parser_oacc_kernels, cp_parser_oacc_loop)
+ (cp_parser_oacc_parallel, cp_parser_oacc_update)
+ (cp_parser_oacc_wait): New functions.
+ * cp-tree.h (finish_oacc_data, finish_oacc_kernels)
+ (finish_oacc_parallel): New prototypes.
+ * semantics.c: Include "gomp-constants.h".
+ (handle_omp_array_sections): Handle GOMP_MAP_FORCE_DEVICEPTR. Use
+ GOMP_MAP_* instead of OMP_CLAUSE_MAP_*. Use
+ OMP_CLAUSE_SET_MAP_KIND.
+ (finish_omp_clauses): Handle OMP_CLAUSE_ASYNC,
+ OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_WAIT, OMP_CLAUSE__CACHE_.
+ Use GOMP_MAP_* instead of OMP_CLAUSE_MAP_*.
+ (finish_oacc_data, finish_oacc_kernels, finish_oacc_parallel): New
+ functions.
+
2015-01-14 Paolo Carlini <paolo.carlini@oracle.com>
PR c++/58671
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 1ce7de7..bfa3d81 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -27653,8 +27653,6 @@ cp_parser_omp_clause_name (cp_parser *parser)
else if (!strcmp ("present_or_create", p)
|| !strcmp ("pcreate", p))
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
- else if (!strcmp ("private", p))
- result = PRAGMA_OMP_CLAUSE_PRIVATE;
else if (!strcmp ("proc_bind", p))
result = PRAGMA_OMP_CLAUSE_PROC_BIND;
break;
@@ -29236,8 +29234,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "self";
break;
case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
- clauses =
- cp_parser_oacc_clause_vector_length (parser, clauses);
+ clauses = cp_parser_oacc_clause_vector_length (parser, clauses);
c_name = "vector_length";
break;
case PRAGMA_OACC_CLAUSE_WAIT:
@@ -29245,7 +29242,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "wait";
break;
default:
- cp_parser_error (parser, "expected clause");
+ cp_parser_error (parser, "expected %<#pragma acc%> clause");
goto saw_error;
}
@@ -29496,7 +29493,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "simdlen";
break;
default:
- cp_parser_error (parser, "expected clause");
+ cp_parser_error (parser, "expected %<#pragma omp%> clause");
goto saw_error;
}
diff --git gcc/doc/invoke.texi gcc/doc/invoke.texi
index 112eb81..510201a 100644
--- gcc/doc/invoke.texi
+++ gcc/doc/invoke.texi
@@ -1895,6 +1895,10 @@ Programming Interface v2.0 @w{@uref{http://www.openacc.org/}}. This option
implies @option{-pthread}, and thus is only supported on targets that
have support for @option{-pthread}.
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC. See
+@w{@uref{https://gcc.gnu.org/wiki/OpenACC}} for more information.
+
@item -fopenmp
@opindex fopenmp
@cindex OpenMP parallel
diff --git gcc/fortran/ChangeLog gcc/fortran/ChangeLog
index df4a2f3..d8b72a2 100644
--- gcc/fortran/ChangeLog
+++ gcc/fortran/ChangeLog
@@ -1,3 +1,199 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Ilmir Usmanov <i.usmanov@samsung.com>
+ Tobias Burnus <burnus@net-b.de>
+
+ * lang.opt (fopenacc): New option.
+ * cpp.c (cpp_define_builtins): Conditionally define _OPENACC.
+ * dump-parse-tree.c (show_omp_node): Split part of it into...
+ (show_omp_clauses): ... this new function.
+ (show_omp_node, show_code_node): Handle EXEC_OACC_PARALLEL_LOOP,
+ EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS,
+ EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
+ EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
+ EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
+ (show_namespace): Update for OpenACC.
+ * f95-lang.c (DEF_FUNCTION_TYPE_VAR_2, DEF_FUNCTION_TYPE_VAR_8)
+ (DEF_FUNCTION_TYPE_VAR_12, DEF_GOACC_BUILTIN)
+ (DEF_GOACC_BUILTIN_COMPILER): New macros.
+ * types.def (BT_FN_VOID_INT_INT_VAR)
+ (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR)
+ (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR):
+ New function types.
+ * gfortran.h (gfc_statement): Add ST_OACC_PARALLEL_LOOP,
+ ST_OACC_END_PARALLEL_LOOP, ST_OACC_PARALLEL, ST_OACC_END_PARALLEL,
+ ST_OACC_KERNELS, ST_OACC_END_KERNELS, ST_OACC_DATA,
+ ST_OACC_END_DATA, ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA,
+ ST_OACC_LOOP, ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE,
+ ST_OACC_WAIT, ST_OACC_CACHE, ST_OACC_KERNELS_LOOP,
+ ST_OACC_END_KERNELS_LOOP, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA,
+ ST_OACC_ROUTINE.
+ (struct gfc_expr_list): New data type.
+ (gfc_get_expr_list): New macro.
+ (gfc_omp_map_op): Add OMP_MAP_FORCE_ALLOC, OMP_MAP_FORCE_DEALLOC,
+ OMP_MAP_FORCE_TO, OMP_MAP_FORCE_FROM, OMP_MAP_FORCE_TOFROM,
+ OMP_MAP_FORCE_PRESENT, OMP_MAP_FORCE_DEVICEPTR.
+ (OMP_LIST_FIRST, OMP_LIST_DEVICE_RESIDENT, OMP_LIST_USE_DEVICE)
+ (OMP_LIST_CACHE): New enumerators.
+ (struct gfc_omp_clauses): Add async_expr, gang_expr, worker_expr,
+ vector_expr, num_gangs_expr, num_workers_expr, vector_length_expr,
+ wait_list, tile_list, async, gang, worker, vector, seq,
+ independent, wait, par_auto, gang_static, and loc members.
+ (struct gfc_namespace): Add oacc_declare_clauses member.
+ (gfc_exec_op): Add EXEC_OACC_KERNELS_LOOP,
+ EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS,
+ EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
+ EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
+ EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
+ (gfc_free_expr_list, gfc_resolve_oacc_directive)
+ (gfc_resolve_oacc_declare, gfc_resolve_oacc_parallel_loop_blocks)
+ (gfc_resolve_oacc_blocks): New prototypes.
+ * match.c (match_exit_cycle): Handle EXEC_OACC_LOOP and
+ EXEC_OACC_PARALLEL_LOOP.
+ * match.h (gfc_match_oacc_cache, gfc_match_oacc_wait)
+ (gfc_match_oacc_update, gfc_match_oacc_declare)
+ (gfc_match_oacc_loop, gfc_match_oacc_host_data)
+ (gfc_match_oacc_data, gfc_match_oacc_kernels)
+ (gfc_match_oacc_kernels_loop, gfc_match_oacc_parallel)
+ (gfc_match_oacc_parallel_loop, gfc_match_oacc_enter_data)
+ (gfc_match_oacc_exit_data, gfc_match_oacc_routine): New
+ prototypes.
+ * openmp.c: Include "diagnostic.h" and "gomp-constants.h".
+ (gfc_free_omp_clauses): Update for members added to struct
+ gfc_omp_clauses.
+ (gfc_match_omp_clauses): Change mask paramter to uint64_t. Add
+ openacc parameter.
+ (resolve_omp_clauses): Add openacc parameter. Update for OpenACC.
+ (struct fortran_omp_context): Add is_openmp member.
+ (gfc_resolve_omp_parallel_blocks): Initialize it.
+ (gfc_resolve_do_iterator): Update for OpenACC.
+ (gfc_resolve_omp_directive): Call
+ resolve_omp_directive_inside_oacc_region.
+ (OMP_CLAUSE_PRIVATE, OMP_CLAUSE_FIRSTPRIVATE)
+ (OMP_CLAUSE_LASTPRIVATE, OMP_CLAUSE_COPYPRIVATE)
+ (OMP_CLAUSE_SHARED, OMP_CLAUSE_COPYIN, OMP_CLAUSE_REDUCTION)
+ (OMP_CLAUSE_IF, OMP_CLAUSE_NUM_THREADS, OMP_CLAUSE_SCHEDULE)
+ (OMP_CLAUSE_DEFAULT, OMP_CLAUSE_ORDERED, OMP_CLAUSE_COLLAPSE)
+ (OMP_CLAUSE_UNTIED, OMP_CLAUSE_FINAL, OMP_CLAUSE_MERGEABLE)
+ (OMP_CLAUSE_ALIGNED, OMP_CLAUSE_DEPEND, OMP_CLAUSE_INBRANCH)
+ (OMP_CLAUSE_LINEAR, OMP_CLAUSE_NOTINBRANCH, OMP_CLAUSE_PROC_BIND)
+ (OMP_CLAUSE_SAFELEN, OMP_CLAUSE_SIMDLEN, OMP_CLAUSE_UNIFORM)
+ (OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP, OMP_CLAUSE_TO)
+ (OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_THREAD_LIMIT)
+ (OMP_CLAUSE_DIST_SCHEDULE): Use uint64_t.
+ (OMP_CLAUSE_ASYNC, OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS)
+ (OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_COPY, OMP_CLAUSE_COPYOUT)
+ (OMP_CLAUSE_CREATE, OMP_CLAUSE_PRESENT)
+ (OMP_CLAUSE_PRESENT_OR_COPY, OMP_CLAUSE_PRESENT_OR_COPYIN)
+ (OMP_CLAUSE_PRESENT_OR_COPYOUT, OMP_CLAUSE_PRESENT_OR_CREATE)
+ (OMP_CLAUSE_DEVICEPTR, OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER)
+ (OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ, OMP_CLAUSE_INDEPENDENT)
+ (OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_DEVICE_RESIDENT)
+ (OMP_CLAUSE_HOST_SELF, OMP_CLAUSE_OACC_DEVICE, OMP_CLAUSE_WAIT)
+ (OMP_CLAUSE_DELETE, OMP_CLAUSE_AUTO, OMP_CLAUSE_TILE): New macros.
+ (gfc_match_omp_clauses): Handle those.
+ (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES)
+ (OACC_LOOP_CLAUSES, OACC_PARALLEL_LOOP_CLAUSES)
+ (OACC_KERNELS_LOOP_CLAUSES, OACC_HOST_DATA_CLAUSES)
+ (OACC_DECLARE_CLAUSES, OACC_UPDATE_CLAUSES)
+ (OACC_ENTER_DATA_CLAUSES, OACC_EXIT_DATA_CLAUSES)
+ (OACC_WAIT_CLAUSES): New macros.
+ (gfc_free_expr_list, match_oacc_expr_list, match_oacc_clause_gang)
+ (gfc_match_omp_map_clause, gfc_match_oacc_parallel_loop)
+ (gfc_match_oacc_parallel, gfc_match_oacc_kernels_loop)
+ (gfc_match_oacc_kernels, gfc_match_oacc_data)
+ (gfc_match_oacc_host_data, gfc_match_oacc_loop)
+ (gfc_match_oacc_declare, gfc_match_oacc_update)
+ (gfc_match_oacc_enter_data, gfc_match_oacc_exit_data)
+ (gfc_match_oacc_wait, gfc_match_oacc_cache)
+ (gfc_match_oacc_routine, oacc_is_loop)
+ (resolve_oacc_scalar_int_expr, resolve_oacc_positive_int_expr)
+ (check_symbol_not_pointer, check_array_not_assumed)
+ (resolve_oacc_data_clauses, resolve_oacc_deviceptr_clause)
+ (oacc_compatible_clauses, oacc_is_parallel, oacc_is_kernels)
+ (omp_code_to_statement, oacc_code_to_statement)
+ (resolve_oacc_directive_inside_omp_region)
+ (resolve_omp_directive_inside_oacc_region)
+ (resolve_oacc_nested_loops, resolve_oacc_params_in_parallel)
+ (resolve_oacc_loop_blocks, gfc_resolve_oacc_blocks)
+ (resolve_oacc_loop, resolve_oacc_cache, gfc_resolve_oacc_declare)
+ (gfc_resolve_oacc_directive): New functions.
+ * parse.c (next_free): Update for OpenACC. Move some code into...
+ (verify_token_free): ... this new function.
+ (next_fixed): Update for OpenACC. Move some code into...
+ (verify_token_fixed): ... this new function.
+ (case_executable): Add ST_OACC_UPDATE, ST_OACC_WAIT,
+ ST_OACC_CACHE, ST_OACC_ENTER_DATA, and ST_OACC_EXIT_DATA.
+ (case_exec_markers): Add ST_OACC_PARALLEL_LOOP, ST_OACC_PARALLEL,
+ ST_OACC_KERNELS, ST_OACC_DATA, ST_OACC_HOST_DATA, ST_OACC_LOOP,
+ ST_OACC_KERNELS_LOOP.
+ (case_decl): Add ST_OACC_ROUTINE.
+ (push_state, parse_critical_block, parse_progunit): Update for
+ OpenACC.
+ (gfc_ascii_statement): Handle ST_OACC_PARALLEL_LOOP,
+ ST_OACC_END_PARALLEL_LOOP, ST_OACC_PARALLEL, ST_OACC_END_PARALLEL,
+ ST_OACC_KERNELS, ST_OACC_END_KERNELS, ST_OACC_KERNELS_LOOP,
+ ST_OACC_END_KERNELS_LOOP, ST_OACC_DATA, ST_OACC_END_DATA,
+ ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA, ST_OACC_LOOP,
+ ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT,
+ ST_OACC_CACHE, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA,
+ ST_OACC_ROUTINE.
+ (verify_st_order, parse_spec): Handle ST_OACC_DECLARE.
+ (parse_executable): Handle ST_OACC_PARALLEL_LOOP,
+ ST_OACC_KERNELS_LOOP, ST_OACC_LOOP, ST_OACC_PARALLEL,
+ ST_OACC_KERNELS, ST_OACC_DATA, ST_OACC_HOST_DATA.
+ (decode_oacc_directive, parse_oacc_structured_block)
+ (parse_oacc_loop, is_oacc): New functions.
+ * parse.h (struct gfc_state_data): Add oacc_declare_clauses
+ member.
+ (is_oacc): New prototype.
+ * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Handle
+ EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_PARALLEL,
+ EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS, EXEC_OACC_DATA,
+ EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP, EXEC_OACC_UPDATE,
+ EXEC_OACC_WAIT, EXEC_OACC_CACHE, EXEC_OACC_ENTER_DATA,
+ EXEC_OACC_EXIT_DATA.
+ (resolve_codes): Call gfc_resolve_oacc_declare.
+ * scanner.c (openacc_flag, openacc_locus): New variables.
+ (skip_free_comments): Update for OpenACC. Move some code into...
+ (skip_omp_attribute): ... this new function.
+ (skip_oacc_attribute): New function.
+ (skip_fixed_comments, gfc_next_char_literal): Update for OpenACC.
+ * st.c (gfc_free_statement): Handle EXEC_OACC_PARALLEL_LOOP,
+ EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS_LOOP, EXEC_OACC_KERNELS,
+ EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP,
+ EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
+ EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA.
+ * trans-decl.c (gfc_generate_function_code): Update for OpenACC.
+ * trans-openmp.c: Include "gomp-constants.h".
+ (gfc_omp_finish_clause, gfc_trans_omp_clauses): Use GOMP_MAP_*
+ instead of OMP_CLAUSE_MAP_*. Use OMP_CLAUSE_SET_MAP_KIND.
+ (gfc_trans_omp_clauses): Handle OMP_LIST_USE_DEVICE,
+ OMP_LIST_DEVICE_RESIDENT, OMP_LIST_CACHE, and OMP_MAP_FORCE_ALLOC,
+ OMP_MAP_FORCE_DEALLOC, OMP_MAP_FORCE_TO, OMP_MAP_FORCE_FROM,
+ OMP_MAP_FORCE_TOFROM, OMP_MAP_FORCE_PRESENT,
+ OMP_MAP_FORCE_DEVICEPTR, and gfc_omp_clauses' async, seq,
+ independent, wait_list, num_gangs_expr, num_workers_expr,
+ vector_length_expr, vector, vector_expr, worker, worker_expr,
+ gang, gang_expr members.
+ (gfc_trans_omp_do): Handle EXEC_OACC_LOOP.
+ (gfc_convert_expr_to_tree, gfc_trans_oacc_construct)
+ (gfc_trans_oacc_executable_directive)
+ (gfc_trans_oacc_wait_directive, gfc_trans_oacc_combined_directive)
+ (gfc_trans_oacc_declare, gfc_trans_oacc_directive): New functions.
+ * trans-stmt.c (gfc_trans_block_construct): Update for OpenACC.
+ * trans-stmt.h (gfc_trans_oacc_directive, gfc_trans_oacc_declare):
+ New prototypes.
+ * trans.c (tranc_code): Handle EXEC_OACC_CACHE, EXEC_OACC_WAIT,
+ EXEC_OACC_UPDATE, EXEC_OACC_LOOP, EXEC_OACC_HOST_DATA,
+ EXEC_OACC_DATA, EXEC_OACC_KERNELS, EXEC_OACC_KERNELS_LOOP,
+ EXEC_OACC_PARALLEL, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_ENTER_DATA,
+ EXEC_OACC_EXIT_DATA.
+ * gfortran.texi: Update for OpenACC.
+ * intrinsic.texi: Likewise.
+ * invoke.texi: Likewise.
+
2015-01-15 Janus Weil <janus@gcc.gnu.org>
PR fortran/58023
diff --git gcc/fortran/gfortran.texi gcc/fortran/gfortran.texi
index b9ea217..300b8b8 100644
--- gcc/fortran/gfortran.texi
+++ gcc/fortran/gfortran.texi
@@ -474,10 +474,13 @@ The GNU Fortran compiler is able to compile nearly all
standard-compliant Fortran 95, Fortran 90, and Fortran 77 programs,
including a number of standard and non-standard extensions, and can be
used on real-world programs. In particular, the supported extensions
-include OpenACC, OpenMP, Cray-style pointers, and several Fortran 2003
-and Fortran
+include OpenMP, Cray-style pointers, and several Fortran 2003 and Fortran
2008 features, including TR 15581. However, it is still under
development and has a few remaining rough edges.
+There also is initial support for OpenACC.
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC. See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
At present, the GNU Fortran compiler passes the
@uref{http://www.fortran-2000.com/ArnaudRecipes/fcvs21_f95.html,
@@ -532,9 +535,13 @@ The current status of the support is can be found in the
@ref{Fortran 2003 status}, @ref{Fortran 2008 status} and
@ref{TS 29113 status} sections of the documentation.
-Additionally, the GNU Fortran compilers supports the OpenACC specification
-(version 2.0, @url{http://www.openacc.org/}), and OpenMP specification
+Additionally, the GNU Fortran compilers supports the OpenMP specification
(version 4.0, @url{http://openmp.org/@/wp/@/openmp-specifications/}).
+There also is initial support for the OpenACC specification (targeting
+version 2.0, @uref{http://www.openacc.org/}).
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC. See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
@node Varying Length Character Strings
@subsection Varying Length Character Strings
@@ -1377,8 +1384,8 @@ without warning.
* Hollerith constants support::
* Cray pointers::
* CONVERT specifier::
-* OpenACC::
* OpenMP::
+* OpenACC::
* Argument list functions::
@end menu
@@ -1893,33 +1900,6 @@ carries a significant speed overhead. If speed in this area matters
to you, it is best if you use this only for data that needs to be
portable.
-@node OpenACC
-@subsection OpenACC
-@cindex OpenACC
-
-OpenACC is an application programming interface (API) that supports
-offloading of code to accelerator devices. It consists of a set of
-compiler directives, library routines, and environment variables that
-influence run-time behavior.
-
-GNU Fortran strives to be compatible to the
-@uref{http://www.openacc.org/, OpenACC Application Programming
-Interface v2.0}.
-
-To enable the processing of the OpenACC directive @code{!$acc} in
-free-form source code; the @code{c$acc}, @code{*$acc} and @code{!$acc}
-directives in fixed form; the @code{!$} conditional compilation
-sentinels in free form; and the @code{c$}, @code{*$} and @code{!$}
-sentinels in fixed form, @command{gfortran} needs to be invoked with
-the @option{-fopenacc}. This also arranges for automatic linking of
-the GNU Offloading and Multi Processing Runtime Library
-@ref{Top,,libgomp,libgomp,GNU Offloading and Multi Processing Runtime
-Library}.
-
-The OpenACC Fortran runtime library routines are provided both in a
-form of a Fortran 90 module named @code{openacc} and in a form of a
-Fortran @code{include} file named @file{openacc_lib.h}.
-
@node OpenMP
@subsection OpenMP
@cindex OpenMP
@@ -1980,6 +1960,37 @@ to the command line. However, this is not supported by @command{gcc} and
thus not recommended.
@end itemize
+@node OpenACC
+@subsection OpenACC
+@cindex OpenACC
+
+OpenACC is an application programming interface (API) that supports
+offloading of code to accelerator devices. It consists of a set of
+compiler directives, library routines, and environment variables that
+influence run-time behavior.
+
+GNU Fortran strives to be compatible to the
+@uref{http://www.openacc.org/, OpenACC Application Programming
+Interface v2.0}.
+
+To enable the processing of the OpenACC directive @code{!$acc} in
+free-form source code; the @code{c$acc}, @code{*$acc} and @code{!$acc}
+directives in fixed form; the @code{!$} conditional compilation
+sentinels in free form; and the @code{c$}, @code{*$} and @code{!$}
+sentinels in fixed form, @command{gfortran} needs to be invoked with
+the @option{-fopenacc}. This also arranges for automatic linking of
+the GNU Offloading and Multi Processing Runtime Library
+@ref{Top,,libgomp,libgomp,GNU Offloading and Multi Processing Runtime
+Library}.
+
+The OpenACC Fortran runtime library routines are provided both in a
+form of a Fortran 90 module named @code{openacc} and in a form of a
+Fortran @code{include} file named @file{openacc_lib.h}.
+
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC. See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
+
@node Argument list functions
@subsection Argument list functions @code{%VAL}, @code{%REF} and @code{%LOC}
@cindex argument list functions
diff --git gcc/fortran/intrinsic.texi gcc/fortran/intrinsic.texi
index 34c525b..06bce15 100644
--- gcc/fortran/intrinsic.texi
+++ gcc/fortran/intrinsic.texi
@@ -13773,8 +13773,8 @@ Fortran 95 elemental function: @ref{IEOR}
* ISO_FORTRAN_ENV::
* ISO_C_BINDING::
* IEEE modules::
-* OpenACC Module OPENACC::
* OpenMP Modules OMP_LIB and OMP_LIB_KINDS::
+* OpenACC Module OPENACC::
@end menu
@node ISO_FORTRAN_ENV
@@ -14020,33 +14020,6 @@ with the following options: @code{-fno-unsafe-math-optimizations
-@node OpenACC Module OPENACC
-@section OpenACC Module @code{OPENACC}
-@table @asis
-@item @emph{Standard}:
-OpenACC Application Programming Interface v2.0
-@end table
-
-
-The OpenACC Fortran runtime library routines are provided both in a
-form of a Fortran 90 module, named @code{OPENACC}, and in form of a
-Fortran @code{include} file named @file{openacc_lib.h}. The
-procedures provided by @code{OPENACC} can be found in the
-@ref{Top,,Introduction,libgomp,GNU Offloading and Multi Processing
-Runtime Library} manual, the named constants defined in the modules
-are listed below.
-
-For details refer to the actual
-@uref{http://www.openacc.org/,
-OpenACC Application Programming Interface v2.0}.
-
-@code{OPENACC} provides the scalar default-integer
-named constant @code{openacc_version} with a value of the form
-@var{yyyymm}, where @code{yyyy} is the year and @var{mm} the month
-of the OpenACC version; for OpenACC v2.0 the value is @code{201306}.
-
-
-
@node OpenMP Modules OMP_LIB and OMP_LIB_KINDS
@section OpenMP Modules @code{OMP_LIB} and @code{OMP_LIB_KINDS}
@table @asis
@@ -14103,3 +14076,30 @@ kind @code{omp_proc_bind_kind}:
@item @code{omp_proc_bind_close}
@item @code{omp_proc_bind_spread}
@end table
+
+
+
+@node OpenACC Module OPENACC
+@section OpenACC Module @code{OPENACC}
+@table @asis
+@item @emph{Standard}:
+OpenACC Application Programming Interface v2.0
+@end table
+
+
+The OpenACC Fortran runtime library routines are provided both in a
+form of a Fortran 90 module, named @code{OPENACC}, and in form of a
+Fortran @code{include} file named @file{openacc_lib.h}. The
+procedures provided by @code{OPENACC} can be found in the
+@ref{Top,,Introduction,libgomp,GNU Offloading and Multi Processing
+Runtime Library} manual, the named constants defined in the modules
+are listed below.
+
+For details refer to the actual
+@uref{http://www.openacc.org/,
+OpenACC Application Programming Interface v2.0}.
+
+@code{OPENACC} provides the scalar default-integer
+named constant @code{openacc_version} with a value of the form
+@var{yyyymm}, where @code{yyyy} is the year and @var{mm} the month
+of the OpenACC version; for OpenACC v2.0 the value is @code{201306}.
diff --git gcc/fortran/invoke.texi gcc/fortran/invoke.texi
index 60ad5df..9228c78 100644
--- gcc/fortran/invoke.texi
+++ gcc/fortran/invoke.texi
@@ -312,6 +312,10 @@ compilation sentinels in free form and @code{c$}, @code{*$} and
@code{!$} sentinels in fixed form, and when linking arranges for the
OpenACC runtime library to be linked in.
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC. See
+@w{@uref{https://gcc.gnu.org/wiki/OpenACC}} for more information.
+
@item -fopenmp
@opindex @code{fopenmp}
@cindex OpenMP
diff --git gcc/fortran/types.def gcc/fortran/types.def
index 40cbc94..fdae28d 100644
--- gcc/fortran/types.def
+++ gcc/fortran/types.def
@@ -218,5 +218,6 @@ DEF_FUNCTION_TYPE_VAR_8 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
BT_PTR, BT_INT, BT_INT)
DEF_FUNCTION_TYPE_VAR_12 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE, BT_PTR, BT_PTR,
- BT_PTR, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_PTR, BT_SIZE,
+ BT_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT, BT_INT,
+ BT_INT, BT_INT)
diff --git gcc/lto/ChangeLog gcc/lto/ChangeLog
index 4eafacc..9b98edd 100644
--- gcc/lto/ChangeLog
+++ gcc/lto/ChangeLog
@@ -1,3 +1,10 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * lto-lang.c (DEF_FUNCTION_TYPE_VAR_8, DEF_FUNCTION_TYPE_VAR_12):
+ New macros.
+ * lto.c: Include "gomp-constants.h".
+
2015-01-14 Ilya Verbin <ilya.verbin@intel.com>
* lto-partition.c (lto_promote_cross_file_statics): Remove argument
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index d9525d5..885a710 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,133 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ Ilmir Usmanov <i.usmanov@samsung.com>
+
+ * lib/target-supports.exp (check_effective_target_fopenacc): New
+ procedure.
+ * g++.dg/goacc-gomp/goacc-gomp.exp: New file.
+ * g++.dg/goacc/goacc.exp: Likewise.
+ * gcc.dg/goacc-gomp/goacc-gomp.exp: Likewise.
+ * gcc.dg/goacc/goacc.exp: Likewise.
+ * gfortran.dg/goacc/goacc.exp: Likewise.
+ * c-c++-common/cpp/openacc-define-1.c: New file.
+ * c-c++-common/cpp/openacc-define-2.c: Likewise.
+ * c-c++-common/cpp/openacc-define-3.c: Likewise.
+ * c-c++-common/goacc-gomp/nesting-1.c: Likewise.
+ * c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
+ * c-c++-common/goacc/acc_on_device-2-off.c: Likewise.
+ * c-c++-common/goacc/acc_on_device-2.c: Likewise.
+ * c-c++-common/goacc/asyncwait-1.c: Likewise.
+ * c-c++-common/goacc/cache-1.c: Likewise.
+ * c-c++-common/goacc/clauses-fail.c: Likewise.
+ * c-c++-common/goacc/collapse-1.c: Likewise.
+ * c-c++-common/goacc/data-1.c: Likewise.
+ * c-c++-common/goacc/data-2.c: Likewise.
+ * c-c++-common/goacc/data-clause-duplicate-1.c: Likewise.
+ * c-c++-common/goacc/deviceptr-1.c: Likewise.
+ * c-c++-common/goacc/deviceptr-2.c: Likewise.
+ * c-c++-common/goacc/deviceptr-3.c: Likewise.
+ * c-c++-common/goacc/if-clause-1.c: Likewise.
+ * c-c++-common/goacc/if-clause-2.c: Likewise.
+ * c-c++-common/goacc/kernels-1.c: Likewise.
+ * c-c++-common/goacc/loop-1.c: Likewise.
+ * c-c++-common/goacc/loop-private-1.c: Likewise.
+ * c-c++-common/goacc/nesting-1.c: Likewise.
+ * c-c++-common/goacc/nesting-data-1.c: Likewise.
+ * c-c++-common/goacc/nesting-fail-1.c: Likewise.
+ * c-c++-common/goacc/parallel-1.c: Likewise.
+ * c-c++-common/goacc/pcopy.c: Likewise.
+ * c-c++-common/goacc/pcopyin.c: Likewise.
+ * c-c++-common/goacc/pcopyout.c: Likewise.
+ * c-c++-common/goacc/pcreate.c: Likewise.
+ * c-c++-common/goacc/pragma_context.c: Likewise.
+ * c-c++-common/goacc/present-1.c: Likewise.
+ * c-c++-common/goacc/reduction-1.c: Likewise.
+ * c-c++-common/goacc/reduction-2.c: Likewise.
+ * c-c++-common/goacc/reduction-3.c: Likewise.
+ * c-c++-common/goacc/reduction-4.c: Likewise.
+ * c-c++-common/goacc/sb-1.c: Likewise.
+ * c-c++-common/goacc/sb-2.c: Likewise.
+ * c-c++-common/goacc/sb-3.c: Likewise.
+ * c-c++-common/goacc/update-1.c: Likewise.
+ * gcc.dg/goacc/acc_on_device-1.c: Likewise.
+ * gfortran.dg/goacc/acc_on_device-1.f95: Likewise.
+ * gfortran.dg/goacc/acc_on_device-2-off.f95: Likewise.
+ * gfortran.dg/goacc/acc_on_device-2.f95: Likewise.
+ * gfortran.dg/goacc/assumed.f95: Likewise.
+ * gfortran.dg/goacc/asyncwait-1.f95: Likewise.
+ * gfortran.dg/goacc/asyncwait-2.f95: Likewise.
+ * gfortran.dg/goacc/asyncwait-3.f95: Likewise.
+ * gfortran.dg/goacc/asyncwait-4.f95: Likewise.
+ * gfortran.dg/goacc/branch.f95: Likewise.
+ * gfortran.dg/goacc/cache-1.f95: Likewise.
+ * gfortran.dg/goacc/coarray.f95: Likewise.
+ * gfortran.dg/goacc/continuation-free-form.f95: Likewise.
+ * gfortran.dg/goacc/cray.f95: Likewise.
+ * gfortran.dg/goacc/critical.f95: Likewise.
+ * gfortran.dg/goacc/data-clauses.f95: Likewise.
+ * gfortran.dg/goacc/data-tree.f95: Likewise.
+ * gfortran.dg/goacc/declare-1.f95: Likewise.
+ * gfortran.dg/goacc/enter-exit-data.f95: Likewise.
+ * gfortran.dg/goacc/fixed-1.f: Likewise.
+ * gfortran.dg/goacc/fixed-2.f: Likewise.
+ * gfortran.dg/goacc/fixed-3.f: Likewise.
+ * gfortran.dg/goacc/fixed-4.f: Likewise.
+ * gfortran.dg/goacc/host_data-tree.f95: Likewise.
+ * gfortran.dg/goacc/if.f95: Likewise.
+ * gfortran.dg/goacc/kernels-tree.f95: Likewise.
+ * gfortran.dg/goacc/list.f95: Likewise.
+ * gfortran.dg/goacc/literal.f95: Likewise.
+ * gfortran.dg/goacc/loop-1.f95: Likewise.
+ * gfortran.dg/goacc/loop-2.f95: Likewise.
+ * gfortran.dg/goacc/loop-3.f95: Likewise.
+ * gfortran.dg/goacc/loop-tree-1.f90: Likewise.
+ * gfortran.dg/goacc/omp.f95: Likewise.
+ * gfortran.dg/goacc/parallel-kernels-clauses.f95: Likewise.
+ * gfortran.dg/goacc/parallel-kernels-regions.f95: Likewise.
+ * gfortran.dg/goacc/parallel-tree.f95: Likewise.
+ * gfortran.dg/goacc/parameter.f95: Likewise.
+ * gfortran.dg/goacc/private-1.f95: Likewise.
+ * gfortran.dg/goacc/private-2.f95: Likewise.
+ * gfortran.dg/goacc/private-3.f95: Likewise.
+ * gfortran.dg/goacc/pure-elemental-procedures.f95: Likewise.
+ * gfortran.dg/goacc/reduction-2.f95: Likewise.
+ * gfortran.dg/goacc/reduction.f95: Likewise.
+ * gfortran.dg/goacc/routine-1.f90: Likewise.
+ * gfortran.dg/goacc/routine-2.f90: Likewise.
+ * gfortran.dg/goacc/sentinel-free-form.f95: Likewise.
+ * gfortran.dg/goacc/several-directives.f95: Likewise.
+ * gfortran.dg/goacc/sie.f95: Likewise.
+ * gfortran.dg/goacc/subarrays.f95: Likewise.
+ * gfortran.dg/gomp/map-1.f90: Likewise.
+ * gfortran.dg/openacc-define-1.f90: Likewise.
+ * gfortran.dg/openacc-define-2.f90: Likewise.
+ * gfortran.dg/openacc-define-3.f90: Likewise.
+ * g++.dg/gomp/block-1.C: Update for changed compiler output.
+ * g++.dg/gomp/block-2.C: Likewise.
+ * g++.dg/gomp/block-3.C: Likewise.
+ * g++.dg/gomp/block-5.C: Likewise.
+ * g++.dg/gomp/target-1.C: Likewise.
+ * g++.dg/gomp/target-2.C: Likewise.
+ * g++.dg/gomp/taskgroup-1.C: Likewise.
+ * g++.dg/gomp/teams-1.C: Likewise.
+ * gcc.dg/cilk-plus/jump-openmp.c: Likewise.
+ * gcc.dg/cilk-plus/jump.c: Likewise.
+ * gcc.dg/gomp/block-1.c: Likewise.
+ * gcc.dg/gomp/block-10.c: Likewise.
+ * gcc.dg/gomp/block-2.c: Likewise.
+ * gcc.dg/gomp/block-3.c: Likewise.
+ * gcc.dg/gomp/block-4.c: Likewise.
+ * gcc.dg/gomp/block-5.c: Likewise.
+ * gcc.dg/gomp/block-6.c: Likewise.
+ * gcc.dg/gomp/block-7.c: Likewise.
+ * gcc.dg/gomp/block-8.c: Likewise.
+ * gcc.dg/gomp/block-9.c: Likewise.
+ * gcc.dg/gomp/target-1.c: Likewise.
+ * gcc.dg/gomp/target-2.c: Likewise.
+ * gcc.dg/gomp/taskgroup-1.c: Likewise.
+ * gcc.dg/gomp/teams-1.c: Likewise.
+
2015-01-15 David Malcolm <dmalcolm@redhat.com>
* jit.dg/test-error-mismatching-types-in-assignment-op.c: New
diff --git gcc/testsuite/c-c++-common/goacc/asyncwait-1.c gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
index 9d69dcf..ccc0106 100644
--- gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
+++ gcc/testsuite/c-c++-common/goacc/asyncwait-1.c
@@ -1,224 +1,159 @@
-/* { dg-skip-if "not yet" { c++ } } */
-
-#include <stdlib.h>
-
-int
-main (int argc, char **argv)
+void
+f (int N, float *a, float *b)
{
- int N = 64;
- float *a, *b;
- int i;
-
- a = (float *) malloc (N * sizeof (float));
- b = (float *) malloc (N * sizeof (float));
-
- for (i = 0; i < N; i++)
- {
- a[i] = 3.0;
- b[i] = 0.0;
- }
+ int ii;
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (,1) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,,) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 /* { dg-error "expected '\\)' before end of line" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (*) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (a)
- /* { dg-error "expected integer expression before" "" { target c } 85 } */
- /* { dg-error "'async' expression must be integral" "" { target c++ } 85 } */
+ /* { dg-error "expected integer expression before" "" { target c } 54 } */
+ /* { dg-error "'async' expression must be integral" "" { target c++ } 54 } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1.0)
- /* { dg-error "expected integer expression before" "" { target c } 95 } */
- /* { dg-error "'async' expression must be integral" "" { target c++ } 95 } */
+ /* { dg-error "expected integer expression before" "" { target c } 62 } */
+ /* { dg-error "'async' expression must be integral" "" { target c++ } 62 } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async () /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (,1) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,,) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 /* { dg-error "expected '\\\)' before end of line" } */
+ /* { dg-error "expected integer expression before '\\\)'" "" { target c++ } 118 } */
{
- /* { dg-error "expected integer expression list before" "" { target c++ } 169 } */
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,*) /* { dg-error "expected (primary-|)expression before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,a) /*{ dg-error "must be integral" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (a) /* { dg-error "must be integral" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1.0) /* { dg-error "must be integral" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait () /* { dg-error "expected (integer |)expression (list |)before" } */
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait
{
- int ii;
-
for (ii = 0; ii < N; ii++)
b[ii] = a[ii];
}
@@ -236,7 +171,7 @@ main (int argc, char **argv)
#pragma acc wait (1,2,,) /* { dg-error "expected (primary-|)expression before" } */
#pragma acc wait (1 /* { dg-error "expected '\\\)' before end of line" } */
- /* { dg-error "expected integer expression list before" "" { target c++ } 238 } */
+ /* { dg-error "expected integer expression before '\\\)'" "" { target c++ } 173 } */
#pragma acc wait (1,*) /* { dg-error "expected (primary-|)expression before" } */
@@ -246,9 +181,9 @@ main (int argc, char **argv)
#pragma acc wait (1.0) /* { dg-error "expression must be integral" } */
-#pragma acc wait 1 /* { dg-error "expected clause before numeric constant" } */
+#pragma acc wait 1 /* { dg-error "expected '#pragma acc' clause before numeric constant" } */
-#pragma acc wait N /* { dg-error "expected clause before 'N'" } */
+#pragma acc wait N /* { dg-error "expected '#pragma acc' clause before 'N'" } */
#pragma acc wait async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */
@@ -269,12 +204,10 @@ main (int argc, char **argv)
#pragma acc wait async (*) /* { dg-error "expected (primary-|)expression before " } */
#pragma acc wait async (a)
- /* { dg-error "expected integer expression before" "" { target c } 271 } */
- /* { dg-error "expression must be integral" "" { target c++ } 271 } */
+ /* { dg-error "expected integer expression before" "" { target c } 206 } */
+ /* { dg-error "expression must be integral" "" { target c++ } 206 } */
#pragma acc wait async (1.0)
- /* { dg-error "expected integer expression before" "" { target c } 275 } */
- /* { dg-error "expression must be integral" "" { target c++ } 275 } */
-
- return 0;
+ /* { dg-error "expected integer expression before" "" { target c } 210 } */
+ /* { dg-error "expression must be integral" "" { target c++ } 210 } */
}
diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c
index e8e1278..8990180 100644
--- gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -3,16 +3,16 @@ f (void)
{
int i;
-#pragma acc parallel one /* { dg-error "expected clause before 'one'" } */
+#pragma acc parallel one /* { dg-error "expected '#pragma acc' clause before 'one'" } */
;
-#pragma acc kernels eins /* { dg-error "expected clause before 'eins'" } */
+#pragma acc kernels eins /* { dg-error "expected '#pragma acc' clause before 'eins'" } */
;
-#pragma acc data two /* { dg-error "expected clause before 'two'" } */
+#pragma acc data two /* { dg-error "expected '#pragma acc' clause before 'two'" } */
;
-#pragma acc loop deux /* { dg-error "expected clause before 'deux'" } */
+#pragma acc loop deux /* { dg-error "expected '#pragma acc' clause before 'deux'" } */
for (i = 0; i < 2; ++i)
;
}
diff --git gcc/testsuite/c-c++-common/goacc/data-2.c gcc/testsuite/c-c++-common/goacc/data-2.c
index 9c0a185..a67d8a4 100644
--- gcc/testsuite/c-c++-common/goacc/data-2.c
+++ gcc/testsuite/c-c++-common/goacc/data-2.c
@@ -5,7 +5,7 @@ foo (void)
int n;
#pragma acc enter data copyin (a, b) async wait
#pragma acc enter data create (b[20:30]) async wait
-#pragma acc enter data (a) /* { dg-error "expected clause before '\\\(' token" } */
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
#pragma acc exit data delete (a) if (0)
#pragma acc exit data copyout (b) if (a)
diff --git include/ChangeLog include/ChangeLog
index c1011b9..0917d94 100644
--- include/ChangeLog
+++ include/ChangeLog
@@ -1,3 +1,9 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * gomp-constants.h: New file.
+
2015-12-14 Jan-Benedict Glaw <jbglaw@lug-owl.de>
* libiberty.h: Merge Copyright year update from Binutils.
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 6e1e141..9b003cb 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,354 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Tom de Vries <tom@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ Nathan Sidwell <nathan@codesourcery.com>
+ Tobias Burnus <burnus@net-b.de>
+
+ * Makefile.am (search_path): Add $(top_srcdir)/../include.
+ (libgomp_la_SOURCES): Add splay-tree.c, libgomp-plugin.c,
+ oacc-parallel.c, oacc-host.c, oacc-init.c, oacc-mem.c,
+ oacc-async.c, oacc-plugin.c, oacc-cuda.c.
+ [USE_FORTRAN] (libgomp_la_SOURCES): Add openacc.f90.
+ Include $(top_srcdir)/plugin/Makefrag.am.
+ (nodist_libsubinclude_HEADERS): Add openacc.h.
+ [USE_FORTRAN] (nodist_finclude_HEADERS): Add openacc_lib.h,
+ openacc.f90, openacc.mod, openacc_kinds.mod.
+ (omp_lib.mod): Generalize into...
+ (%.mod): ... this new rule.
+ (openacc_kinds.mod, openacc.mod): New rules.
+ * plugin/configfrag.ac: New file.
+ * configure.ac: Move plugin/offloading support into it. Include
+ it. Instantiate testsuite/libgomp-test-support.pt.exp.
+ * plugin/Makefrag.am: New file.
+ * testsuite/Makefile.am (OFFLOAD_TARGETS)
+ (OFFLOAD_ADDITIONAL_OPTIONS, OFFLOAD_ADDITIONAL_LIB_PATHS): Don't
+ export.
+ (libgomp-test-support.exp): New rule.
+ (all-local): Depend on it.
+ * Makefile.in: Regenerate.
+ * testsuite/Makefile.in: Regenerate.
+ * config.h.in: Likewise.
+ * configure: Likewise.
+ * configure.tgt: Harden shell syntax.
+ * env.c: Include "oacc-int.h".
+ (parse_acc_device_type): New function.
+ (gomp_debug_var, goacc_device_type, goacc_device_num): New
+ variables.
+ (initialize_env): Initialize those. Call
+ goacc_runtime_initialize.
+ * error.c (gomp_vdebug, gomp_debug, gomp_vfatal): New functions.
+ (gomp_fatal): Call gomp_vfatal.
+ * libgomp.h: Include "libgomp-plugin.h" and <stdarg.h>.
+ (gomp_debug_var, goacc_device_type, goacc_device_num, gomp_vdebug)
+ (gomp_debug, gomp_verror, gomp_vfatal, gomp_init_targets_once)
+ (splay_tree_node, splay_tree, splay_tree_key)
+ (struct target_mem_desc, struct splay_tree_key_s)
+ (struct gomp_memory_mapping, struct acc_dispatch_t)
+ (struct gomp_device_descr, gomp_acc_insert_pointer)
+ (gomp_acc_remove_pointer, target_mem_desc, gomp_copy_from_async)
+ (gomp_unmap_vars, gomp_init_device, gomp_init_tables)
+ (gomp_free_memmap, gomp_fini_device): New declarations.
+ (gomp_vdebug, gomp_debug): New macros.
+ Include "splay-tree.h".
+ * libgomp.map (OACC_2.0): New symbol version. Use for
+ acc_get_num_devices, acc_get_num_devices_h_, acc_set_device_type,
+ acc_set_device_type_h_, acc_get_device_type,
+ acc_get_device_type_h_, acc_set_device_num, acc_set_device_num_h_,
+ acc_get_device_num, acc_get_device_num_h_, acc_async_test,
+ acc_async_test_h_, acc_async_test_all, acc_async_test_all_h_,
+ acc_wait, acc_wait_h_, acc_wait_async, acc_wait_async_h_,
+ acc_wait_all, acc_wait_all_h_, acc_wait_all_async,
+ acc_wait_all_async_h_, acc_init, acc_init_h_, acc_shutdown,
+ acc_shutdown_h_, acc_on_device, acc_on_device_h_, acc_malloc,
+ acc_free, acc_copyin, acc_copyin_32_h_, acc_copyin_64_h_,
+ acc_copyin_array_h_, acc_present_or_copyin,
+ acc_present_or_copyin_32_h_, acc_present_or_copyin_64_h_,
+ acc_present_or_copyin_array_h_, acc_create, acc_create_32_h_,
+ acc_create_64_h_, acc_create_array_h_, acc_present_or_create,
+ acc_present_or_create_32_h_, acc_present_or_create_64_h_,
+ acc_present_or_create_array_h_, acc_copyout, acc_copyout_32_h_,
+ acc_copyout_64_h_, acc_copyout_array_h_, acc_delete,
+ acc_delete_32_h_, acc_delete_64_h_, acc_delete_array_h_,
+ acc_update_device, acc_update_device_32_h_,
+ acc_update_device_64_h_, acc_update_device_array_h_,
+ acc_update_self, acc_update_self_32_h_, acc_update_self_64_h_,
+ acc_update_self_array_h_, acc_map_data, acc_unmap_data,
+ acc_deviceptr, acc_hostptr, acc_is_present, acc_is_present_32_h_,
+ acc_is_present_64_h_, acc_is_present_array_h_,
+ acc_memcpy_to_device, acc_memcpy_from_device,
+ acc_get_current_cuda_device, acc_get_current_cuda_context,
+ acc_get_cuda_stream, acc_set_cuda_stream.
+ (GOACC_2.0): New symbol version. Use for GOACC_data_end,
+ GOACC_data_start, GOACC_enter_exit_data, GOACC_parallel,
+ GOACC_update, GOACC_wait, GOACC_get_thread_num,
+ GOACC_get_num_threads.
+ (GOMP_PLUGIN_1.0): New symbol version. Use for
+ GOMP_PLUGIN_malloc, GOMP_PLUGIN_malloc_cleared,
+ GOMP_PLUGIN_realloc, GOMP_PLUGIN_debug, GOMP_PLUGIN_error,
+ GOMP_PLUGIN_fatal, GOMP_PLUGIN_async_unmap_vars,
+ GOMP_PLUGIN_acc_thread.
+ * libgomp.texi: Update for OpenACC changes, and GOMP_DEBUG
+ environment variable.
+ * libgomp_g.h (GOACC_data_start, GOACC_data_end)
+ (GOACC_enter_exit_data, GOACC_parallel, GOACC_update, GOACC_wait)
+ (GOACC_get_num_threads, GOACC_get_thread_num): New declarations.
+ * splay-tree.h (splay_tree_lookup, splay_tree_insert)
+ (splay_tree_remove): New declarations.
+ (rotate_left, rotate_right, splay_tree_splay, splay_tree_insert)
+ (splay_tree_remove, splay_tree_lookup): Move into...
+ * splay-tree.c: ... this new file.
+ * target.c: Include "oacc-plugin.h", "oacc-int.h", <assert.h>.
+ (splay_tree_node, splay_tree, splay_tree_key)
+ (struct target_mem_desc, struct splay_tree_key_s)
+ (struct gomp_device_descr): Don't declare.
+ (num_devices_openmp): New variable.
+ (gomp_get_num_devices ): Use it.
+ (gomp_init_targets_once): New function.
+ (gomp_get_num_devices ): Use it.
+ (get_kind, gomp_copy_from_async, gomp_free_memmap)
+ (gomp_fini_device, gomp_register_image_for_device): New functions.
+ (gomp_map_vars): Add devaddrs parameter.
+ (gomp_update): Add mm parameter.
+ (gomp_init_device): Move most of it into...
+ (gomp_init_tables): ... this new function.
+ (gomp_register_images_for_device): Remove function.
+ (splay_compare, gomp_map_vars, gomp_unmap_vars, gomp_init_device):
+ Make them hidden instead of static.
+ (gomp_map_vars_existing, gomp_map_vars, gomp_unmap_vars)
+ (gomp_update, gomp_init_device, GOMP_target, GOMP_target_data)
+ (GOMP_target_end_data, GOMP_target_update)
+ (gomp_load_plugin_for_device, gomp_target_init): Update for
+ OpenACC changes.
+ * oacc-async.c: New file.
+ * oacc-cuda.c: Likewise.
+ * oacc-host.c: Likewise.
+ * oacc-init.c: Likewise.
+ * oacc-int.h: Likewise.
+ * oacc-mem.c: Likewise.
+ * oacc-parallel.c: Likewise.
+ * oacc-plugin.c: Likewise.
+ * oacc-plugin.h: Likewise.
+ * oacc-ptx.h: Likewise.
+ * openacc.f90: Likewise.
+ * openacc.h: Likewise.
+ * openacc_lib.h: Likewise.
+ * plugin/plugin-host.c: Likewise.
+ * plugin/plugin-nvptx.c: Likewise.
+ * libgomp-plugin.c: Likewise.
+ * libgomp-plugin.h: Likewise.
+ * libgomp_target.h: Remove file after merging content into the
+ former file. Update all users.
+ * testsuite/lib/libgomp.exp: Load libgomp-test-support.exp.
+ (offload_targets_s, offload_targets_s_openacc): New variables.
+ (check_effective_target_openacc_nvidia_accel_present)
+ (check_effective_target_openacc_nvidia_accel_selected): New
+ procedures.
+ (libgomp_init): Update for OpenACC changes.
+ * testsuite/libgomp-test-support.exp.in: New file.
+ * testsuite/libgomp.oacc-c++/c++.exp: Likewise.
+ * testsuite/libgomp.oacc-c/c.exp: Likewise.
+ * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/abort-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/abort-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/abort-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/abort-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/cache-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/clauses-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/collapse-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/kernels-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/kernels-empty.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-10.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-11.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-12.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-19.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-26.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-27.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-31.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-32.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-33.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-35.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-36.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-37.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-39.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-40.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-41.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-45.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-46.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-49.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-50.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-51.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-55.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-56.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-57.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-58.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-59.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-6.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-60.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-61.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-62.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-63.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-64.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-65.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-66.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-67.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-68.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-69.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-7.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-71.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-72.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-74.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-75.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-76.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-77.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-78.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-80.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-81.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-83.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-84.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-86.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-87.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-88.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-89.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-9.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-90.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-92.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/nested-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/offset-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/parallel-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/parallel-empty.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/pointer-align-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/present-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/subr.h: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/subr.ptx: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/timer.h: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/update-1.c: Likewise.
+ * testsuite/libgomp.oacc-fortran/abort-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/abort-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/asyncwait-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/asyncwait-3.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-3.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-4.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-5.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-6.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-7.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/collapse-8.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-3.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-4.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-10.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-3.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-4.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-5.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-6.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-7.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/lib-8.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/map-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/openacc_version-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/openacc_version-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/pointer-align-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/pset-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-3.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-4.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/routine-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/routine-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/routine-3.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/routine-4.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/subarrays-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/subarrays-2.f90: Likewise.
+
2015-01-10 Thomas Schwinge <thomas@codesourcery.com>
Julian Brown <julian@codesourcery.com>
David Malcolm <dmalcolm@redhat.com>
diff --git libgomp/Makefile.am libgomp/Makefile.am
index 4471fab..5411278 100644
--- libgomp/Makefile.am
+++ libgomp/Makefile.am
@@ -61,9 +61,9 @@ libgomp_la_LINK = $(LINK) $(libgomp_la_LDFLAGS)
libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c single.c \
task.c team.c work.c lock.c mutex.c proc.c sem.c bar.c ptrlock.c \
- time.c fortran.c affinity.c target.c oacc-parallel.c splay-tree.c \
- oacc-host.c oacc-init.c oacc-mem.c oacc-async.c oacc-plugin.c \
- oacc-cuda.c libgomp-plugin.c
+ time.c fortran.c affinity.c target.c splay-tree.c libgomp-plugin.c \
+ oacc-parallel.c oacc-host.c oacc-init.c oacc-mem.c oacc-async.c \
+ oacc-plugin.c oacc-cuda.c
include $(top_srcdir)/plugin/Makefrag.am
diff --git libgomp/Makefile.in libgomp/Makefile.in
index 8758989..b61b108 100644
--- libgomp/Makefile.in
+++ libgomp/Makefile.in
@@ -151,9 +151,9 @@ am_libgomp_la_OBJECTS = alloc.lo barrier.lo critical.lo env.lo \
error.lo iter.lo iter_ull.lo loop.lo loop_ull.lo ordered.lo \
parallel.lo sections.lo single.lo task.lo team.lo work.lo \
lock.lo mutex.lo proc.lo sem.lo bar.lo ptrlock.lo time.lo \
- fortran.lo affinity.lo target.lo oacc-parallel.lo \
- splay-tree.lo oacc-host.lo oacc-init.lo oacc-mem.lo \
- oacc-async.lo oacc-plugin.lo oacc-cuda.lo libgomp-plugin.lo \
+ fortran.lo affinity.lo target.lo splay-tree.lo \
+ libgomp-plugin.lo oacc-parallel.lo oacc-host.lo oacc-init.lo \
+ oacc-mem.lo oacc-async.lo oacc-plugin.lo oacc-cuda.lo \
$(am__objects_1)
libgomp_la_OBJECTS = $(am_libgomp_la_OBJECTS)
DEFAULT_INCLUDES = -I.@am__isrc@
@@ -396,9 +396,9 @@ libgomp_la_SOURCES = alloc.c barrier.c critical.c env.c error.c iter.c \
iter_ull.c loop.c loop_ull.c ordered.c parallel.c sections.c \
single.c task.c team.c work.c lock.c mutex.c proc.c sem.c \
bar.c ptrlock.c time.c fortran.c affinity.c target.c \
- oacc-parallel.c splay-tree.c oacc-host.c oacc-init.c \
- oacc-mem.c oacc-async.c oacc-plugin.c oacc-cuda.c \
- libgomp-plugin.c $(am__append_2)
+ splay-tree.c libgomp-plugin.c oacc-parallel.c oacc-host.c \
+ oacc-init.c oacc-mem.c oacc-async.c oacc-plugin.c oacc-cuda.c \
+ $(am__append_2)
# Nvidia PTX OpenACC plugin.
@PLUGIN_NVPTX_TRUE@libgomp_plugin_nvptx_version_info = -version-info $(libtool_VERSION)
diff --git libgomp/env.c libgomp/env.c
index 130c52c..6b5e963 100644
--- libgomp/env.c
+++ libgomp/env.c
@@ -79,8 +79,8 @@ unsigned long gomp_bind_var_list_len;
void **gomp_places_list;
unsigned long gomp_places_list_len;
int gomp_debug_var;
+char *goacc_device_type;
int goacc_device_num;
-char* goacc_device_type;
/* Parse the OMP_SCHEDULE environment variable. */
@@ -1017,10 +1017,10 @@ parse_affinity (bool ignore)
}
static void
-goacc_parse_device_type (void)
+parse_acc_device_type (void)
{
const char *env = getenv ("ACC_DEVICE_TYPE");
-
+
if (env && *env != '\0')
goacc_device_type = strdup (env);
else
@@ -1287,14 +1287,14 @@ initialize_env (void)
}
handle_omp_display_env (stacksize, wait_policy);
-
- /* Look for OpenACC-specific environment variables. */
+
+ /* OpenACC. */
+
if (!parse_int ("ACC_DEVICE_NUM", &goacc_device_num, true))
goacc_device_num = 0;
- goacc_parse_device_type ();
+ parse_acc_device_type ();
- /* Initialize OpenACC-specific internal state. */
goacc_runtime_initialize ();
}
diff --git libgomp/error.c libgomp/error.c
index fb96fb3..094c24a 100644
--- libgomp/error.c
+++ libgomp/error.c
@@ -38,7 +38,7 @@
#undef gomp_vdebug
void
-gomp_vdebug (int kind __attribute__((unused)), const char *msg, va_list list)
+gomp_vdebug (int kind __attribute__ ((unused)), const char *msg, va_list list)
{
if (gomp_debug_var)
vfprintf (stderr, msg, list);
diff --git libgomp/libgomp-plugin.c libgomp/libgomp-plugin.c
index 1dd33f5..ffb22e9 100644
--- libgomp/libgomp-plugin.c
+++ libgomp/libgomp-plugin.c
@@ -53,7 +53,7 @@ void
GOMP_PLUGIN_debug (int kind, const char *msg, ...)
{
va_list ap;
-
+
va_start (ap, msg);
gomp_debug (kind, msg, ap);
va_end (ap);
@@ -63,7 +63,7 @@ void
GOMP_PLUGIN_error (const char *msg, ...)
{
va_list ap;
-
+
va_start (ap, msg);
gomp_verror (msg, ap);
va_end (ap);
@@ -73,11 +73,8 @@ void
GOMP_PLUGIN_fatal (const char *msg, ...)
{
va_list ap;
-
+
va_start (ap, msg);
gomp_vfatal (msg, ap);
va_end (ap);
-
- /* Unreachable. */
- abort ();
}
diff --git libgomp/libgomp-plugin.h libgomp/libgomp-plugin.h
index c8383e1..d9cbff5 100644
--- libgomp/libgomp-plugin.h
+++ libgomp/libgomp-plugin.h
@@ -62,16 +62,16 @@ struct mapping_table
};
/* Miscellaneous functions. */
-extern void *GOMP_PLUGIN_malloc (size_t) __attribute__((malloc));
-extern void *GOMP_PLUGIN_malloc_cleared (size_t) __attribute__((malloc));
+extern void *GOMP_PLUGIN_malloc (size_t) __attribute__ ((malloc));
+extern void *GOMP_PLUGIN_malloc_cleared (size_t) __attribute__ ((malloc));
extern void *GOMP_PLUGIN_realloc (void *, size_t);
extern void GOMP_PLUGIN_debug (int, const char *, ...)
- __attribute__((format (printf, 2, 3)));
+ __attribute__ ((format (printf, 2, 3)));
extern void GOMP_PLUGIN_error (const char *, ...)
- __attribute__((format (printf, 1, 2)));
+ __attribute__ ((format (printf, 1, 2)));
extern void GOMP_PLUGIN_fatal (const char *, ...)
- __attribute__((noreturn, format (printf, 1, 2)));
+ __attribute__ ((noreturn, format (printf, 1, 2)));
#ifdef __cplusplus
}
diff --git libgomp/libgomp.h libgomp/libgomp.h
index 97732a5..ab8e229 100644
--- libgomp/libgomp.h
+++ libgomp/libgomp.h
@@ -260,7 +260,7 @@ extern void **gomp_places_list;
extern unsigned long gomp_places_list_len;
extern int gomp_debug_var;
extern int goacc_device_num;
-extern char* goacc_device_type;
+extern char *goacc_device_type;
enum gomp_task_kind
{
@@ -542,7 +542,7 @@ extern void *gomp_realloc (void *, size_t);
extern void gomp_vdebug (int, const char *, va_list);
extern void gomp_debug (int, const char *, ...)
- __attribute__((format (printf, 2, 3)));
+ __attribute__ ((format (printf, 2, 3)));
#define gomp_vdebug(KIND, FMT, VALIST) \
do { \
if (__builtin_expect (gomp_debug_var, 0)) \
@@ -555,11 +555,11 @@ extern void gomp_debug (int, const char *, ...)
} while (0)
extern void gomp_verror (const char *, va_list);
extern void gomp_error (const char *, ...)
- __attribute__((format (printf, 1, 2)));
+ __attribute__ ((format (printf, 1, 2)));
extern void gomp_vfatal (const char *, va_list)
- __attribute__((noreturn));
+ __attribute__ ((noreturn));
extern void gomp_fatal (const char *, ...)
- __attribute__((noreturn, format (printf, 1, 2)));
+ __attribute__ ((noreturn, format (printf, 1, 2)));
/* iter.c */
@@ -633,7 +633,9 @@ extern void gomp_free_thread (void *);
extern void gomp_init_targets_once (void);
extern int gomp_get_num_devices (void);
-#include "splay-tree.h"
+typedef struct splay_tree_node_s *splay_tree_node;
+typedef struct splay_tree_s *splay_tree;
+typedef struct splay_tree_key_s *splay_tree_key;
struct target_mem_desc {
/* Reference count. */
@@ -662,18 +664,37 @@ struct target_mem_desc {
splay_tree_key list[];
};
+struct splay_tree_key_s {
+ /* Address of the host object. */
+ uintptr_t host_start;
+ /* Address immediately after the host object. */
+ uintptr_t host_end;
+ /* Descriptor of the target memory. */
+ struct target_mem_desc *tgt;
+ /* Offset from tgt->tgt_start to the start of the target object. */
+ uintptr_t tgt_offset;
+ /* Reference count. */
+ uintptr_t refcount;
+ /* Asynchronous reference count. */
+ uintptr_t async_refcount;
+ /* True if data should be copied from device to host at the end. */
+ bool copy_from;
+};
+
+#include "splay-tree.h"
+
/* Information about mapped memory regions (per device/context). */
struct gomp_memory_mapping
{
- /* Splay tree containing information about mapped memory regions. */
- struct splay_tree_s splay_tree;
-
/* Mutex for operating with the splay tree and other shared structures. */
gomp_mutex_t lock;
/* True when tables have been added to this memory map. */
bool is_initialized;
+
+ /* Splay tree containing information about mapped memory regions. */
+ struct splay_tree_s splay_tree;
};
typedef struct acc_dispatch_t
@@ -747,12 +768,6 @@ struct gomp_device_descr
/* This is the TYPE of device. */
enum offload_target_type type;
- /* Set to true when device is initialized. */
- bool is_initialized;
-
- /* True when offload regions have been registered with this device. */
- bool offload_regions_registered;
-
/* Function handlers. */
const char *(*get_name_func) (void);
unsigned int (*get_caps_func) (void);
@@ -768,17 +783,23 @@ struct gomp_device_descr
void *(*host2dev_func) (int, void *, const void *, size_t);
void (*run_func) (int, void *, void *);
- /* OpenACC-specific functions. */
- /* This is mutable because of its mutable data_environ and target_data
- members. */
- acc_dispatch_t openacc;
-
/* Memory-mapping info for this device instance. */
/* Uses a separate lock. */
struct gomp_memory_mapping mem_map;
/* Mutex for the mutable data. */
gomp_mutex_t lock;
+
+ /* Set to true when device is initialized. */
+ bool is_initialized;
+
+ /* True when offload regions have been registered with this device. */
+ bool offload_regions_registered;
+
+ /* OpenACC-specific data and functions. */
+ /* This is mutable because of its mutable data_environ and target_data
+ members. */
+ acc_dispatch_t openacc;
};
extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *);
@@ -787,18 +808,12 @@ extern void gomp_acc_remove_pointer (void *, bool, int, int);
extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
size_t, void **, void **,
size_t *, void *, bool, bool);
-
extern void gomp_copy_from_async (struct target_mem_desc *);
-
extern void gomp_unmap_vars (struct target_mem_desc *, bool);
-
extern void gomp_init_device (struct gomp_device_descr *);
-
extern void gomp_init_tables (struct gomp_device_descr *,
struct gomp_memory_mapping *);
-
extern void gomp_free_memmap (struct gomp_memory_mapping *);
-
extern void gomp_fini_device (struct gomp_device_descr *);
/* work.c */
diff --git libgomp/libgomp.map libgomp/libgomp.map
index bfdb78c..59ed4ee 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -318,7 +318,6 @@ GOACC_2.0 {
GOACC_data_end;
GOACC_data_start;
GOACC_enter_exit_data;
- GOACC_kernels;
GOACC_parallel;
GOACC_update;
GOACC_wait;
@@ -326,13 +325,18 @@ GOACC_2.0 {
GOACC_get_num_threads;
};
+GOACC_2.0.GOMP_4_BRANCH {
+ global:
+ GOACC_kernels;
+} GOACC_2.0;
+
GOMP_PLUGIN_1.0 {
global:
GOMP_PLUGIN_malloc;
GOMP_PLUGIN_malloc_cleared;
GOMP_PLUGIN_realloc;
- GOMP_PLUGIN_error;
GOMP_PLUGIN_debug;
+ GOMP_PLUGIN_error;
GOMP_PLUGIN_fatal;
GOMP_PLUGIN_async_unmap_vars;
GOMP_PLUGIN_acc_thread;
diff --git libgomp/libgomp.texi libgomp/libgomp.texi
index c4713ad..ff093d7 100644
--- libgomp/libgomp.texi
+++ libgomp/libgomp.texi
@@ -143,6 +143,11 @@ A complete description of all OpenACC directives accepted may be found in
the @uref{http://www.openacc.org/, OpenMP Application Programming
Interface} manual, version 2.0.
+Note that this is an experimental feature, incomplete, and subject to
+change in future versions of GCC. See
+@uref{https://gcc.gnu.org/wiki/OpenACC} for more information.
+
+
@c ---------------------------------------------------------------------
@c OpenACC Runtime Library Routines
diff --git libgomp/oacc-cuda.c libgomp/oacc-cuda.c
index 6f1a06f..c8ef376 100644
--- libgomp/oacc-cuda.c
+++ libgomp/oacc-cuda.c
@@ -74,7 +74,7 @@ acc_set_cuda_stream (int async, void *stream)
if (async < 0 || stream == NULL)
return 0;
-
+
goacc_lazy_initialize ();
if (base_dev && base_dev->openacc.cuda.set_stream_func)
diff --git libgomp/oacc-host.c libgomp/oacc-host.c
index 21c1ead..6aeb1e7 100644
--- libgomp/oacc-host.c
+++ libgomp/oacc-host.c
@@ -33,32 +33,31 @@
static struct gomp_device_descr host_dispatch =
{
.name = "host",
-
- .type = OFFLOAD_TARGET_TYPE_HOST,
.capabilities = (GOMP_OFFLOAD_CAP_OPENACC_200
| GOMP_OFFLOAD_CAP_NATIVE_EXEC
| GOMP_OFFLOAD_CAP_SHARED_MEM),
-
- .is_initialized = false,
- .offload_regions_registered = false,
+ .target_id = 0,
+ .type = OFFLOAD_TARGET_TYPE_HOST,
.get_name_func = GOMP_OFFLOAD_get_name,
- .get_type_func = GOMP_OFFLOAD_get_type,
.get_caps_func = GOMP_OFFLOAD_get_caps,
-
- .init_device_func = GOMP_OFFLOAD_init_device,
- .fini_device_func = GOMP_OFFLOAD_fini_device,
+ .get_type_func = GOMP_OFFLOAD_get_type,
.get_num_devices_func = GOMP_OFFLOAD_get_num_devices,
.register_image_func = GOMP_OFFLOAD_register_image,
+ .init_device_func = GOMP_OFFLOAD_init_device,
+ .fini_device_func = GOMP_OFFLOAD_fini_device,
.get_table_func = GOMP_OFFLOAD_get_table,
-
.alloc_func = GOMP_OFFLOAD_alloc,
.free_func = GOMP_OFFLOAD_free,
- .host2dev_func = GOMP_OFFLOAD_host2dev,
.dev2host_func = GOMP_OFFLOAD_dev2host,
-
+ .host2dev_func = GOMP_OFFLOAD_host2dev,
.run_func = GOMP_OFFLOAD_run,
+ .mem_map.is_initialized = false,
+ .mem_map.splay_tree.root = NULL,
+ .is_initialized = false,
+ .offload_regions_registered = false,
+
.openacc = {
.open_device_func = GOMP_OFFLOAD_openacc_open_device,
.close_device_func = GOMP_OFFLOAD_openacc_close_device,
diff --git libgomp/oacc-init.c libgomp/oacc-init.c
index 6f4a32c..166eb55 100644
--- libgomp/oacc-init.c
+++ libgomp/oacc-init.c
@@ -33,13 +33,12 @@
#include <stdlib.h>
#include <strings.h>
#include <stdbool.h>
-#include <stdio.h>
#include <string.h>
static gomp_mutex_t acc_device_lock;
/* The dispatch table for the current accelerator device. This is global, so
- you can only have one type of device open at any given time in a program.
+ you can only have one type of device open at any given time in a program.
This is the "base" device in that several devices that use the same
dispatch table may be active concurrently: this one (the "zeroth") is used
for overall initialisation/shutdown, and other instances -- not necessarily
@@ -129,7 +128,7 @@ resolve_device (acc_device_t d)
if (dispatchers[d] && dispatchers[d]->get_num_devices_func () > 0)
goto found;
if (d_arg == acc_device_default)
- {
+ {
d = acc_device_host;
goto found;
}
@@ -171,7 +170,7 @@ acc_init_1 (acc_device_t d)
gomp_fatal ("device already active");
/* We need to remember what we were intialized as, to check shutdown etc. */
- init_key = d;
+ init_key = d;
gomp_init_device (acc_dev);
@@ -203,9 +202,9 @@ static void
goacc_destroy_thread (void *data)
{
struct goacc_thread *thr = data, *walk, *prev;
-
+
gomp_mutex_lock (&goacc_thread_lock);
-
+
if (thr)
{
if (base_dev && thr->target_tls)
@@ -500,7 +499,7 @@ acc_get_device_num (acc_device_t d)
num = dev->openacc.get_device_num_func ();
if (num < 0)
num = goacc_device_num;
-
+
return num;
}
@@ -514,11 +513,11 @@ acc_set_device_num (int n, acc_device_t d)
if (!base_dev)
gomp_init_targets_once ();
-
+
if ((int) d == 0)
{
int i;
-
+
/* A device setting of zero sets all device types on the system to use
the Nth instance of that device type. Only attempt it for initialized
devices though. */
@@ -571,6 +570,7 @@ acc_on_device (acc_device_t dev)
/* Just rely on the compiler builtin. */
return __builtin_acc_on_device (dev);
}
+
ialias (acc_on_device)
attribute_hidden void
diff --git libgomp/oacc-int.h libgomp/oacc-int.h
index 700577d..85619c8 100644
--- libgomp/oacc-int.h
+++ libgomp/oacc-int.h
@@ -35,8 +35,8 @@
that are part of the external ABI, and the lower case prefix "goacc"
is used group items that are completely private to the library. */
-#ifndef _OACC_INT_H
-#define _OACC_INT_H 1
+#ifndef OACC_INT_H
+#define OACC_INT_H 1
#include "openacc.h"
#include "config.h"
@@ -58,16 +58,16 @@ struct goacc_thread
{
/* The device for the current thread. */
struct gomp_device_descr *dev;
-
+
struct gomp_device_descr *saved_bound_dev;
/* This is a linked list of data mapped by the "acc data" pragma, following
strictly push/pop semantics according to lexical scope. */
struct target_mem_desc *mapped_data;
-
+
/* These structures form a list: this is the next thread in that list. */
struct goacc_thread *next;
-
+
/* Target-specific data (used by plugin). */
void *target_tls;
};
@@ -88,8 +88,6 @@ goacc_thread (void)
}
#endif
-struct gomp_device_descr;
-
void goacc_register (struct gomp_device_descr *) __GOACC_NOTHROW;
/* Current dispatcher. */
@@ -104,4 +102,4 @@ void goacc_lazy_initialize (void);
# pragma GCC visibility pop
#endif
-#endif /* _OACC_INT_H */
+#endif
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 674fb76..0096d51 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -31,15 +31,13 @@
#include "libgomp.h"
#include "gomp-constants.h"
#include "oacc-int.h"
-#include <stdio.h>
+#include "splay-tree.h"
#include <stdint.h>
#include <assert.h>
-#include "splay-tree.h"
-
/* Return block containing [H->S), or NULL if not contained. */
-attribute_hidden splay_tree_key
+static splay_tree_key
lookup_host (struct gomp_memory_mapping *mem_map, void *h, size_t s)
{
struct splay_tree_key_s node;
@@ -342,11 +340,11 @@ acc_unmap_data (void *h)
gomp_unmap_vars (t, true);
}
-#define PCC_Present (1 << 0)
-#define PCC_Create (1 << 1)
-#define PCC_Copy (1 << 2)
+#define FLAG_PRESENT (1 << 0)
+#define FLAG_CREATE (1 << 1)
+#define FLAG_COPY (1 << 2)
-attribute_hidden void *
+static void *
present_create_copy (unsigned f, void *h, size_t s)
{
void *d;
@@ -366,13 +364,13 @@ present_create_copy (unsigned f, void *h, size_t s)
/* Present. */
d = (void *) (n->tgt->tgt_start + n->tgt_offset);
- if (!(f & PCC_Present))
+ if (!(f & FLAG_PRESENT))
gomp_fatal ("[%p,+%d] already mapped to [%p,+%d]",
(void *)h, (int)s, (void *)d, (int)s);
if ((h + s) > (void *)n->host_end)
gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
}
- else if (!(f & PCC_Create))
+ else if (!(f & FLAG_CREATE))
{
gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
}
@@ -383,7 +381,7 @@ present_create_copy (unsigned f, void *h, size_t s)
unsigned short kinds;
void *hostaddrs = h;
- if (f & PCC_Copy)
+ if (f & FLAG_COPY)
kinds = GOMP_MAP_TO;
else
kinds = GOMP_MAP_ALLOC;
@@ -406,28 +404,28 @@ present_create_copy (unsigned f, void *h, size_t s)
void *
acc_create (void *h, size_t s)
{
- return present_create_copy (PCC_Create, h, s);
+ return present_create_copy (FLAG_CREATE, h, s);
}
void *
acc_copyin (void *h, size_t s)
{
- return present_create_copy (PCC_Create | PCC_Copy, h, s);
+ return present_create_copy (FLAG_CREATE | FLAG_COPY, h, s);
}
void *
acc_present_or_create (void *h, size_t s)
{
- return present_create_copy (PCC_Present | PCC_Create, h, s);
+ return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s);
}
void *
acc_present_or_copyin (void *h, size_t s)
{
- return present_create_copy (PCC_Present | PCC_Create | PCC_Copy, h, s);
+ return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
}
-#define DC_Copyout (1 << 0)
+#define FLAG_COPYOUT (1 << 0)
static void
delete_copyout (unsigned f, void *h, size_t s)
@@ -454,7 +452,7 @@ delete_copyout (unsigned f, void *h, size_t s)
gomp_fatal ("[%p,%d] surrounds2 [%p,+%d]",
(void *) n->host_start, (int) host_size, (void *) h, (int) s);
- if (f & DC_Copyout)
+ if (f & FLAG_COPYOUT)
acc_dev->dev2host_func (acc_dev->target_id, h, d, s);
acc_unmap_data (h);
@@ -470,7 +468,7 @@ acc_delete (void *h , size_t s)
void acc_copyout (void *h, size_t s)
{
- delete_copyout (DC_Copyout, h, s);
+ delete_copyout (FLAG_COPYOUT, h, s);
}
static void
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index b6ee7c1..3bbf8de 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -31,7 +31,6 @@
#include "libgomp_g.h"
#include "gomp-constants.h"
#include "oacc-int.h"
-#include <stdio.h>
#include <string.h>
#include <stdarg.h>
#include <assert.h>
diff --git libgomp/oacc-plugin.c libgomp/oacc-plugin.c
index 1fd6b2d..44cd6d6 100644
--- libgomp/oacc-plugin.c
+++ libgomp/oacc-plugin.c
@@ -34,7 +34,7 @@ void
GOMP_PLUGIN_async_unmap_vars (void *ptr)
{
struct target_mem_desc *tgt = ptr;
-
+
gomp_unmap_vars (tgt, false);
}
diff --git libgomp/oacc-plugin.h libgomp/oacc-plugin.h
index 33cec79..c60eb9c 100644
--- libgomp/oacc-plugin.h
+++ libgomp/oacc-plugin.h
@@ -24,10 +24,10 @@
see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
<http://www.gnu.org/licenses/>. */
-#ifndef _OACC_PLUGIN_H
-#define _OACC_PLUGIN_H 1
+#ifndef OACC_PLUGIN_H
+#define OACC_PLUGIN_H 1
-extern void GOMP_PLUGIN_async_unmap_vars (void *ptr);
+extern void GOMP_PLUGIN_async_unmap_vars (void *);
extern void *GOMP_PLUGIN_acc_thread (void);
#endif
diff --git libgomp/openacc.h libgomp/openacc.h
index daf7c95..3343241 100644
--- libgomp/openacc.h
+++ libgomp/openacc.h
@@ -29,7 +29,7 @@
#ifndef _OPENACC_H
#define _OPENACC_H 1
-/* The OpenACC std is silent on whether or not including openacc.h
+/* The OpenACC standard is silent on whether or not including <openacc.h>
might or must not include other header files. We chose to include
some. */
#include <stddef.h>
@@ -46,80 +46,71 @@ extern "C" {
# define __GOACC_NOTHROW __attribute__ ((__nothrow__))
#endif
- /* Types */
- typedef enum acc_device_t
- {
- /* Keep in sync with include/gomp-constants.h. */
- acc_device_none = 0,
- acc_device_default = 1,
- acc_device_host = 2,
- acc_device_host_nonshm = 3,
- acc_device_not_host = 4,
- acc_device_nvidia = 5,
- _ACC_device_hwm
- } acc_device_t;
+/* Types */
+typedef enum acc_device_t
+ {
+ /* Keep in sync with include/gomp-constants.h. */
+ acc_device_none = 0,
+ acc_device_default = 1,
+ acc_device_host = 2,
+ acc_device_host_nonshm = 3,
+ acc_device_not_host = 4,
+ acc_device_nvidia = 5,
+ _ACC_device_hwm
+ } acc_device_t;
- typedef enum acc_async_t
- {
- /* Keep in sync with include/gomp-constants.h. */
- acc_async_noval = -1,
- acc_async_sync = -2
- } acc_async_t;
+typedef enum acc_async_t
+ {
+ /* Keep in sync with include/gomp-constants.h. */
+ acc_async_noval = -1,
+ acc_async_sync = -2
+ } acc_async_t;
- int acc_get_num_devices (acc_device_t __dev) __GOACC_NOTHROW;
- void acc_set_device_type (acc_device_t __dev) __GOACC_NOTHROW;
- acc_device_t acc_get_device_type (void) __GOACC_NOTHROW;
- void acc_set_device_num (int __num, acc_device_t __dev) __GOACC_NOTHROW;
- int acc_get_device_num (acc_device_t __dev) __GOACC_NOTHROW;
- int acc_async_test (int __async) __GOACC_NOTHROW;
- int acc_async_test_all (void) __GOACC_NOTHROW;
- void acc_wait (int __async) __GOACC_NOTHROW;
- void acc_wait_async (int __async1, int __async2) __GOACC_NOTHROW;
- void acc_wait_all (void) __GOACC_NOTHROW;
- void acc_wait_all_async (int __async) __GOACC_NOTHROW;
- void acc_init (acc_device_t __dev) __GOACC_NOTHROW;
- void acc_shutdown (acc_device_t __dev) __GOACC_NOTHROW;
- int acc_on_device (acc_device_t __dev) __GOACC_NOTHROW;
- void *acc_malloc (size_t) __GOACC_NOTHROW;
- void acc_free (void *) __GOACC_NOTHROW;
- /* Some of these would be more correct with const qualifiers, but
- the standard specifies otherwise. */
- void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
- void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
- void *acc_create (void *, size_t) __GOACC_NOTHROW;
- void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
- void acc_copyout (void *, size_t) __GOACC_NOTHROW;
- void acc_delete (void *, size_t) __GOACC_NOTHROW;
- void acc_update_device (void *, size_t) __GOACC_NOTHROW;
- void acc_update_self (void *, size_t) __GOACC_NOTHROW;
- void acc_map_data (void *, void *, size_t) __GOACC_NOTHROW;
- void acc_unmap_data (void *) __GOACC_NOTHROW;
- void *acc_deviceptr (void *) __GOACC_NOTHROW;
- void *acc_hostptr (void *) __GOACC_NOTHROW;
- int acc_is_present (void *, size_t) __GOACC_NOTHROW;
- void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
- void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
+int acc_get_num_devices (acc_device_t) __GOACC_NOTHROW;
+void acc_set_device_type (acc_device_t) __GOACC_NOTHROW;
+acc_device_t acc_get_device_type (void) __GOACC_NOTHROW;
+void acc_set_device_num (int, acc_device_t) __GOACC_NOTHROW;
+int acc_get_device_num (acc_device_t) __GOACC_NOTHROW;
+int acc_async_test (int) __GOACC_NOTHROW;
+int acc_async_test_all (void) __GOACC_NOTHROW;
+void acc_wait (int) __GOACC_NOTHROW;
+void acc_wait_async (int, int) __GOACC_NOTHROW;
+void acc_wait_all (void) __GOACC_NOTHROW;
+void acc_wait_all_async (int) __GOACC_NOTHROW;
+void acc_init (acc_device_t) __GOACC_NOTHROW;
+void acc_shutdown (acc_device_t) __GOACC_NOTHROW;
+int acc_on_device (acc_device_t) __GOACC_NOTHROW;
+void *acc_malloc (size_t) __GOACC_NOTHROW;
+void acc_free (void *) __GOACC_NOTHROW;
+/* Some of these would be more correct with const qualifiers, but
+ the standard specifies otherwise. */
+void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
+void *acc_create (void *, size_t) __GOACC_NOTHROW;
+void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
+void acc_copyout (void *, size_t) __GOACC_NOTHROW;
+void acc_delete (void *, size_t) __GOACC_NOTHROW;
+void acc_update_device (void *, size_t) __GOACC_NOTHROW;
+void acc_update_self (void *, size_t) __GOACC_NOTHROW;
+void acc_map_data (void *, void *, size_t) __GOACC_NOTHROW;
+void acc_unmap_data (void *) __GOACC_NOTHROW;
+void *acc_deviceptr (void *) __GOACC_NOTHROW;
+void *acc_hostptr (void *) __GOACC_NOTHROW;
+int acc_is_present (void *, size_t) __GOACC_NOTHROW;
+void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
+void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
- void ACC_target (int, void (*) (void *), const void *,
- size_t, void **, size_t *, unsigned char *, int *) __GOACC_NOTHROW;
- void ACC_parallel (int, void (*) (void *), const void *,
- size_t, void **, size_t *, unsigned char *) __GOACC_NOTHROW;
- void ACC_add_device_code (void const *, char const *) __GOACC_NOTHROW;
+/* Old names. OpenACC does not specify whether these can or must
+ not be macros, inlines or aliases for the new names. */
+#define acc_pcreate acc_present_or_create
+#define acc_pcopyin acc_present_or_copyin
- void ACC_async_copy (int) __GOACC_NOTHROW;
- void ACC_async_kern (int) __GOACC_NOTHROW;
+/* CUDA-specific routines. */
+void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
+void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
+void *acc_get_cuda_stream (int) __GOACC_NOTHROW;
+int acc_set_cuda_stream (int, void *) __GOACC_NOTHROW;
- /* Old names. OpenACC does not specify whether these can or must
- not be macros, inlines or aliases for the new names. */
- #define acc_pcreate acc_present_or_create
- #define acc_pcopyin acc_present_or_copyin
-
- /* CUDA-specific routines. */
- void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
- void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
- void *acc_get_cuda_stream (int __async) __GOACC_NOTHROW;
- int acc_set_cuda_stream (int __async, void *__stream) __GOACC_NOTHROW;
-
#ifdef __cplusplus
}
#endif
diff --git libgomp/plugin/plugin-host.c libgomp/plugin/plugin-host.c
index 7437407..ebf7f11 100644
--- libgomp/plugin/plugin-host.c
+++ libgomp/plugin/plugin-host.c
@@ -55,10 +55,6 @@
#define SELF "host: "
#endif
-#ifndef HOST_NONSHM_PLUGIN
-static struct gomp_device_descr host_dispatch;
-#endif
-
STATIC const char *
GOMP_OFFLOAD_get_name (void)
{
@@ -99,24 +95,24 @@ GOMP_OFFLOAD_get_num_devices (void)
}
STATIC void
-GOMP_OFFLOAD_register_image (void *host_table __attribute__((unused)),
- void *target_data __attribute__((unused)))
+GOMP_OFFLOAD_register_image (void *host_table __attribute__ ((unused)),
+ void *target_data __attribute__ ((unused)))
{
}
STATIC void
-GOMP_OFFLOAD_init_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_init_device (int n __attribute__ ((unused)))
{
}
STATIC void
-GOMP_OFFLOAD_fini_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_fini_device (int n __attribute__ ((unused)))
{
}
STATIC int
-GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
- struct mapping_table **table __attribute__((unused)))
+GOMP_OFFLOAD_get_table (int n __attribute__ ((unused)),
+ struct mapping_table **table __attribute__ ((unused)))
{
return 0;
}
@@ -143,23 +139,23 @@ STATIC void
GOMP_OFFLOAD_openacc_set_device_num (int n)
{
if (n > 0)
- GOMP(fatal) ("device number %u out of range for host execution", n);
+ GOMP (fatal) ("device number %u out of range for host execution", n);
}
STATIC void *
-GOMP_OFFLOAD_alloc (int n __attribute__((unused)), size_t s)
+GOMP_OFFLOAD_alloc (int n __attribute__ ((unused)), size_t s)
{
- return GOMP(malloc) (s);
+ return GOMP (malloc) (s);
}
STATIC void
-GOMP_OFFLOAD_free (int n __attribute__((unused)), void *p)
+GOMP_OFFLOAD_free (int n __attribute__ ((unused)), void *p)
{
free (p);
}
STATIC void *
-GOMP_OFFLOAD_host2dev (int n __attribute__((unused)), void *d, const void *h,
+GOMP_OFFLOAD_host2dev (int n __attribute__ ((unused)), void *d, const void *h,
size_t s)
{
#ifdef HOST_NONSHM_PLUGIN
@@ -170,7 +166,7 @@ GOMP_OFFLOAD_host2dev (int n __attribute__((unused)), void *d, const void *h,
}
STATIC void *
-GOMP_OFFLOAD_dev2host (int n __attribute__((unused)), void *h, const void *d,
+GOMP_OFFLOAD_dev2host (int n __attribute__ ((unused)), void *h, const void *d,
size_t s)
{
#ifdef HOST_NONSHM_PLUGIN
@@ -181,7 +177,7 @@ GOMP_OFFLOAD_dev2host (int n __attribute__((unused)), void *h, const void *d,
}
STATIC void
-GOMP_OFFLOAD_run (int n __attribute__((unused)), void *fn_ptr, void *vars)
+GOMP_OFFLOAD_run (int n __attribute__ ((unused)), void *fn_ptr, void *vars)
{
void (*fn)(void *) = (void (*)(void *)) fn_ptr;
@@ -190,16 +186,16 @@ GOMP_OFFLOAD_run (int n __attribute__((unused)), void *fn_ptr, void *vars)
STATIC void
GOMP_OFFLOAD_openacc_parallel (void (*fn) (void *),
- size_t mapnum __attribute__((unused)),
- void **hostaddrs __attribute__((unused)),
- void **devaddrs __attribute__((unused)),
- size_t *sizes __attribute__((unused)),
- unsigned short *kinds __attribute__((unused)),
- int num_gangs __attribute__((unused)),
- int num_workers __attribute__((unused)),
- int vector_length __attribute__((unused)),
- int async __attribute__((unused)),
- void *targ_mem_desc __attribute__((unused)))
+ size_t mapnum __attribute__ ((unused)),
+ void **hostaddrs __attribute__ ((unused)),
+ void **devaddrs __attribute__ ((unused)),
+ size_t *sizes __attribute__ ((unused)),
+ unsigned short *kinds __attribute__ ((unused)),
+ int num_gangs __attribute__ ((unused)),
+ int num_workers __attribute__ ((unused)),
+ int vector_length __attribute__ ((unused)),
+ int async __attribute__ ((unused)),
+ void *targ_mem_desc __attribute__ ((unused)))
{
#ifdef HOST_NONSHM_PLUGIN
fn (devaddrs);
@@ -219,12 +215,12 @@ GOMP_OFFLOAD_openacc_register_async_cleanup (void *targ_mem_desc)
}
STATIC void
-GOMP_OFFLOAD_openacc_async_set_async (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_set_async (int async __attribute__ ((unused)))
{
}
STATIC int
-GOMP_OFFLOAD_openacc_async_test (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_test (int async __attribute__ ((unused)))
{
return 1;
}
@@ -236,7 +232,7 @@ GOMP_OFFLOAD_openacc_async_test_all (void)
}
STATIC void
-GOMP_OFFLOAD_openacc_async_wait (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_wait (int async __attribute__ ((unused)))
{
}
@@ -246,25 +242,25 @@ GOMP_OFFLOAD_openacc_async_wait_all (void)
}
STATIC void
-GOMP_OFFLOAD_openacc_async_wait_async (int async1 __attribute__((unused)),
- int async2 __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_wait_async (int async1 __attribute__ ((unused)),
+ int async2 __attribute__ ((unused)))
{
}
STATIC void
-GOMP_OFFLOAD_openacc_async_wait_all_async (int async __attribute__((unused)))
+GOMP_OFFLOAD_openacc_async_wait_all_async (int async __attribute__ ((unused)))
{
}
STATIC void *
GOMP_OFFLOAD_openacc_create_thread_data (void *targ_data
- __attribute__((unused)))
+ __attribute__ ((unused)))
{
return NULL;
}
STATIC void
GOMP_OFFLOAD_openacc_destroy_thread_data (void *tls_data
- __attribute__((unused)))
+ __attribute__ ((unused)))
{
}
diff --git libgomp/plugin/plugin-nvptx.c libgomp/plugin/plugin-nvptx.c
index ee0c818..483cb75 100644
--- libgomp/plugin/plugin-nvptx.c
+++ libgomp/plugin/plugin-nvptx.c
@@ -738,7 +738,7 @@ link_ptx (CUmodule *module, char *ptx_code)
CUlinkState linkstate;
CUresult r;
void *linkout;
- size_t linkoutsize __attribute__((unused));
+ size_t linkoutsize __attribute__ ((unused));
GOMP_PLUGIN_debug (0, "attempting to load:\n---\n%s\n---\n", ptx_code);
@@ -1518,19 +1518,19 @@ GOMP_OFFLOAD_register_image (void *host_table, void *target_data)
}
void
-GOMP_OFFLOAD_init_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_init_device (int n __attribute__ ((unused)))
{
(void) nvptx_init ();
}
void
-GOMP_OFFLOAD_fini_device (int n __attribute__((unused)))
+GOMP_OFFLOAD_fini_device (int n __attribute__ ((unused)))
{
nvptx_fini ();
}
int
-GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
+GOMP_OFFLOAD_get_table (int n __attribute__ ((unused)),
struct mapping_table **tablep)
{
CUmodule module;
@@ -1589,26 +1589,26 @@ GOMP_OFFLOAD_get_table (int n __attribute__((unused)),
}
void *
-GOMP_OFFLOAD_alloc (int n __attribute__((unused)), size_t size)
+GOMP_OFFLOAD_alloc (int n __attribute__ ((unused)), size_t size)
{
return nvptx_alloc (size);
}
void
-GOMP_OFFLOAD_free (int n __attribute__((unused)), void *ptr)
+GOMP_OFFLOAD_free (int n __attribute__ ((unused)), void *ptr)
{
nvptx_free (ptr);
}
void *
-GOMP_OFFLOAD_dev2host (int ord __attribute__((unused)), void *dst,
+GOMP_OFFLOAD_dev2host (int ord __attribute__ ((unused)), void *dst,
const void *src, size_t n)
{
return nvptx_dev2host (dst, src, n);
}
void *
-GOMP_OFFLOAD_host2dev (int ord __attribute__((unused)), void *dst,
+GOMP_OFFLOAD_host2dev (int ord __attribute__ ((unused)), void *dst,
const void *src, size_t n)
{
return nvptx_host2dev (dst, src, n);
diff --git libgomp/splay-tree.c libgomp/splay-tree.c
index 3e4a904..030ca8f 100644
--- libgomp/splay-tree.c
+++ libgomp/splay-tree.c
@@ -26,14 +26,7 @@
/* The splay tree code copied from include/splay-tree.h and adjusted,
so that all the data lives directly in splay_tree_node_s structure
- and no extra allocations are needed.
-
- Files including this header should before including it add:
-typedef struct splay_tree_node_s *splay_tree_node;
-typedef struct splay_tree_s *splay_tree;
-typedef struct splay_tree_key_s *splay_tree_key;
- define splay_tree_key_s structure, and define
- splay_compare inline function. */
+ and no extra allocations are needed. */
/* For an easily readable description of splay-trees, see:
diff --git libgomp/splay-tree.h libgomp/splay-tree.h
index d04e13c..085021c 100644
--- libgomp/splay-tree.h
+++ libgomp/splay-tree.h
@@ -46,27 +46,6 @@ typedef struct splay_tree_key_s *splay_tree_key;
#ifndef _SPLAY_TREE_H
#define _SPLAY_TREE_H 1
-typedef struct splay_tree_node_s *splay_tree_node;
-typedef struct splay_tree_s *splay_tree;
-typedef struct splay_tree_key_s *splay_tree_key;
-
-struct splay_tree_key_s {
- /* Address of the host object. */
- uintptr_t host_start;
- /* Address immediately after the host object. */
- uintptr_t host_end;
- /* Descriptor of the target memory. */
- struct target_mem_desc *tgt;
- /* Offset from tgt->tgt_start to the start of the target object. */
- uintptr_t tgt_offset;
- /* Reference count. */
- uintptr_t refcount;
- /* Asynchronous reference count. */
- uintptr_t async_refcount;
- /* True if data should be copied from device to host at the end. */
- bool copy_from;
-};
-
/* The nodes in the splay tree. */
struct splay_tree_node_s {
struct splay_tree_key_s key;
diff --git libgomp/target.c libgomp/target.c
index 6871e7b..83ad511 100644
--- libgomp/target.c
+++ libgomp/target.c
@@ -34,7 +34,6 @@
#include <stdbool.h>
#include <stdlib.h>
#include <string.h>
-#include <stdio.h>
#include <assert.h>
#ifdef PLUGIN_SUPPORT
@@ -168,6 +167,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
tgt_align = align;
tgt_size = mapnum * sizeof (void *);
}
+
gomp_mutex_lock (&mm->lock);
for (i = 0; i < mapnum; i++)
@@ -288,9 +288,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
tgt->refcount++;
array->left = NULL;
array->right = NULL;
-
splay_tree_insert (&mm->splay_tree, array);
-
switch (kind & typemask)
{
case GOMP_MAP_ALLOC:
@@ -302,7 +300,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
case GOMP_MAP_TOFROM:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_TOFROM:
- /* Copy from host to device memory. */
/* FIXME: Perhaps add some smarts, like if copying
several adjacent fields from host to target, use some
host buffer to avoid sending each var individually. */
@@ -318,7 +315,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
if (cur_node.host_start == (uintptr_t) NULL)
{
cur_node.tgt_offset = (uintptr_t) NULL;
- /* Copy from host to device memory. */
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -346,7 +342,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
if (n == NULL)
gomp_fatal ("Pointer target of array section "
"wasn't mapped");
-
cur_node.host_start -= n->host_start;
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start;
@@ -354,7 +349,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
array section. Now subtract bias to get what we want
to initialize the pointer with. */
cur_node.tgt_offset -= sizes[i];
- /* Copy from host to device memory. */
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -363,7 +357,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
sizeof (void *));
break;
case GOMP_MAP_TO_PSET:
- /* Copy from host to device memory. */
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -388,7 +381,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
if (cur_node.host_start == (uintptr_t) NULL)
{
cur_node.tgt_offset = (uintptr_t) NULL;
- /* Copy from host to device memory. */
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start + k->tgt_offset
@@ -428,7 +420,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
array section. Now subtract bias to get what we
want to initialize the pointer with. */
cur_node.tgt_offset -= sizes[j];
- /* Copy from host to device memory. */
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start + k->tgt_offset
@@ -476,7 +467,6 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
else
cur_node.tgt_offset = tgt->list[i]->tgt->tgt_start
+ tgt->list[i]->tgt_offset;
- /* Copy from host to device memory. */
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -528,7 +518,6 @@ gomp_copy_from_async (struct target_mem_desc *tgt)
{
splay_tree_key k = tgt->list[i];
if (k->copy_from)
- /* Copy from device to host memory. */
devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
(void *) (k->tgt->tgt_start + k->tgt_offset),
k->host_end - k->host_start);
@@ -567,7 +556,6 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
{
splay_tree_key k = tgt->list[i];
if (k->copy_from && do_copyfrom)
- /* Copy from device to host memory. */
devicep->dev2host_func (devicep->target_id, (void *) k->host_start,
(void *) (k->tgt->tgt_start + k->tgt_offset),
k->host_end - k->host_start);
@@ -621,7 +609,6 @@ gomp_update (struct gomp_device_descr *devicep, struct gomp_memory_mapping *mm,
(void *) n->host_start,
(void *) n->host_end);
if (GOMP_MAP_COPY_TO_P (kind & typemask))
- /* Copy from host to device memory. */
devicep->host2dev_func (devicep->target_id,
(void *) (n->tgt->tgt_start
+ n->tgt_offset
@@ -630,7 +617,6 @@ gomp_update (struct gomp_device_descr *devicep, struct gomp_memory_mapping *mm,
(void *) cur_node.host_start,
cur_node.host_end - cur_node.host_start);
if (GOMP_MAP_COPY_FROM_P (kind & typemask))
- /* Copy from device to host memory. */
devicep->dev2host_func (devicep->target_id,
(void *) cur_node.host_start,
(void *) (n->tgt->tgt_start
@@ -647,9 +633,6 @@ gomp_update (struct gomp_device_descr *devicep, struct gomp_memory_mapping *mm,
gomp_mutex_unlock (&mm->lock);
}
-static void gomp_register_image_for_device (struct gomp_device_descr *device,
- struct offload_image_descr *image);
-
/* This function should be called from every offload image.
It gets the descriptor of the host func and var tables HOST_TABLE, TYPE of
the target, and TARGET_DATA needed by target plugin. */
@@ -662,9 +645,6 @@ GOMP_offload_register (void *host_table, enum offload_target_type target_type,
(num_offload_images + 1)
* sizeof (struct offload_image_descr));
- if (offload_images == NULL)
- return;
-
offload_images[num_offload_images].type = target_type;
offload_images[num_offload_images].host_table = host_table;
offload_images[num_offload_images].target_data = target_data;
@@ -691,9 +671,10 @@ gomp_init_tables (struct gomp_device_descr *devicep,
{
/* Get address mapping table for device. */
struct mapping_table *table = NULL;
- int i, num_entries = devicep->get_table_func (devicep->target_id, &table);
+ int num_entries = devicep->get_table_func (devicep->target_id, &table);
/* Insert host-target address mapping into dev_splay_tree. */
+ int i;
for (i = 0; i < num_entries; i++)
{
struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt));
@@ -1116,12 +1097,13 @@ gomp_target_init (void)
break;
}
- current_device.type = current_device.get_type_func ();
current_device.name = current_device.get_name_func ();
+ /* current_device.capabilities has already been set. */
+ current_device.type = current_device.get_type_func ();
+ current_device.mem_map.is_initialized = false;
+ current_device.mem_map.splay_tree.root = NULL;
current_device.is_initialized = false;
current_device.offload_regions_registered = false;
- current_device.mem_map.splay_tree.root = NULL;
- current_device.mem_map.is_initialized = false;
current_device.openacc.data_environ = NULL;
current_device.openacc.target_data = NULL;
for (i = 0; i < new_num_devices; i++)
diff --git libgomp/testsuite/lib/libgomp.exp libgomp/testsuite/lib/libgomp.exp
index ca62897..dded36a 100644
--- libgomp/testsuite/lib/libgomp.exp
+++ libgomp/testsuite/lib/libgomp.exp
@@ -36,12 +36,17 @@ load_gcc_lib fortran-modules.exp
load_file libgomp-test-support.exp
# Populate offload_targets_s (offloading targets separated by a space), and
-# offload_targets_s_openacc (the same, but with OpenACC names).
+# offload_targets_s_openacc (the same, but with OpenACC names; OpenACC spells
+# some of them a little differently).
set offload_targets_s [split $offload_targets ","]
-# OpenACC spells some of them a little differently.
set offload_targets_s_openacc {}
foreach offload_target_openacc $offload_targets_s {
switch $offload_target_openacc {
+ intelmic {
+ # TODO. Skip; will all FAIL because of missing
+ # GOMP_OFFLOAD_CAP_OPENACC_200.
+ continue
+ }
nvptx {
set offload_target_openacc "nvidia"
}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/abort.c libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c
similarity index 100%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/abort.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
index 80c51d5..747109f 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
@@ -140,7 +140,7 @@ main (int argc, char **argv)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
-#pragma acc delete (N)
+#pragma acc exit data delete (N)
#pragma acc wait (1)
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
index 2a77936..184b355 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
@@ -538,7 +538,76 @@ main(int argc, char **argv)
{
if (b[i] != 18.0)
abort ();
- }
+ }
+
+#pragma acc enter data copyin (b[0:N]) if (0)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (0)
+
+#pragma acc enter data copyin (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+ if (!acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (zero)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (zero)
+
+#pragma acc enter data copyin (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+ if (!acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (one == 0)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 0)
+
+#pragma acc enter data copyin (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (!acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (b, N * sizeof (float)))
+ abort ();
+#endif
return 0;
}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
index 0579185..c164598 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
@@ -6,30 +6,136 @@ int
main (int argc, char *argv[])
{
#define N 10
- char a[N];
-
- {
+ char a[N];
int i;
+
for (i = 0; i < N; ++i)
- a[i] = 0;
- }
+ a[i] = 0;
-#pragma acc data copyout (a)
- {
-#pragma acc parallel /* will result in a "dummy frame" */ present (a)
+#pragma acc data copy (a)
{
- int i;
- for (i = 0; i < N; ++i)
- a[i] = i;
+#pragma acc parallel present (a)
+ {
+ int j;
+
+ for (j = 0; j < N; ++j)
+ a[j] = j;
+ }
}
- }
- {
- int i;
for (i = 0; i < N; ++i)
- if (a[i] != i)
- abort ();
- }
+ {
+ if (a[i] != i)
+ abort ();
+ }
+
+ for (i = 0; i < N; ++i)
+ a[i] = 0;
+
+#pragma acc data copy (a)
+ {
+#pragma acc kernels present (a)
+ {
+ int j;
+
+ for (j = 0; j < N; ++j)
+ a[j] = j;
+ }
+ }
+
+ for (i = 0; i < N; ++i)
+ {
+ if (a[i] != i)
+ abort ();
+ }
+
+ for (i = 0; i < N; ++i)
+ a[i] = 0;
+
+#pragma acc data copy (a)
+ {
+#pragma acc data present (a)
+ {
+#pragma acc parallel present (a)
+ {
+ int j;
+
+ for (j = 0; j < N; ++j)
+ a[j] = j;
+ }
+ }
+ }
+
+ for (i = 0; i < N; ++i)
+ {
+ if (a[i] != i)
+ abort ();
+ }
+
+#pragma acc data copy (a)
+ {
+#pragma acc data present (a)
+ {
+#pragma acc kernels present (a)
+ {
+ int j;
+
+ for (j = 0; j < N; ++j)
+ a[j] = j;
+ }
+ }
+ }
+
+ for (i = 0; i < N; ++i)
+ {
+ if (a[i] != i)
+ abort ();
+ }
+
+ for (i = 0; i < N; ++i)
+ a[i] = 0;
+
+#pragma acc enter data copyin (a)
+
+#pragma acc data present (a)
+ {
+#pragma acc parallel present (a)
+ {
+ int j;
+
+ for (j = 0; j < N; ++j)
+ a[j] = j;
+ }
+ }
+
+#pragma acc exit data copyout (a)
+
+ for (i = 0; i < N; ++i)
+ {
+ if (a[i] != i)
+ abort ();
+ }
+
+#pragma acc enter data copyin (a)
+
+#pragma acc data present (a)
+ {
+#pragma acc kernels present (a)
+ {
+ int j;
+
+ for (j = 0; j < N; ++j)
+ a[j] = j;
+ }
+ }
+
+#pragma acc exit data copyout (a)
+
+ for (i = 0; i < N; ++i)
+ {
+ if (a[i] != i)
+ abort ();
+ }
- return 0;
+ return 0;
}
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-11.f90 libgomp/testsuite/libgomp.oacc-fortran/lib-11.f90
deleted file mode 100644
index a54d6a7..0000000
--- libgomp/testsuite/libgomp.oacc-fortran/lib-11.f90
+++ /dev/null
@@ -1,82 +0,0 @@
-! { dg-do run }
-
-program main
- implicit none
- include "openacc_lib.h"
-
- integer, target :: a_3d_i(10, 10, 10)
- complex a_3d_c(10, 10, 10)
- real a_3d_r(10, 10, 10)
-
- integer i, j, k
- complex c
- real r
- integer, parameter :: i_size = sizeof (i)
- integer, parameter :: c_size = sizeof (c)
- integer, parameter :: r_size = sizeof (r)
-
- if (acc_get_num_devices (acc_device_nvidia) .eq. 0) call exit
-
- call acc_init (acc_device_nvidia)
-
- call set3d (.FALSE., a_3d_i, a_3d_c, a_3d_r)
-
- call acc_copyin (a_3d_i)
- call acc_copyin (a_3d_c)
- call acc_copyin (a_3d_r)
-
- if (acc_is_present (a_3d_i) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_c) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_r) .neqv. .TRUE.) call abort
-
- do i = 1, 10
- do j = 1, 10
- do k = 1, 10
- if (acc_is_present (a_3d_i(i, j, k), i_size) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_c(i, j, k), i_size) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_r(i, j, k), i_size) .neqv. .TRUE.) call abort
- end do
- end do
- end do
-
- call acc_shutdown (acc_device_nvidia)
-
-contains
-
- subroutine set3d (clear, a_i, a_c, a_r)
- logical clear
- integer, dimension (:,:,:), intent (inout) :: a_i
- complex, dimension (:,:,:), intent (inout) :: a_c
- real, dimension (:,:,:), intent (inout) :: a_r
-
- integer i, j, k
- integer lb1, ub1, lb2, ub2, lb3, ub3
-
- lb1 = lbound (a_i, 1)
- ub1 = ubound (a_i, 1)
-
- lb2 = lbound (a_i, 2)
- ub2 = ubound (a_i, 2)
-
- lb3 = lbound (a_i, 3)
- ub3 = ubound (a_i, 3)
-
- do i = lb1, ub1
- do j = lb2, ub2
- do k = lb3, ub3
- if (clear) then
- a_i(i, j, k) = 0
- a_c(i, j, k) = cmplx (0.0, 0.0)
- a_r(i, j, k) = 0.0
- else
- a_i(i, j, k) = i
- a_c(i, j, k) = cmplx (i, j)
- a_r(i, j, k) = i
- end if
- end do
- end do
- end do
-
- end subroutine
-
-end program
diff --git libgomp/testsuite/libgomp.oacc-fortran/lib-9.f90 libgomp/testsuite/libgomp.oacc-fortran/lib-9.f90
deleted file mode 100644
index ad758b2..0000000
--- libgomp/testsuite/libgomp.oacc-fortran/lib-9.f90
+++ /dev/null
@@ -1,83 +0,0 @@
-! { dg-do run }
-
-program main
- use openacc
- use iso_c_binding
- implicit none
-
- integer, target :: a_3d_i(10, 10, 10)
- complex a_3d_c(10, 10, 10)
- real a_3d_r(10, 10, 10)
-
- integer i, j, k
- complex c
- real r
- integer, parameter :: i_size = sizeof (i)
- integer, parameter :: c_size = sizeof (c)
- integer, parameter :: r_size = sizeof (r)
-
- if (acc_get_num_devices (acc_device_nvidia) .eq. 0) call exit
-
- call acc_init (acc_device_nvidia)
-
- call set3d (.FALSE., a_3d_i, a_3d_c, a_3d_r)
-
- call acc_copyin (a_3d_i)
- call acc_copyin (a_3d_c)
- call acc_copyin (a_3d_r)
-
- if (acc_is_present (a_3d_i) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_c) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_r) .neqv. .TRUE.) call abort
-
- do i = 1, 10
- do j = 1, 10
- do k = 1, 10
- if (acc_is_present (a_3d_i(i, j, k), i_size) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_c(i, j, k), i_size) .neqv. .TRUE.) call abort
- if (acc_is_present (a_3d_r(i, j, k), i_size) .neqv. .TRUE.) call abort
- end do
- end do
- end do
-
- call acc_shutdown (acc_device_nvidia)
-
-contains
-
- subroutine set3d (clear, a_i, a_c, a_r)
- logical clear
- integer, dimension (:,:,:), intent (inout) :: a_i
- complex, dimension (:,:,:), intent (inout) :: a_c
- real, dimension (:,:,:), intent (inout) :: a_r
-
- integer i, j, k
- integer lb1, ub1, lb2, ub2, lb3, ub3
-
- lb1 = lbound (a_i, 1)
- ub1 = ubound (a_i, 1)
-
- lb2 = lbound (a_i, 2)
- ub2 = ubound (a_i, 2)
-
- lb3 = lbound (a_i, 3)
- ub3 = ubound (a_i, 3)
-
- do i = lb1, ub1
- do j = lb2, ub2
- do k = lb3, ub3
- if (clear) then
- a_i(i, j, k) = 0
- a_c(i, j, k) = cmplx (0.0, 0.0)
- a_r(i, j, k) = 0.0
- else
- a_i(i, j, k) = i
- a_c(i, j, k) = cmplx (i, j)
- a_r(i, j, k) = i
- end if
- end do
- end do
- end do
-
- end subroutine
-
-end program
diff --git liboffloadmic/ChangeLog liboffloadmic/ChangeLog
index 9faa452..074926e 100644
--- liboffloadmic/ChangeLog
+++ liboffloadmic/ChangeLog
@@ -1,3 +1,8 @@
+2015-01-15 Thomas Schwinge <thomas@codesourcery.com>
+
+ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_name)
+ (GOMP_OFFLOAD_get_caps, GOMP_OFFLOAD_fini_device): New functions.
+
2014-11-13 Ilya Verbin <ilya.verbin@intel.com>
Andrey Turetskiy <andrey.turetskiy@intel.com>
GrÃÃe,
Thomas
Attachment:
pgpO5IQ0sO3Kc.pgp
Description: PGP signature
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |