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 NULL_TREE;
}
+ 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 NULL_TREE;
}
+ 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)
+ if (this_req == OMP_REQUIRES_DYNAMIC_ALLOCATORS)
sorry_at (cloc, "%qs clause on %<requires%> directive not "
"supported yet", p);
if (p)
/* 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 ();
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 NULL_TREE;
}
+ 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 NULL_TREE;
}
+ 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)
+ if (this_req == OMP_REQUIRES_DYNAMIC_ALLOCATORS)
sorry_at (cloc, "%qs clause on %<requires%> directive not "
"supported yet", p);
if (p)
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)
+ /* Currently, everything except 'dynamic_allocators' is allowed. */
+ if (requires_clause == 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))
#include "system.h"
#include "coretypes.h"
#include "options.h"
+#include "tree.h"
#include "gfortran.h"
#include <setjmp.h>
#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. */
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);
+ }
+
/* Do the parse tree dump. */
gfc_current_ns = flag_dump_fortran_original ? gfc_global_ns_list : NULL;
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
{
#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
#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. */
__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
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" } */
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
