+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'
+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.
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
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");
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");
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");
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)
+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
/* 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;
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);
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);
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);
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)
+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
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))
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;
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)
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
{
+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.
#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" } */
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
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
+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.
#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. */
--- /dev/null
+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.
__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]
__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[];
&__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
+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
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 **);
__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;
return true;
}
+static bool
+host_supported_features (unsigned int *n)
+{
+ return (*n == 0);
+}
+
static unsigned
host_version (void)
{
.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,
free (data);
}
+/* Indicate which GOMP_REQUIRES_* features are supported, currently none. */
+
+bool
+GOMP_OFFLOAD_supported_features (unsigned int *mask)
+{
+ return (*mask == 0);
+}
+
/* }}} */
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. */
#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. */
/* 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 *
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++)
{
DLSYM (get_num_devices);
DLSYM (init_device);
DLSYM (fini_device);
+ DLSYM (supported_features);
DLSYM (load_image);
DLSYM (unload_image);
DLSYM (alloc);
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
--- /dev/null
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires reverse_offload
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
--- /dev/null
+/* { 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" } */
--- /dev/null
+/* { dg-skip-if "" { *-*-* } } */
+
+#pragma omp requires reverse_offload
+
+int x;
+
+void foo (void)
+{
+ #pragma omp target
+ x = 1;
+}
--- /dev/null
+/* { 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" } */
--- /dev/null
+2021-02-02 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_features):
+ New function.
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)
{
projects / thirdparty / gcc.git / commitdiff
