From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id CCC2138515E3 for ; Thu, 25 Mar 2021 11:18:57 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org CCC2138515E3 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: ItJf/SxGv76AMWj1xEI7gGMD7Bh446UQfhT4/3523tSoTLlKZ13hh6vlgoEYi9B11o244PI/hp vsWe9WVFbkzFTdMP8ceU/SdRPMioLP4FVVdAP4OxzmfjFVaULC6XLNXuW6vkuspUcfP4T+JDbZ mmijR1ndgM7hVpqgqPDeVRNGhHSzMj86jf2hltnBClMRP2AE4euNXjfT+iIv0L4jBgaIhU8NcZ kfdsY5P8UJ+2s6/D4NBeJY22hcteJRMnhNgCyuGASqxR2h1oJVWzKzsBohr7p2ctOXp7clCKQI DvM= X-IronPort-AV: E=Sophos;i="5.81,277,1610438400"; d="scan'208,223";a="59600400" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 25 Mar 2021 03:18:56 -0800 IronPort-SDR: YAA1t3a3zcK+hg5BP1A35GHSviGkCbBsS8Bz7XP3sVzxhMiPshG7M5LDdbopqhoERnUltq40FW v9bWR/sBg8gprsQAlSzfjMraJgrJypnZWdefrJR221fYPuUQALq3wScGWede9Oh+UAag1PFxDP yhghXbtOcI4p03I1zTV4rK4bvD+9oGHQogS8euKeKJL+XxMnjJSA25dh0JZuLymCep2kK87GHn zNg3A8dNJbsm1ZSp3dROwKeBOULyLQIqOkYcwPevSeF8aZoRXApHPTYOU/vDTrLgIEYMgljrK2 D4M= From: Thomas Schwinge To: Chung-Lin Tang , CC: Jakub Jelinek , Catherine Moore , Andrew Stubbs , "Tobias Burnus" Subject: Re: [PATCH, OpenMP 5.0] More implementation of the requires directive In-Reply-To: <4273bf27-3f0e-0066-393b-2a561a7b9e12@codesourcery.com> References: <4273bf27-3f0e-0066-393b-2a561a7b9e12@codesourcery.com> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Thu, 25 Mar 2021 12:18:45 +0100 Message-ID: <87czvnzdka.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, KAM_STOCKGEN, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 25 Mar 2021 11:19:01 -0000 --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Hi! On 2021-01-13T23:07:44+0800, Chung-Lin Tang wrote= : > this patch provides more implementation of the requires directive, basica= lly: > > (1) The collection of the reverse_offload, unified_address, and unified_s= hared_memory > clauses into a .gnu.gomp_requires section > > (2) libgomp checking of consistency across the entire .gnu.gomp_requires = section, > and querying into the offload plugin to see if the offload target support= s the required > features (as of now, the setting is that none of those features are suppo= rted by any > of the plugins). > > We currently emit errors, but do not fatally cause exit of the program if= those > are not met. We're still unsure if complete block-out of program executio= n is the right > thing for the user. This can be discussed later. > > Is this okay for trunk after stage1 re-opens? (As posted, per a quick check) this got pushed to devel/omp/gcc-10 branch in commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0: requires directive". Building the libgomp Intel MIC plugin fails: make[3]: Entering directory '[...]/build-gcc-offload-x86_64-intelmicemu= l-linux-gnu/x86_64-intelmicemul-linux-gnu/liboffloadmic/plugin' [...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/./gcc/xg++ [...] = -loffloadmic_target -lcoi_device -lgomp -rdynamic ../ofldbegin.o offload_ta= rget_main.o ../ofldend.o -o offload_target_main ./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_ma= sk_table_end' ./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_ma= sk_table' collect2: error: ld returned 1 exit status Makefile:806: recipe for target 'offload_target_main' failed make[3]: *** [offload_target_main] Error 1 I've pushed "[WIP] OpenMP 5.0: requires directive: workaround to fix libgomp IntelMIC plugin build" to devel/omp/gcc-10 branch in commit ff77b4a0db75bc82a5519e31a882f9a25a02cd56, see attached. This seemed like a safe default, to get this un-stuck, but I suppose this will need further work. I haven't read up what this OpenMP functionality exactly is, and haven't thought about how it ought to be implemented -- but from a quick look, instead of libgomp directly referring to '__requires_mask_table', shouldn't this use some "dynamic indirection scheme" (like we have for the dynamic offloading code registering/loading function calls via constructors, synthesized by the 'mkoffload's?), so that it also works for shared objects ('*.so', etc.) containing OpenMP code? But maybe I just have no clue what I'm talking about, and this is not applicable here. ;-) 'make check-target-libgomp': libgomp: while loading libgomp-plugin-hsa.so.1: [...]/libgomp-plugin-hs= a.so.1: undefined symbol: GOMP_OFFLOAD_supported_features I've pushed "OpenMP 5.0: requires directive: adjust libgomp HSA plugin" to devel/omp/gcc-10 branch in commit 4ef4921cb10693c59b488002179db131683af8bc, see attached. (The libgomp HSA plugin has been removed in master branch, so not applicable there.) Gr=C3=BC=C3=9Fe Thomas > 2021-01-13 Chung-Lin Tang > > gcc/c/ > * c-parser.c (c_parser_declaration_or_fndef): Set > OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has > "omp declare target" attribute. > (c_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in > omp_requires_mask. > (c_parser_omp_target_enter_data): Likewise. > (c_parser_omp_target_exit_data): Likewise. > (c_parser_omp_requires): Adjust to only mention "not implemented ye= t" > for OMP_REQUIRES_DYNAMIC_ALLOCATORS. > > gcc/cp/ > * parser.c (cp_parser_simple_declaration): Set > OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has > "omp declare target" attribute. > (cp_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in > omp_requires_mask. > (cp_parser_omp_target_enter_data): Likewise. > (cp_parser_omp_target_exit_data): Likewise. > (cp_parser_omp_requires): Adjust to only mention "not implemented y= et" > for OMP_REQUIRES_DYNAMIC_ALLOCATORS. > > gcc/fortran/ > * openmp.c (gfc_check_omp_requires): Fix REVERSE_OFFLOAD typo. > (gfc_match_omp_requires): Adjust to only mention "not implemented y= et" > for OMP_REQUIRES_DYNAMIC_ALLOCATORS. > * parse.c ("tree.h"): Add include. > ("omp-general.h"): Likewise. > (gfc_parse_file): Add code to merge omp_requires to omp_requires_ma= sk. > > gcc/ > * omp-offload.c (omp_finish_file): Add code to reate OpenMP require= s > mask variable in .gnu.gomp_requires section if needed. > > gcc/testsuite/ > * c-c++-common/gomp/requires-4.c: Remove prune of "not supported ye= t". > * gcc/testsuite/gfortran.dg/gomp/requires-4.f90: Fix REVERSE_OFFLOA= D typo. > * gcc/testsuite/gfortran.dg/gomp/requires-8.f90: Likewise. > > include/ > * gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS): New symbol. > (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY): Likewise. > (GOMP_REQUIRES_REVERSE_OFFLOAD): Likewise. > > libgcc/ > * offloadstuff.c (__requires_mask_table): New symbol to mark start = of > .gnu.gomp_requires section. > (__requires_mask_table_end): New symbol to mark end of > .gnu.gomp_requires section. > > libgomp/ > * libgomp-plugin.h (GOMP_OFFLOAD_supported_features): New declarati= on. > * libgomp.h (struct gomp_device_descr): New 'supported_features_fun= c' > plugin hook field. > * oacc-host.c (host_supported_features): New host hook function. > (host_dispatch): Initialize 'supported_features_func' host hook. > * plugin/plugin-gcn.c (GOMP_OFFLOAD_supported_features): New functi= on. > * plugin/plugin-nvptx.c (GOMP_OFFLOAD_supported_features): Likewise= . > * target.c (): Add include of standard header. > (gomp_requires_mask): New static variable. > (__requires_mask_table): New declaration. > (__requires_mask_table_end): Likewise. > (gomp_load_plugin_for_device): Add loading of 'supported_features' = hook. > (gomp_target_init): Add code to summarize .gnu._gomp_requires secti= on > mask values, emit error if inconsistency found. > > * testsuite/libgomp.c-c++-common/requires-1.c: New test. > * testsuite/libgomp.c-c++-common/requires-1-aux.c: New file linked = with > above test. > * testsuite/libgomp.c-c++-common/requires-2.c: New test. > * testsuite/libgomp.c-c++-common/requires-2-aux.c: New file linked = with > above test. > > liboffloadmic/ > * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_featur= es): > New function. > diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c > index c77d9fccdc2..e685b26746e 100644 > --- a/gcc/c/c-parser.c > +++ b/gcc/c/c-parser.c > @@ -2475,6 +2475,12 @@ c_parser_declaration_or_fndef (c_parser *parser, b= ool fndef_ok, > break; > } > > + if (flag_openmp > + && lookup_attribute ("omp declare target", > + DECL_ATTRIBUTES (current_function_decl))) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_= USED); > + > if (DECL_DECLARED_INLINE_P (current_function_decl)) > tv =3D TV_PARSE_INLINE; > else > @@ -19556,6 +19562,10 @@ c_parser_omp_teams (location_t loc, c_parser *pa= rser, > static tree > c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) > { > + if (flag_openmp) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_U= SED); > + > tree clauses > =3D c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, > "#pragma omp target data"); > @@ -19698,6 +19708,10 @@ c_parser_omp_target_enter_data (location_t loc, = c_parser *parser, > return NULL_TREE; > } > > + if (flag_openmp) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_U= SED); > + > tree clauses > =3D c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_M= ASK, > "#pragma omp target enter data"); > @@ -19784,6 +19798,10 @@ c_parser_omp_target_exit_data (location_t loc, c= _parser *parser, > return NULL_TREE; > } > > + if (flag_openmp) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_U= SED); > + > tree clauses > =3D c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MA= SK, > "#pragma omp target exit data"); > @@ -21371,7 +21389,7 @@ c_parser_omp_requires (c_parser *parser) > c_parser_skip_to_pragma_eol (parser, false); > return; > } > - if (p) > + if (this_req =3D=3D OMP_REQUIRES_DYNAMIC_ALLOCATORS) > sorry_at (cloc, "%qs clause on % directive not " > "supported yet", p); > if (p) > diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c > index c713852fe93..afbc4e551d4 100644 > --- a/gcc/cp/parser.c > +++ b/gcc/cp/parser.c > @@ -14455,6 +14455,11 @@ cp_parser_simple_declaration (cp_parser* parser, > /* Otherwise, we're done with the list of declarators. */ > else > { > + if (flag_openmp && lookup_attribute ("omp declare target", > + DECL_ATTRIBUTES (decl))) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask > + | OMP_REQUIRES_TARGET_USED); > pop_deferring_access_checks (); > return; > } > @@ -41432,6 +41437,10 @@ cp_parser_omp_teams (cp_parser *parser, cp_token= *pragma_tok, > static tree > cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool= *if_p) > { > + if (flag_openmp) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_U= SED); > + > tree clauses > =3D cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, > "#pragma omp target data", pragma_tok); > @@ -41535,6 +41544,10 @@ cp_parser_omp_target_enter_data (cp_parser *pars= er, cp_token *pragma_tok, > return NULL_TREE; > } > > + if (flag_openmp) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_U= SED); > + > tree clauses > =3D cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_= MASK, > "#pragma omp target enter data", pragma_to= k); > @@ -41625,6 +41638,10 @@ cp_parser_omp_target_exit_data (cp_parser *parse= r, cp_token *pragma_tok, > return NULL_TREE; > } > > + if (flag_openmp) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_U= SED); > + > tree clauses > =3D cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_M= ASK, > "#pragma omp target exit data", pragma_tok= ); > @@ -43819,7 +43836,7 @@ cp_parser_omp_requires (cp_parser *parser, cp_tok= en *pragma_tok) > cp_parser_skip_to_pragma_eol (parser, pragma_tok); > return false; > } > - if (p) > + if (this_req =3D=3D OMP_REQUIRES_DYNAMIC_ALLOCATORS) > sorry_at (cloc, "%qs clause on % directive not " > "supported yet", p); > if (p) > diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c > index cb166f956b7..c25531a4989 100644 > --- a/gcc/fortran/openmp.c > +++ b/gcc/fortran/openmp.c > @@ -3668,7 +3668,7 @@ gfc_check_omp_requires (gfc_namespace *ns, int ref_= omp_requires) > if ((ref_omp_requires & OMP_REQ_REVERSE_OFFLOAD) > && !(ns->omp_requires & OMP_REQ_REVERSE_OFFLOAD)) > gfc_error ("Program unit at %L has OpenMP device constructs/routine= s " > - "but does not set !$OMP REQUIRES REVERSE_OFFSET but othe= r " > + "but does not set !$OMP REQUIRES REVERSE_OFFLOAD but oth= er " > "program units do", &ns->proc_name->declared_at); > if ((ref_omp_requires & OMP_REQ_UNIFIED_ADDRESS) > && !(ns->omp_requires & OMP_REQ_UNIFIED_ADDRESS)) > @@ -3855,7 +3855,8 @@ gfc_match_omp_requires (void) > else > goto error; > > - if (requires_clause & ~OMP_REQ_ATOMIC_MEM_ORDER_MASK) > + /* Currently, everything except 'dynamic_allocators' is allowed. = */ > + if (requires_clause =3D=3D OMP_REQ_DYNAMIC_ALLOCATORS) > gfc_error_now ("Sorry, %qs clause at %L on REQUIRES directive is no= t " > "yet supported", clause, &old_loc); > if (!gfc_omp_requires_add_clause (requires_clause, clause, &old_lo= c, NULL)) > diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c > index 1549f8e1635..4731bca2cf7 100644 > --- a/gcc/fortran/parse.c > +++ b/gcc/fortran/parse.c > @@ -22,10 +22,12 @@ along with GCC; see the file COPYING3. If not see > #include "system.h" > #include "coretypes.h" > #include "options.h" > +#include "tree.h" > #include "gfortran.h" > #include > #include "match.h" > #include "parse.h" > +#include "omp-general.h" > > /* Current statement label. Zero means no statement label. Because new= _st > can get wiped during statement matching, we have to keep it separate.= */ > @@ -6572,6 +6574,23 @@ done: > gfc_current_ns =3D gfc_current_ns->sibling) > gfc_check_omp_requires (gfc_current_ns, omp_requires); > > + if (omp_requires) > + { > + omp_requires_mask =3D (enum omp_requires) OMP_REQUIRES_TARGET_USED= ; > + if (omp_requires & OMP_REQ_REVERSE_OFFLOAD) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask > + | OMP_REQUIRES_REVERSE_OFFLOAD); > + if (omp_requires & OMP_REQ_UNIFIED_ADDRESS) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask > + | OMP_REQUIRES_UNIFIED_ADDRESS); > + if (omp_requires & OMP_REQ_UNIFIED_SHARED_MEMORY) > + omp_requires_mask > + =3D (enum omp_requires) (omp_requires_mask > + | OMP_REQUIRES_UNIFIED_SHARED_MEMORY); > + } > + > /* Do the parse tree dump. */ > gfc_current_ns =3D flag_dump_fortran_original ? gfc_global_ns_list : N= ULL; > > diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c > index ba0937fba94..9cc7d2945fc 100644 > --- a/gcc/omp-offload.c > +++ b/gcc/omp-offload.c > @@ -437,6 +437,24 @@ omp_finish_file (void) > > varpool_node::finalize_decl (vars_decl); > varpool_node::finalize_decl (funcs_decl); > + > + if (flag_openmp && (omp_requires_mask & OMP_REQUIRES_TARGET_USED) = !=3D 0) > + { > + const char *requires_section =3D ".gnu.gomp_requires"; > + tree maskvar =3D build_decl (UNKNOWN_LOCATION, VAR_DECL, > + get_identifier (".gomp_requires_mask")= , > + unsigned_type_node); > + SET_DECL_ALIGN (maskvar, TYPE_ALIGN (unsigned_type_node)); > + TREE_STATIC (maskvar) =3D 1; > + DECL_INITIAL (maskvar) > + =3D build_int_cst (unsigned_type_node, > + ((unsigned int) omp_requires_mask > + & (OMP_REQUIRES_UNIFIED_ADDRESS > + | OMP_REQUIRES_UNIFIED_SHARED_MEMORY > + | OMP_REQUIRES_REVERSE_OFFLOAD))); > + set_decl_section_name (maskvar, requires_section); > + varpool_node::finalize_decl (maskvar); > + } > } > else > { > diff --git a/gcc/testsuite/c-c++-common/gomp/requires-4.c b/gcc/testsuite= /c-c++-common/gomp/requires-4.c > index 88ba7746cf8..8f45d83ea6e 100644 > --- a/gcc/testsuite/c-c++-common/gomp/requires-4.c > +++ b/gcc/testsuite/c-c++-common/gomp/requires-4.c > @@ -9,5 +9,3 @@ foo (void) > #pragma omp requires unified_shared_memory /* { dg-error "'unified_sha= red_memory' clause used lexically after first target construct or offloadin= g API" } */ > #pragma omp requires unified_address /* { dg-error "'unified_address' cl= ause used lexically after first target construct or offloading API" } */ > #pragma omp requires reverse_offload /* { dg-error "'reverse_offload' cl= ause used lexically after first target construct or offloading API" } */ > - > -/* { dg-prune-output "not supported yet" } */ > diff --git a/gcc/testsuite/gfortran.dg/gomp/requires-4.f90 b/gcc/testsuit= e/gfortran.dg/gomp/requires-4.f90 > index b17aceb898b..c870a2840d3 100644 > --- a/gcc/testsuite/gfortran.dg/gomp/requires-4.f90 > +++ b/gcc/testsuite/gfortran.dg/gomp/requires-4.f90 > @@ -9,7 +9,7 @@ end module m > subroutine foo > !$omp target > !$omp end target > -! { dg-error "OpenMP device constructs/routines but does not set !.OMP R= EQUIRES REVERSE_OFFSET but other program units do" "" { target *-*-* } 9 } > +! { dg-error "OpenMP device constructs/routines but does not set !.OMP R= EQUIRES REVERSE_OFFLOAD but other program units do" "" { target *-*-* } 9 } > ! { dg-error "OpenMP device constructs/routines but does not set !.OMP R= EQUIRES UNIFIED_ADDRESS but other program units do" "" { target *-*-* } 9 } > ! { dg-error "OpenMP device constructs/routines but does not set !.OMP R= EQUIRES UNIFIED_SHARED_MEMORY but other program units do" "" { target *-*-*= } 9 } > end > diff --git a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 b/gcc/testsuit= e/gfortran.dg/gomp/requires-8.f90 > index 3c32ae9860e..3819b0c28cc 100644 > --- a/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 > +++ b/gcc/testsuite/gfortran.dg/gomp/requires-8.f90 > @@ -13,7 +13,7 @@ contains > end subroutine foo > end module m > > -subroutine bar ! { dg-error "has OpenMP device constructs/routines but = does not set !.OMP REQUIRES REVERSE_OFFSET but other program units do" } > +subroutine bar ! { dg-error "has OpenMP device constructs/routines but = does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" } > !use m > !$omp requires unified_shared_memory > !$omp declare target > diff --git a/include/gomp-constants.h b/include/gomp-constants.h > index 11a9308e3d2..d5a0b2c5ea7 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -301,6 +301,12 @@ enum gomp_map_kind > #define GOMP_DEPEND_INOUT 3 > #define GOMP_DEPEND_MUTEXINOUTSET 4 > > +/* Flag values for requires-directive features, must match corresponding > + OMP_REQUIRES_* values in gcc/omp-general.h. */ > +#define GOMP_REQUIRES_UNIFIED_ADDRESS 0x10 > +#define GOMP_REQUIRES_UNIFIED_SHARED_MEMORY 0x20 > +#define GOMP_REQUIRES_REVERSE_OFFLOAD 0x80 > + > /* HSA specific data structures. */ > > /* Identifiers of device-specific target arguments. */ > diff --git a/libgcc/offloadstuff.c b/libgcc/offloadstuff.c > index b19428af6d8..78210a88f15 100644 > --- a/libgcc/offloadstuff.c > +++ b/libgcc/offloadstuff.c > @@ -54,6 +54,9 @@ const void *const __offload_var_table[0] > __attribute__ ((__used__, visibility ("hidden"), > section (OFFLOAD_VAR_TABLE_SECTION_NAME))) =3D { }; > > +const unsigned int const __requires_mask_table[0] > + __attribute__ ((__used__, section (".gnu.gomp_requires"))) =3D { }; > + > #elif defined CRT_END > > const void *const __offload_funcs_end[0] > @@ -63,6 +66,9 @@ const void *const __offload_vars_end[0] > __attribute__ ((__used__, visibility ("hidden"), > section (OFFLOAD_VAR_TABLE_SECTION_NAME))) =3D { }; > > +const unsigned int const __requires_mask_table_end[0] > + __attribute__ ((__used__, section (".gnu.gomp_requires"))) =3D { }; > + > #elif defined CRT_TABLE > > extern const void *const __offload_func_table[]; > @@ -77,6 +83,9 @@ const void *const __OFFLOAD_TABLE__[] > &__offload_var_table, &__offload_vars_end > }; > > +extern const unsigned int const __requires_mask_table[]; > +extern const unsigned int const __requires_mask_table_end[]; > + > #else /* ! CRT_BEGIN && ! CRT_END && ! CRT_TABLE */ > #error "One of CRT_BEGIN, CRT_END or CRT_TABLE must be defined." > #endif > diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h > index 62645ce9954..f54469fdd6b 100644 > --- a/libgomp/libgomp-plugin.h > +++ b/libgomp/libgomp-plugin.h > @@ -122,6 +122,7 @@ extern int GOMP_OFFLOAD_get_type (void); > extern int GOMP_OFFLOAD_get_num_devices (void); > extern bool GOMP_OFFLOAD_init_device (int); > extern bool GOMP_OFFLOAD_fini_device (int); > +extern bool GOMP_OFFLOAD_supported_features (unsigned *); > extern unsigned GOMP_OFFLOAD_version (void); > extern int GOMP_OFFLOAD_load_image (int, unsigned, const void *, > struct addr_pair **); > diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h > index 305cba3aa02..09f2ac67943 100644 > --- a/libgomp/libgomp.h > +++ b/libgomp/libgomp.h > @@ -1130,6 +1130,7 @@ struct gomp_device_descr > __typeof (GOMP_OFFLOAD_get_num_devices) *get_num_devices_func; > __typeof (GOMP_OFFLOAD_init_device) *init_device_func; > __typeof (GOMP_OFFLOAD_fini_device) *fini_device_func; > + __typeof (GOMP_OFFLOAD_supported_features) *supported_features_func; > __typeof (GOMP_OFFLOAD_version) *version_func; > __typeof (GOMP_OFFLOAD_load_image) *load_image_func; > __typeof (GOMP_OFFLOAD_unload_image) *unload_image_func; > diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c > index f3bbd2b9c61..94a7fac2a39 100644 > --- a/libgomp/oacc-host.c > +++ b/libgomp/oacc-host.c > @@ -71,6 +71,12 @@ host_fini_device (int n __attribute__ ((unused))) > return true; > } > > +static bool > +host_supported_features (unsigned int *n) > +{ > + return (*n =3D=3D 0); > +} > + > static unsigned > host_version (void) > { > @@ -273,6 +279,7 @@ static struct gomp_device_descr host_dispatch =3D > .get_num_devices_func =3D host_get_num_devices, > .init_device_func =3D host_init_device, > .fini_device_func =3D host_fini_device, > + .supported_features_func =3D host_supported_features, > .version_func =3D host_version, > .load_image_func =3D host_load_image, > .unload_image_func =3D host_unload_image, > diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c > index 47f0b6e25f8..718d78173fe 100644 > --- a/libgomp/plugin/plugin-gcn.c > +++ b/libgomp/plugin/plugin-gcn.c > @@ -3991,4 +3991,12 @@ GOMP_OFFLOAD_openacc_destroy_thread_data (void *da= ta) > free (data); > } > > +/* Indicate which GOMP_REQUIRES_* features are supported, currently none= . */ > + > +bool > +GOMP_OFFLOAD_supported_features (unsigned int *mask) > +{ > + return (*mask =3D=3D 0); > +} > + > /* }}} */ > diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.= c > index 681c344b9c2..4cc25fbe232 100644 > --- a/libgomp/plugin/plugin-nvptx.c > +++ b/libgomp/plugin/plugin-nvptx.c > @@ -1236,6 +1236,14 @@ GOMP_OFFLOAD_fini_device (int n) > return true; > } > > +/* Indicate which GOMP_REQUIRES_* features are supported, currently none= . */ > + > +bool > +GOMP_OFFLOAD_supported_features (unsigned int *mask) > +{ > + return (*mask =3D=3D 0); > +} > + > /* Return the libgomp version number we're compatible with. There is > no requirement for cross-version compatibility. */ > > diff --git a/libgomp/target.c b/libgomp/target.c > index 4a4e1f80745..f06df7ba28d 100644 > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -31,6 +31,7 @@ > #include "gomp-constants.h" > #include > #include > +#include > #include > #ifdef HAVE_INTTYPES_H > # include /* For PRIu64. */ > @@ -79,6 +80,16 @@ static int num_devices; > /* Number of GOMP_OFFLOAD_CAP_OPENMP_400 devices. */ > static int num_devices_openmp; > > +/* Mask of requires directive clause values, summarized from .gnu.gomp.r= equires > + section. Offload plugins are queried with this mask to see if all req= uired > + features are supported. */ > +static unsigned int gomp_requires_mask; > + > +/* Start/end of .gnu.gomp.requires section of program, defined in > + crtoffloadbegin/end.o. */ > +extern const unsigned int __requires_mask_table[]; > +extern const unsigned int __requires_mask_table_end[]; > + > /* Similar to gomp_realloc, but release register_lock before gomp_fatal.= */ > > static void * > @@ -1961,6 +1972,20 @@ gomp_init_device (struct gomp_device_descr *device= p) > gomp_fatal ("device initialization failed"); > } > > + unsigned int features =3D gomp_requires_mask; > + if (!devicep->supported_features_func (&features)) > + { > + char buf[64], *end =3D buf + sizeof (buf), *p =3D buf; > + if (features & GOMP_REQUIRES_UNIFIED_ADDRESS) > + p +=3D snprintf (p, end - p, "unified_address"); > + if (features & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY) > + p +=3D snprintf (p, end - p, "%sunified_shared_memory", > + (p =3D=3D buf ? "" : ", ")); > + if (features & GOMP_REQUIRES_REVERSE_OFFLOAD) > + p +=3D snprintf (p, end - p, "%sreverse_offload", (p =3D=3D buf ? "= " : ", ")); > + gomp_error ("device does not support required features: %s", buf); > + } > + > /* Load to device all images registered by the moment. */ > for (i =3D 0; i < num_offload_images; i++) > { > @@ -3200,6 +3225,7 @@ gomp_load_plugin_for_device (struct gomp_device_des= cr *device, > DLSYM (get_num_devices); > DLSYM (init_device); > DLSYM (fini_device); > + DLSYM (supported_features); > DLSYM (load_image); > DLSYM (unload_image); > DLSYM (alloc); > @@ -3310,6 +3336,28 @@ gomp_target_init (void) > if (gomp_target_offload_var =3D=3D GOMP_TARGET_OFFLOAD_DISABLED) > return; > > + gomp_requires_mask =3D 0; > + const unsigned int *mask_ptr =3D __requires_mask_table; > + bool error_emitted =3D false; > + while (mask_ptr !=3D __requires_mask_table_end) > + { > + if (gomp_requires_mask =3D=3D 0) > + gomp_requires_mask =3D *mask_ptr; > + else if (gomp_requires_mask !=3D *mask_ptr) > + { > + if (!error_emitted) > + { > + gomp_error ("requires-directive clause inconsistency between = " > + "compilation units detected"); > + error_emitted =3D true; > + } > + /* This is inconsistent, but still merge to query for all feature= s > + later. */ > + gomp_requires_mask |=3D *mask_ptr; > + } > + mask_ptr++; > + } > + > cur =3D OFFLOAD_PLUGINS; > if (*cur) > do > diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c b/li= bgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c > new file mode 100644 > index 00000000000..8b9341523c6 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1-aux.c > @@ -0,0 +1,11 @@ > +/* { dg-skip-if "" { *-*-* } } */ > + > +#pragma omp requires reverse_offload > + > +int x; > + > +void foo (void) > +{ > + #pragma omp target > + x =3D 1; > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-1.c b/libgom= p/testsuite/libgomp.c-c++-common/requires-1.c > new file mode 100644 > index 00000000000..b5a3c512d28 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-1.c > @@ -0,0 +1,21 @@ > +/* { dg-additional-sources requires-1-aux.c } */ > + > +#pragma omp requires unified_shared_memory > + > +int a[10]; > +extern void foo (void); > + > +int > +main (void) > +{ > + #pragma omp target > + for (int i =3D 0; i < 10; i++) > + a[i] =3D 0; > + > + foo (); > + return 0; > +} > + > +/* { dg-output "libgomp: requires-directive clause inconsistency between= compilation units detected" } */ > +/* { dg-prune-output "device does not support required features" } */ > +/* { dg-shouldfail "" } */ > diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c b/li= bgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c > new file mode 100644 > index 00000000000..8b9341523c6 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2-aux.c > @@ -0,0 +1,11 @@ > +/* { dg-skip-if "" { *-*-* } } */ > + > +#pragma omp requires reverse_offload > + > +int x; > + > +void foo (void) > +{ > + #pragma omp target > + x =3D 1; > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-2.c b/libgom= p/testsuite/libgomp.c-c++-common/requires-2.c > new file mode 100644 > index 00000000000..6fb280baabd > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-2.c > @@ -0,0 +1,20 @@ > +/* { dg-additional-sources requires-2-aux.c } */ > + > +#pragma omp requires reverse_offload > + > +int a[10]; > +extern void foo (void); > + > +int > +main (void) > +{ > + #pragma omp target > + for (int i =3D 0; i < 10; i++) > + a[i] =3D 0; > + > + foo (); > + return 0; > +} > + > +/* { dg-output "libgomp: device does not support required features: reve= rse_offload" } */ > +/* { dg-shouldfail "" } */ > diff --git a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp b/liboffloa= dmic/plugin/libgomp-plugin-intelmic.cpp > index d1678d0514e..f92418fa416 100644 > --- a/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp > +++ b/liboffloadmic/plugin/libgomp-plugin-intelmic.cpp > @@ -233,6 +233,14 @@ GOMP_OFFLOAD_fini_device (int device) > return true; > } > > +/* Indicate which GOMP_REQUIRES_* features are supported, currently none= . */ > + > +extern "C" bool > +GOMP_OFFLOAD_supported_features (unsigned int *mask) > +{ > + return (*mask =3D=3D 0); > +} > + > static bool > get_target_table (int device, int &num_funcs, int &num_vars, void **&tab= le) > { ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 M=C3=BCnchen R= egistergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas H= eurung, Frank Th=C3=BCrauf --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-WIP-OpenMP-5.0-requires-directive-workaround-to.og10.patch" >From ff77b4a0db75bc82a5519e31a882f9a25a02cd56 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 3 Mar 2021 22:37:58 +0100 Subject: [PATCH 1/2] [WIP] OpenMP 5.0: requires directive: workaround to fix libgomp IntelMIC plugin build Fix-up for og10 commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0: requires directive". The GCC offloading target configurations don't build/use 'crtoffloadbegin.o'/'crtoffloadtable.o'/'crtoffloadend.o' ('libgcc/offloadstuff.c'), but the libgomp IntelMIC plugin still does link against libgomp, and the latter unconditionally refers to '__requires_mask_table', '__requires_mask_table_end': make[3]: Entering directory '[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/liboffloadmic/plugin' [...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/./gcc/xg++ [...] -loffloadmic_target -lcoi_device -lgomp -rdynamic ../ofldbegin.o offload_target_main.o ../ofldend.o -o offload_target_main ./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_mask_table_end' ./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_mask_table' collect2: error: ld returned 1 exit status Makefile:806: recipe for target 'offload_target_main' failed make[3]: *** [offload_target_main] Error 1 I have not researched what a proper fix would look like. libgomp/ * target.c (__requires_mask_table, __requires_mask_table_end): Add '__attribute__((weak))'. --- libgomp/ChangeLog.omp | 5 +++++ libgomp/target.c | 2 ++ 2 files changed, 7 insertions(+) diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 0e3fd122f850..03ca74c8f3d5 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,8 @@ +2021-03-25 Thomas Schwinge + + * target.c (__requires_mask_table, __requires_mask_table_end): Add + '__attribute__((weak))'. + 2021-03-01 Kwok Cheung Yeung * testsuite/libgomp.c-c++-common/collapse-4.c: New. diff --git a/libgomp/target.c b/libgomp/target.c index 699dc72cb722..9c7582635aa3 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -104,7 +104,9 @@ static unsigned int gomp_requires_mask; /* Start/end of .gnu.gomp.requires section of program, defined in crtoffloadbegin/end.o. */ +__attribute__((weak)) extern const unsigned int __requires_mask_table[]; +__attribute__((weak)) extern const unsigned int __requires_mask_table_end[]; /* Similar to gomp_realloc, but release register_lock before gomp_fatal. */ -- 2.30.2 --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0002-OpenMP-5.0-requires-directive-adjust-libgomp-HS.og10.patch" >From 4ef4921cb10693c59b488002179db131683af8bc Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 3 Mar 2021 22:51:01 +0100 Subject: [PATCH 2/2] OpenMP 5.0: requires directive: adjust libgomp HSA plugin Fix-up for og10 commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0: requires directive". libgomp: while loading libgomp-plugin-hsa.so.1: [...]/libgomp-plugin-hsa.so.1: undefined symbol: GOMP_OFFLOAD_supported_features libgomp/ * plugin/plugin-hsa.c (GOMP_OFFLOAD_supported_features): New function. --- libgomp/ChangeLog.omp | 3 +++ libgomp/plugin/plugin-hsa.c | 8 ++++++++ 2 files changed, 11 insertions(+) diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 03ca74c8f3d5..19f48dc61202 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,5 +1,8 @@ 2021-03-25 Thomas Schwinge + * plugin/plugin-hsa.c (GOMP_OFFLOAD_supported_features): New + function. + * target.c (__requires_mask_table, __requires_mask_table_end): Add '__attribute__((weak))'. diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c index abd3bc64163b..bddb690ca14f 100644 --- a/libgomp/plugin/plugin-hsa.c +++ b/libgomp/plugin/plugin-hsa.c @@ -1869,3 +1869,11 @@ GOMP_OFFLOAD_dev2dev (int ord, void *dst, const void *src, size_t n) "it should never be called"); return false; } + +/* Indicate which GOMP_REQUIRES_* features are supported, currently none. */ + +bool +GOMP_OFFLOAD_supported_features (unsigned int *mask) +{ + return (*mask == 0); +} -- 2.30.2 --=-=-=--