@@ -451,6 +451,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONG_ULONG_ULONGPTR, BT_BOOL, BT_ULONG,
DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONGLONG_ULONGLONG_ULONGLONGPTR, BT_BOOL,
BT_ULONGLONG, BT_ULONGLONG, BT_PTR_ULONGLONG)
DEF_FUNCTION_TYPE_3 (BT_FN_INT_INT_INT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_UINT, BT_VOID, BT_PTR, BT_INT, BT_UINT)
DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR,
BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR)
@@ -1365,6 +1365,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs,
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
kind = " oacc_enter_exit_data";
break;
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ kind = " oacc_declare";
+ break;
default:
gcc_unreachable ();
}
@@ -100,7 +100,7 @@ enum gf_mask {
GF_OMP_FOR_KIND_CILKSIMD = GF_OMP_FOR_SIMD | 1,
GF_OMP_FOR_COMBINED = 1 << 3,
GF_OMP_FOR_COMBINED_INTO = 1 << 4,
- GF_OMP_TARGET_KIND_MASK = (1 << 3) - 1,
+ GF_OMP_TARGET_KIND_MASK = (1 << 4) - 1,
GF_OMP_TARGET_KIND_REGION = 0,
GF_OMP_TARGET_KIND_DATA = 1,
GF_OMP_TARGET_KIND_UPDATE = 2,
@@ -109,6 +109,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_OACC_DATA = 5,
GF_OMP_TARGET_KIND_OACC_UPDATE = 6,
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 7,
+ GF_OMP_TARGET_KIND_OACC_DECLARE = 8,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -5663,6 +5664,7 @@ is_gimple_omp_oacc (const_gimple stmt)
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
return true;
default:
return false;
@@ -5819,10 +5819,26 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
splay_tree_node n;
unsigned flags = in_code ? GOVD_SEEN : 0;
bool ret = false, shared;
+ bool device_resident = false;
if (error_operand_p (decl))
return false;
+ if (flag_openacc && is_global_var (decl))
+ {
+ tree attr = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ if (attr)
+ {
+ tree t, c;
+ for (t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t))
+ {
+ c = TREE_VALUE (t);
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT)
+ device_resident = true;
+ }
+ }
+ }
+
/* Threadprivate variables are predetermined. */
if (is_global_var (decl))
{
@@ -5899,7 +5915,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
by default are firstprivate (gang-local) in parallel. */
if (!n2 && !AGGREGATE_TYPE_P (type))
{
- if (ctx->acc_region_kind == ARK_PARALLEL)
+ if (device_resident)
+ flags |= GOVD_MAP_TO_ONLY;
+ else if (ctx->acc_region_kind == ARK_PARALLEL)
flags |= (GOVD_GANGLOCAL | GOVD_MAP_TO_ONLY);
/* Scalars under kernels are default 'copy'. */
else if (ctx->acc_region_kind == ARK_KERNELS)
@@ -7729,6 +7747,10 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
+ case OACC_DECLARE:
+ kind = GF_OMP_TARGET_KIND_OACC_DECLARE;
+ ork = ORK_OACC;
+ break;
case OACC_ENTER_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
ork = ORK_OACC;
@@ -8707,11 +8729,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = gimplify_oacc_host_data (expr_p, pre_p);
break;
- case OACC_DECLARE:
- sorry ("directive not yet implemented");
- ret = GS_ALL_DONE;
- break;
-
case OACC_KERNELS:
case OACC_PARALLEL:
case OACC_DATA:
@@ -8724,6 +8741,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_DECLARE:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_UPDATE:
@@ -299,3 +299,7 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update",
BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_STATIC, "GOACC_register_static",
+ BT_FN_VOID_PTR_INT_UINT, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DECLARE, "GOACC_declare",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
@@ -9519,6 +9519,7 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
data_region = false;
break;
case GF_OMP_TARGET_KIND_DATA:
@@ -9825,6 +9826,9 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
break;
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ start_ix = BUILT_IN_GOACC_DECLARE;
+ break;
default:
gcc_unreachable ();
}
@@ -9944,6 +9948,7 @@ expand_omp_target (struct omp_region *region)
args.quick_push (build_zero_cst (ptr_type_node));
break;
case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOACC_DECLARE:
case BUILT_IN_GOACC_ENTER_EXIT_DATA:
case BUILT_IN_GOACC_KERNELS:
case BUILT_IN_GOACC_KERNELS_INTERNAL:
@@ -9960,6 +9965,7 @@ expand_omp_target (struct omp_region *region)
switch (start_ix)
{
case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOACC_DECLARE:
case BUILT_IN_GOMP_TARGET:
case BUILT_IN_GOMP_TARGET_DATA:
case BUILT_IN_GOMP_TARGET_UPDATE:
@@ -10268,6 +10274,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
/* ..., other than for those stand-alone directives... */
region = NULL;
break;
@@ -12771,6 +12778,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
data_region = false;
break;
case GF_OMP_TARGET_KIND_DATA:
@@ -12835,6 +12843,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FORCE_PRESENT:
case GOMP_MAP_FORCE_DEALLOC:
case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_LINK:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
@@ -13888,6 +13898,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
cur_region = cur_region->outer;
break;
default:
@@ -1,3 +1,10 @@
+
+2015-06-04 James Norris <jnorris@codesourcery.com>
+
+ * c-c++-common/goacc/declare-1.c: Update tests.
+ * c-c++-common/goacc/declare-2.c: Likewise.
+ * gfortran.dg/goacc/declare-1.f95: Update tests.
+
2015-06-01 Tom de Vries <tom@codesourcery.com>
Revert:
@@ -1,6 +1,5 @@
/* Test valid uses of declare directive. */
/* { dg-do compile } */
-/* { dg-skip-if "not yet" { c++ } } */
int v0;
#pragma acc declare create(v0)
@@ -1,11 +1,10 @@
/* Test invalid uses of declare directive. */
/* { dg-do compile } */
-/* { dg-skip-if "not yet" { c++ } } */
#pragma acc declare /* { dg-error "no valid clauses" } */
#pragma acc declare create(undeclared) /* { dg-error "undeclared" } */
-/* { dg-error "no valid clauses" "second error" { target *-*-* } 7 } */
+/* { dg-error "no valid clauses" "second error" { target *-*-* } 6 } */
int v0[10];
#pragma acc declare create(v0[1:3]) /* { dg-error "subarray" } */
@@ -42,7 +41,7 @@ void
f (void)
{
int va0;
-#pragma acc declare link(va0) /* { dg-error "invalid variable" } */
+#pragma acc declare link(va0) /* { dg-error "global variable" } */
extern int ve0;
#pragma acc declare copy(ve0) /* { dg-error "invalid use of" } */
@@ -15,5 +15,6 @@ contains
END BLOCK
end function foo
end program test
-! { dg-final { scan-tree-dump-times "pragma acc data map\\(force_tofrom:i\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_to:i\\)" 2 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_from:i\\)" 2 "original" } }
! { dg-final { cleanup-tree-dump "original" } }
@@ -57,6 +57,7 @@ along with GCC; see the file COPYING3. If not see
#include "lto-streamer.h"
#include "context.h"
#include "omp-low.h"
+#include "gomp-constants.h"
const char * const tls_model_names[]={"none", "emulated",
"global-dynamic", "local-dynamic",
@@ -161,6 +162,58 @@ varpool_node::create_empty (void)
return node;
}
+static void
+make_offloadable_1 (varpool_node *node, tree decl)
+{
+ node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+ g->have_offload = true;
+ if (!in_lto_p)
+ vec_safe_push (offload_vars, decl);
+ node->force_output = 1;
+#endif
+}
+
+void
+make_offloadable (varpool_node *node, tree decl)
+{
+ tree attrs;
+
+ if (node->offloadable)
+ return;
+
+ if (flag_openmp)
+ {
+ make_offloadable_1 (node, decl);
+ return;
+ }
+
+ attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ if (attrs)
+ {
+ tree *t;
+ int total = 0, skip = 0;
+
+ gcc_assert (&TREE_VALUE (attrs));
+
+ for (t = &TREE_VALUE (attrs); *t; t = &TREE_CHAIN (*t))
+ {
+ HOST_WIDE_INT kind = OMP_CLAUSE_MAP_KIND (TREE_VALUE (*t));
+
+ total++;
+
+ if (kind == GOMP_MAP_LINK)
+ skip++;
+ }
+
+ if (total - skip > 0)
+ make_offloadable_1 (node, decl);
+
+ DECL_ATTRIBUTES (decl)
+ = remove_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ }
+}
+
/* Return varpool node assigned to DECL. Create new one when needed. */
varpool_node *
varpool_node::get_create (tree decl)
@@ -168,22 +221,18 @@ varpool_node::get_create (tree decl)
varpool_node *node = varpool_node::get (decl);
gcc_checking_assert (TREE_CODE (decl) == VAR_DECL);
if (node)
- return node;
+ {
+ if (flag_openacc && !DECL_EXTERNAL (decl))
+ make_offloadable (node, decl);
+ return node;
+ }
node = varpool_node::create_empty ();
node->decl = decl;
if ((flag_openacc || flag_openmp) && !DECL_EXTERNAL (decl)
&& lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)))
- {
- node->offloadable = 1;
-#ifdef ENABLE_OFFLOADING
- g->have_offload = true;
- if (!in_lto_p)
- vec_safe_push (offload_vars, decl);
- node->force_output = 1;
-#endif
- }
+ make_offloadable (node, decl);
node->register_symbol ();
return node;