From patchwork Mon Feb 9 09:57:02 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 437847 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 08BCC14009B for ; Mon, 9 Feb 2015 21:17:52 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type; q=dns; s= default; b=JAZc59DjuXbrlLzN+iqtYBlicg7cgBXAu2NVyzHY3MqRisjdwlAqr 2JjyMhisv5CDf69D6FqxVLnUS3Jw33Z70DqfTeUNoWYi8LIpiolgBePYW+UqI5yD X/FT0diRNgNwwjLMkE9tJf/W6MPcdx66fj4y2ky5pjtda2xiaSJWdU= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type; s= default; bh=n+XCCHzyqylHmZ7rbUgremZ2N+M=; b=iBEGQOj0KWBwmgAq3Z/J FUtRu65z2Pmi2072nUilh/Y2UrDWd6/lib8BcN24UY3ppbJZ1W60QyaYAXUDqr8p MFJkfajRd/er3ebZYI0DW9Hjnt6nRM3zatAHPwuuhf3YPct5GrZuDKNZidhTsqRq DvgWzCFR2zXRh9XHluqEPuU= Received: (qmail 2157 invoked by alias); 9 Feb 2015 09:57:36 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 2145 invoked by uid 89); 9 Feb 2015 09:57:35 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.2 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 09 Feb 2015 09:57:23 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1YKl5h-0001PH-ON from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Mon, 09 Feb 2015 01:57:19 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Mon, 9 Feb 2015 09:57:15 +0000 From: Thomas Schwinge To: Subject: [gomp4] Merge trunk r219682 (2015-01-15) into gomp-4_0-branch User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (x86_64-pc-linux-gnu) Date: Mon, 9 Feb 2015 10:57:02 +0100 Message-ID: <87pp9jif0x.fsf@schwinge.name> MIME-Version: 1.0 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(-) Grüße, Thomas diff --git contrib/ChangeLog contrib/ChangeLog index 1f02d95..f062ea9 100644 --- contrib/ChangeLog +++ contrib/ChangeLog @@ -1,3 +1,8 @@ +2015-01-15 Thomas Schwinge + + * gcc_update (files_and_dependencies): Update rules for new + libgomp/plugin/Makefrag.am and libgomp/plugin/configfrag.ac files. + 2015-01-12 Yury Gribov * 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 + Bernd Schmidt + Cesar Philippidis + James Norris + Tom de Vries + Ilmir Usmanov + Dmitry Bocharnikov + Evgeny Gavrin + Jakub Jelinek + + * 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 (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 * 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 + + * gcc-interface/utils.c (DEF_FUNCTION_TYPE_VAR_8) + (DEF_FUNCTION_TYPE_VAR_12): New macros. + 2015-01-09 Michael Collison * 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 + Bernd Schmidt + James Norris + Cesar Philippidis + Ilmir Usmanov + Jakub Jelinek + + * 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 * 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 + Bernd Schmidt + Cesar Philippidis + James Norris + Jakub Jelinek + Ilmir Usmanov + + * 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 * 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 *c_parser_expr_list (c_parser *, bool, bool, vec **, location_t *, tree *, vec *, 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 + James Norris + Cesar Philippidis + Ilmir Usmanov + Jakub Jelinek + + * 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 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 + Cesar Philippidis + James Norris + Ilmir Usmanov + Tobias Burnus + + * 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 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 + James Norris + + * 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 * 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 + James Norris + Cesar Philippidis + Ilmir Usmanov + + * 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 * 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 - -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 + Julian Brown + James Norris + + * gomp-constants.h: New file. + 2015-12-14 Jan-Benedict Glaw * 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 + James Norris + Tom de Vries + Julian Brown + Cesar Philippidis + Nathan Sidwell + Tobias Burnus + + * 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 . + (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", . + (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 Julian Brown David Malcolm 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 #include #include -#include #include 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 +#include "splay-tree.h" #include #include -#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 #include #include #include 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 . */ -#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 might or must not include other header files. We chose to include some. */ #include @@ -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 #include #include -#include #include #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 + + * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_name) + (GOMP_OFFLOAD_get_caps, GOMP_OFFLOAD_fini_device): New functions. + 2014-11-13 Ilya Verbin Andrey Turetskiy