OpenMP 5.0: requires directive
This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2021-January/563393.html This patch completes more of the reverse_offload, unified_address, and unified_shared_memory clauses for the OpenMP 5.0 requires directive, including runtime verification of the offload target. (currently no offload devices actually support above features, only warning messages are emitted) This may possibly reverted/updated when a final patch is approved for mainline. 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com> gcc/c/ChangeLog: * c-parser.cc (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 yet" for OMP_REQUIRES_DYNAMIC_ALLOCATORS. gcc/cp/ChangeLog: * parser.cc (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 yet" for OMP_REQUIRES_DYNAMIC_ALLOCATORS. gcc/fortran/ChangeLog: * openmp.cc (gfc_check_omp_requires): Fix REVERSE_OFFLOAD typo. (gfc_match_omp_requires): Adjust to only mention "not implemented yet" for OMP_REQUIRES_DYNAMIC_ALLOCATORS. * parse.cc ("tree.h"): Add include. ("omp-general.h"): Likewise. (gfc_parse_file): Add code to merge omp_requires to omp_requires_mask. gcc/ChangeLog: * omp-offload.cc (omp_finish_file): Add code to create OpenMP requires mask variable in .gnu.gomp_requires section if needed. gcc/testsuite/ChangeLog: * c-c++-common/gomp/requires-4.c: Remove prune of "not supported yet". * gfortran.dg/gomp/requires-4.f90: Fix REVERSE_OFFLOAD typo. * gfortran.dg/gomp/requires-8.f90: Likewise. include/ChangeLog: * gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS): New symbol. (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY): Likewise. (GOMP_REQUIRES_REVERSE_OFFLOAD): Likewise. libgcc/ChangeLog: * 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/ChangeLog: * libgomp-plugin.h (GOMP_OFFLOAD_supported_features): New declaration. * libgomp.h (struct gomp_device_descr): New 'supported_features_func' 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 function. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_supported_features): Likewise. * target.c (<stdio.h>): 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 section 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/ChangeLog: * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_features): New function.
This commit is contained in:
parent
4b1205e579
commit
6ab6303f61
|
@ -1,3 +1,8 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* omp-offload.cc (omp_finish_file): Add code to create OpenMP requires
|
||||
mask variable in .gnu.gomp_requires section if needed.
|
||||
|
||||
2021-02-01 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* omp-low.cc (install_var_field): Add new 'tree key_expr = NULL_TREE'
|
||||
|
|
|
@ -1,3 +1,15 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* c-parser.cc (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 yet"
|
||||
for OMP_REQUIRES_DYNAMIC_ALLOCATORS.
|
||||
|
||||
2020-08-19 Sandra Loosemore <sandra@codesourcery.com>
|
||||
|
||||
Add a "combined" flag for "acc kernels loop" etc directives.
|
||||
|
|
|
@ -2480,6 +2480,12 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
|
|||
break;
|
||||
}
|
||||
|
||||
if (flag_openmp
|
||||
&& lookup_attribute ("omp declare target",
|
||||
DECL_ATTRIBUTES (current_function_decl)))
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
|
||||
|
||||
if (DECL_DECLARED_INLINE_P (current_function_decl))
|
||||
tv = TV_PARSE_INLINE;
|
||||
else
|
||||
|
@ -20781,6 +20787,10 @@ c_parser_omp_teams (location_t loc, c_parser *parser,
|
|||
static tree
|
||||
c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
|
||||
{
|
||||
if (flag_openmp)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
|
||||
|
||||
tree clauses
|
||||
= c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
|
||||
"#pragma omp target data");
|
||||
|
@ -20923,6 +20933,10 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
|
|||
return true;
|
||||
}
|
||||
|
||||
if (flag_openmp)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
|
||||
|
||||
tree clauses
|
||||
= c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
|
||||
"#pragma omp target enter data");
|
||||
|
@ -21009,6 +21023,10 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
|
|||
return true;
|
||||
}
|
||||
|
||||
if (flag_openmp)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
|
||||
|
||||
tree clauses
|
||||
= c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
|
||||
"#pragma omp target exit data");
|
||||
|
@ -22623,9 +22641,6 @@ c_parser_omp_requires (c_parser *parser)
|
|||
c_parser_skip_to_pragma_eol (parser, false);
|
||||
return;
|
||||
}
|
||||
if (p && this_req != OMP_REQUIRES_DYNAMIC_ALLOCATORS)
|
||||
sorry_at (cloc, "%qs clause on %<requires%> directive not "
|
||||
"supported yet", p);
|
||||
if (p)
|
||||
c_parser_consume_token (parser);
|
||||
if (this_req)
|
||||
|
|
|
@ -1,3 +1,15 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* parser.cc (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 yet"
|
||||
for OMP_REQUIRES_DYNAMIC_ALLOCATORS.
|
||||
|
||||
2021-02-01 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* semantics.cc (finish_omp_clauses): Adjust to allow duplicate
|
||||
|
|
|
@ -15354,6 +15354,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
|
||||
= (enum omp_requires) (omp_requires_mask
|
||||
| OMP_REQUIRES_TARGET_USED);
|
||||
pop_deferring_access_checks ();
|
||||
cp_finalize_omp_declare_simd (parser, &odsd);
|
||||
return;
|
||||
|
@ -43988,6 +43993,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
|
||||
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
|
||||
|
||||
tree clauses
|
||||
= cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
|
||||
"#pragma omp target data", pragma_tok);
|
||||
|
@ -44091,6 +44100,10 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
|
|||
return true;
|
||||
}
|
||||
|
||||
if (flag_openmp)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
|
||||
|
||||
tree clauses
|
||||
= cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
|
||||
"#pragma omp target enter data", pragma_tok);
|
||||
|
@ -44182,6 +44195,10 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
|
|||
return true;
|
||||
}
|
||||
|
||||
if (flag_openmp)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
|
||||
|
||||
tree clauses
|
||||
= cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
|
||||
"#pragma omp target exit data", pragma_tok);
|
||||
|
@ -46559,9 +46576,6 @@ cp_parser_omp_requires (cp_parser *parser, cp_token *pragma_tok)
|
|||
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
|
||||
return false;
|
||||
}
|
||||
if (p && this_req != OMP_REQUIRES_DYNAMIC_ALLOCATORS)
|
||||
sorry_at (cloc, "%qs clause on %<requires%> directive not "
|
||||
"supported yet", p);
|
||||
if (p)
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
if (this_req)
|
||||
|
|
|
@ -1,3 +1,12 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* openmp.cc (gfc_check_omp_requires): Fix REVERSE_OFFLOAD typo.
|
||||
(gfc_match_omp_requires): Adjust to only mention "not implemented yet"
|
||||
for OMP_REQUIRES_DYNAMIC_ALLOCATORS.
|
||||
* parse.cc ("tree.h"): Add include.
|
||||
("omp-general.h"): Likewise.
|
||||
(gfc_parse_file): Add code to merge omp_requires to omp_requires_mask.
|
||||
|
||||
2020-08-22 Sandra Loosemore <sandra@codesourcery.com>
|
||||
|
||||
Permit calls to Fortran intrinsics when annotating loops in
|
||||
|
|
|
@ -5225,7 +5225,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/routines "
|
||||
"but does not set !$OMP REQUIRES REVERSE_OFFSET but other "
|
||||
"but does not set !$OMP REQUIRES REVERSE_OFFLOAD but other "
|
||||
"program units do", &ns->proc_name->declared_at);
|
||||
if ((ref_omp_requires & OMP_REQ_UNIFIED_ADDRESS)
|
||||
&& !(ns->omp_requires & OMP_REQ_UNIFIED_ADDRESS))
|
||||
|
@ -5412,10 +5412,6 @@ gfc_match_omp_requires (void)
|
|||
else
|
||||
goto error;
|
||||
|
||||
if (requires_clause & ~(OMP_REQ_ATOMIC_MEM_ORDER_MASK
|
||||
| OMP_REQ_DYNAMIC_ALLOCATORS))
|
||||
gfc_error_now ("Sorry, %qs clause at %L on REQUIRES directive is not "
|
||||
"yet supported", clause, &old_loc);
|
||||
if (!gfc_omp_requires_add_clause (requires_clause, clause, &old_loc, NULL))
|
||||
goto error;
|
||||
requires_clauses |= requires_clause;
|
||||
|
|
|
@ -6890,6 +6890,23 @@ done:
|
|||
gfc_current_ns = gfc_current_ns->sibling)
|
||||
gfc_check_omp_requires (gfc_current_ns, omp_requires);
|
||||
|
||||
if (omp_requires)
|
||||
{
|
||||
omp_requires_mask = (enum omp_requires) OMP_REQUIRES_TARGET_USED;
|
||||
if (omp_requires & OMP_REQ_REVERSE_OFFLOAD)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask
|
||||
| OMP_REQUIRES_REVERSE_OFFLOAD);
|
||||
if (omp_requires & OMP_REQ_UNIFIED_ADDRESS)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask
|
||||
| OMP_REQUIRES_UNIFIED_ADDRESS);
|
||||
if (omp_requires & OMP_REQ_UNIFIED_SHARED_MEMORY)
|
||||
omp_requires_mask
|
||||
= (enum omp_requires) (omp_requires_mask
|
||||
| OMP_REQUIRES_UNIFIED_SHARED_MEMORY);
|
||||
}
|
||||
|
||||
/* Populate omp_requires_mask (needed for resolving OpenMP
|
||||
metadirectives and declare variant). */
|
||||
switch (omp_requires & OMP_REQ_ATOMIC_MEM_ORDER_MASK)
|
||||
|
|
|
@ -439,6 +439,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) != 0)
|
||||
{
|
||||
const char *requires_section = ".gnu.gomp_requires";
|
||||
tree maskvar = 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) = 1;
|
||||
DECL_INITIAL (maskvar)
|
||||
= 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
|
||||
{
|
||||
|
|
|
@ -1,3 +1,9 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* c-c++-common/gomp/requires-4.c: Remove prune of "not supported yet".
|
||||
* gfortran.dg/gomp/requires-4.f90: Fix REVERSE_OFFLOAD typo.
|
||||
* gfortran.dg/gomp/requires-8.f90: Likewise.
|
||||
|
||||
2021-02-01 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* c-c++-common/gomp/clauses-2.c: Adjust testcase.
|
||||
|
|
|
@ -9,5 +9,3 @@ foo (void)
|
|||
#pragma omp requires unified_shared_memory /* { dg-error "'unified_shared_memory' clause used lexically after first target construct or offloading API" } */
|
||||
#pragma omp requires unified_address /* { dg-error "'unified_address' clause used lexically after first target construct or offloading API" } */
|
||||
#pragma omp requires reverse_offload /* { dg-error "'reverse_offload' clause used lexically after first target construct or offloading API" } */
|
||||
|
||||
/* { dg-prune-output "not supported yet" } */
|
||||
|
|
|
@ -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 REQUIRES REVERSE_OFFSET but other program units do" "" { target *-*-* } 9 }
|
||||
! { dg-error "OpenMP device constructs/routines but does not set !.OMP REQUIRES REVERSE_OFFLOAD but other program units do" "" { target *-*-* } 9 }
|
||||
! { dg-error "OpenMP device constructs/routines but does not set !.OMP REQUIRES UNIFIED_ADDRESS but other program units do" "" { target *-*-* } 9 }
|
||||
! { dg-error "OpenMP device constructs/routines but does not set !.OMP REQUIRES UNIFIED_SHARED_MEMORY but other program units do" "" { target *-*-* } 9 }
|
||||
end
|
||||
|
|
|
@ -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
|
||||
|
|
|
@ -1,3 +1,9 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS): New symbol.
|
||||
(GOMP_REQUIRES_UNIFIED_SHARED_MEMORY): Likewise.
|
||||
(GOMP_REQUIRES_REVERSE_OFFLOAD): Likewise.
|
||||
|
||||
2020-07-27 Andrew Stubbs <ams@codesourcery.com>
|
||||
|
||||
* dwarf2.def (DW_OP_LLVM_piece_end): New extension operator.
|
||||
|
|
|
@ -356,6 +356,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. */
|
||||
|
|
|
@ -0,0 +1,6 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* 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.
|
|
@ -54,6 +54,9 @@ const void *const __offload_var_table[0]
|
|||
__attribute__ ((__used__, visibility ("hidden"),
|
||||
section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { };
|
||||
|
||||
const unsigned int const __requires_mask_table[0]
|
||||
__attribute__ ((__used__, section (".gnu.gomp_requires"))) = { };
|
||||
|
||||
#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))) = { };
|
||||
|
||||
const unsigned int const __requires_mask_table_end[0]
|
||||
__attribute__ ((__used__, section (".gnu.gomp_requires"))) = { };
|
||||
|
||||
#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
|
||||
|
|
|
@ -1,3 +1,17 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* libgomp-plugin.h (GOMP_OFFLOAD_supported_features): New declaration.
|
||||
* libgomp.h (struct gomp_device_descr): New 'supported_features_func'
|
||||
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 function.
|
||||
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_supported_features): Likewise.
|
||||
* target.c (<stdio.h>): Add include of standard header.
|
||||
(gomp_requires_mask): New static variable.
|
||||
(__requires_mask_table): New declaration.
|
||||
(__requires_mask_table_end): Likewise.
|
||||
|
||||
2021-01-13 Julian Brown <julian@codesourcery.com>
|
||||
|
||||
* testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: Adjust for loop
|
||||
|
|
|
@ -128,6 +128,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 **);
|
||||
|
|
|
@ -1233,6 +1233,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;
|
||||
|
|
|
@ -71,6 +71,12 @@ host_fini_device (int n __attribute__ ((unused)))
|
|||
return true;
|
||||
}
|
||||
|
||||
static bool
|
||||
host_supported_features (unsigned int *n)
|
||||
{
|
||||
return (*n == 0);
|
||||
}
|
||||
|
||||
static unsigned
|
||||
host_version (void)
|
||||
{
|
||||
|
@ -272,6 +278,7 @@ static struct gomp_device_descr host_dispatch =
|
|||
.get_num_devices_func = host_get_num_devices,
|
||||
.init_device_func = host_init_device,
|
||||
.fini_device_func = host_fini_device,
|
||||
.supported_features_func = host_supported_features,
|
||||
.version_func = host_version,
|
||||
.load_image_func = host_load_image,
|
||||
.unload_image_func = host_unload_image,
|
||||
|
|
|
@ -4089,4 +4089,12 @@ GOMP_OFFLOAD_openacc_destroy_thread_data (void *data)
|
|||
free (data);
|
||||
}
|
||||
|
||||
/* Indicate which GOMP_REQUIRES_* features are supported, currently none. */
|
||||
|
||||
bool
|
||||
GOMP_OFFLOAD_supported_features (unsigned int *mask)
|
||||
{
|
||||
return (*mask == 0);
|
||||
}
|
||||
|
||||
/* }}} */
|
||||
|
|
|
@ -1232,6 +1232,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 == 0);
|
||||
}
|
||||
|
||||
/* Return the libgomp version number we're compatible with. There is
|
||||
no requirement for cross-version compatibility. */
|
||||
|
||||
|
|
|
@ -31,6 +31,7 @@
|
|||
#include "gomp-constants.h"
|
||||
#include <limits.h>
|
||||
#include <stdbool.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#ifdef HAVE_INTTYPES_H
|
||||
# include <inttypes.h> /* For PRIu64. */
|
||||
|
@ -96,6 +97,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.requires
|
||||
section. Offload plugins are queried with this mask to see if all required
|
||||
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 *
|
||||
|
@ -2564,6 +2575,20 @@ gomp_init_device (struct gomp_device_descr *devicep)
|
|||
gomp_fatal ("device initialization failed");
|
||||
}
|
||||
|
||||
unsigned int features = gomp_requires_mask;
|
||||
if (!devicep->supported_features_func (&features))
|
||||
{
|
||||
char buf[64], *end = buf + sizeof (buf), *p = buf;
|
||||
if (features & GOMP_REQUIRES_UNIFIED_ADDRESS)
|
||||
p += snprintf (p, end - p, "unified_address");
|
||||
if (features & GOMP_REQUIRES_UNIFIED_SHARED_MEMORY)
|
||||
p += snprintf (p, end - p, "%sunified_shared_memory",
|
||||
(p == buf ? "" : ", "));
|
||||
if (features & GOMP_REQUIRES_REVERSE_OFFLOAD)
|
||||
p += snprintf (p, end - p, "%sreverse_offload", (p == buf ? "" : ", "));
|
||||
gomp_error ("device does not support required features: %s", buf);
|
||||
}
|
||||
|
||||
/* Load to device all images registered by the moment. */
|
||||
for (i = 0; i < num_offload_images; i++)
|
||||
{
|
||||
|
@ -3914,6 +3939,7 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
|
|||
DLSYM (get_num_devices);
|
||||
DLSYM (init_device);
|
||||
DLSYM (fini_device);
|
||||
DLSYM (supported_features);
|
||||
DLSYM (load_image);
|
||||
DLSYM (unload_image);
|
||||
DLSYM (alloc);
|
||||
|
@ -4024,6 +4050,28 @@ gomp_target_init (void)
|
|||
if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_DISABLED)
|
||||
return;
|
||||
|
||||
gomp_requires_mask = 0;
|
||||
const unsigned int *mask_ptr = __requires_mask_table;
|
||||
bool error_emitted = false;
|
||||
while (mask_ptr != __requires_mask_table_end)
|
||||
{
|
||||
if (gomp_requires_mask == 0)
|
||||
gomp_requires_mask = *mask_ptr;
|
||||
else if (gomp_requires_mask != *mask_ptr)
|
||||
{
|
||||
if (!error_emitted)
|
||||
{
|
||||
gomp_error ("requires-directive clause inconsistency between "
|
||||
"compilation units detected");
|
||||
error_emitted = true;
|
||||
}
|
||||
/* This is inconsistent, but still merge to query for all features
|
||||
later. */
|
||||
gomp_requires_mask |= *mask_ptr;
|
||||
}
|
||||
mask_ptr++;
|
||||
}
|
||||
|
||||
cur = OFFLOAD_PLUGINS;
|
||||
if (*cur)
|
||||
do
|
||||
|
|
|
@ -0,0 +1,11 @@
|
|||
/* { dg-skip-if "" { *-*-* } } */
|
||||
|
||||
#pragma omp requires reverse_offload
|
||||
|
||||
int x;
|
||||
|
||||
void foo (void)
|
||||
{
|
||||
#pragma omp target
|
||||
x = 1;
|
||||
}
|
|
@ -0,0 +1,20 @@
|
|||
/* { 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 = 0; i < 10; i++)
|
||||
a[i] = 0;
|
||||
|
||||
foo ();
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-output "libgomp: requires-directive clause inconsistency between compilation units detected" } */
|
||||
/* { dg-prune-output "device does not support required features" } */
|
|
@ -0,0 +1,11 @@
|
|||
/* { dg-skip-if "" { *-*-* } } */
|
||||
|
||||
#pragma omp requires reverse_offload
|
||||
|
||||
int x;
|
||||
|
||||
void foo (void)
|
||||
{
|
||||
#pragma omp target
|
||||
x = 1;
|
||||
}
|
|
@ -0,0 +1,19 @@
|
|||
/* { 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 = 0; i < 10; i++)
|
||||
a[i] = 0;
|
||||
|
||||
foo ();
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-output "libgomp: device does not support required features: reverse_offload" } */
|
|
@ -0,0 +1,4 @@
|
|||
2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
|
||||
|
||||
* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_features):
|
||||
New function.
|
|
@ -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 == 0);
|
||||
}
|
||||
|
||||
static bool
|
||||
get_target_table (int device, int &num_funcs, int &num_vars, void **&table)
|
||||
{
|
||||
|
|
Loading…
Reference in New Issue