extern void finish_lambda_function (tree body);
extern bool regenerated_lambda_fn_p (tree);
extern tree most_general_lambda (tree);
-extern tree finish_omp_target (location_t, tree, tree, bool);
-extern void set_omp_target_this_expr (tree);
/* in tree.c */
extern int cp_tree_operand_length (const_tree);
type cast (_expr.cast_ 5.4) to the type of 'this'. [ The cast
ensures that the transformed expression is an rvalue. ] */
result = rvalue (result);
-
- /* Acknowledge to OpenMP target that 'this' was referenced. */
- set_omp_target_this_expr (result);
}
return result;
cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
enum pragma_context context, bool *if_p)
{
+ tree *pc = NULL, stmt;
+
if (flag_openmp)
omp_requires_mask
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
keep_next_level (true);
tree sb = begin_omp_structured_block (), ret;
unsigned save = cp_parser_begin_omp_structured_block (parser);
- set_omp_target_this_expr (NULL_TREE);
switch (ccode)
{
case OMP_TEAMS:
cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
}
}
- finish_omp_target (pragma_tok->location,
- cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true);
- return true;
+ tree stmt = make_node (OMP_TARGET);
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+ OMP_TARGET_BODY (stmt) = body;
+ OMP_TARGET_COMBINED (stmt) = 1;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+ pc = &OMP_TARGET_CLAUSES (stmt);
+ goto check_clauses;
}
else if (!flag_openmp) /* flag_openmp_simd */
{
return false;
}
- tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
- "#pragma omp target", pragma_tok);
- c_omp_adjust_map_clauses (clauses, true);
+ stmt = make_node (OMP_TARGET);
+ TREE_TYPE (stmt) = void_type_node;
+
+ OMP_TARGET_CLAUSES (stmt)
+ = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+ "#pragma omp target", pragma_tok);
+ c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
+ pc = &OMP_TARGET_CLAUSES (stmt);
keep_next_level (true);
- set_omp_target_this_expr (NULL_TREE);
- tree body = cp_parser_omp_structured_block (parser, if_p);
+ OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p);
- finish_omp_target (pragma_tok->location, clauses, body, false);
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+
+check_clauses:
+ while (*pc)
+ {
+ if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (*pc))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
+ break;
+ default:
+ error_at (OMP_CLAUSE_LOCATION (*pc),
+ "%<#pragma omp target%> with map-type other "
+ "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+ "on %<map%> clause");
+ *pc = OMP_CLAUSE_CHAIN (*pc);
+ continue;
+ }
+ pc = &OMP_CLAUSE_CHAIN (*pc);
+ }
return true;
}
static vec<tree> omp_private_member_vec;
static bool omp_private_member_ignore_next;
-/* Used for OpenMP target region 'this' references. */
-static tree omp_target_this_expr = NULL_TREE;
-
-static hash_map<tree, tree> omp_target_ptr_members_accessed;
/* Deferred Access Checking Overview
---------------------------------
finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
{
gcc_assert (TREE_CODE (decl) == FIELD_DECL);
- tree orig_object = object;
bool try_omp_private = !object && omp_private_member_map;
tree ret;
return error_mark_node;
}
- if (orig_object == NULL_TREE)
- {
- tree this_expr = TREE_OPERAND (object, 0);
-
- /* Acknowledge to OpenMP target that 'this' was referenced. */
- set_omp_target_this_expr (this_expr);
- }
-
if (current_class_ptr)
TREE_USED (current_class_ptr) = 1;
if (processing_template_decl)
if (v)
ret = convert_from_reference (*v);
}
- else if (omp_target_this_expr
- && TREE_TYPE (ret)
- && POINTER_TYPE_P (TREE_TYPE (ret)))
- {
- if (omp_target_ptr_members_accessed.get (decl) == NULL)
- omp_target_ptr_members_accessed.put (decl, ret);
- }
-
return ret;
}
}
if (result)
- {
- /* The keyword 'this' is a prvalue expression. */
- result = rvalue (result);
-
- /* Acknowledge to OpenMP target that 'this' was referenced. */
- set_omp_target_this_expr (result);
-
- return result;
- }
+ /* The keyword 'this' is a prvalue expression. */
+ return rvalue (result);
tree fn = current_nonlambda_function ();
if (fn && DECL_STATIC_FUNCTION_P (fn))
{
if (error_operand_p (t))
return error_mark_node;
- if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
- && TREE_CODE (t) == FIELD_DECL)
- t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && TREE_CODE (t) == FIELD_DECL)
+ t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
ret = t;
if (TREE_CODE (t) == COMPONENT_REF
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- if (TREE_CODE (t) == FIELD_DECL)
- t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
else if (TREE_CODE (t) == COMPONENT_REF)
bool order_seen = false;
bool schedule_seen = false;
bool oacc_async = false;
- bool indir_component_ref_p = false;
tree last_iterators = NULL_TREE;
bool last_iterators_remove = false;
/* 1 if normal/task reduction has been seen, -1 if inscan reduction
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
}
- indir_component_ref_p = false;
if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
&& TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
- {
- t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
- indir_component_ref_p = true;
- STRIP_NOPS (t);
- }
+ t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
if (TREE_CODE (t) == COMPONENT_REF
&& ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
|| ort == C_ORT_ACC)
&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
- && !indir_component_ref_p
&& !cxx_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
return add_stmt (stmt);
}
-void
-set_omp_target_this_expr (tree this_val)
-{
- omp_target_this_expr = this_val;
-
- if (omp_target_this_expr == NULL_TREE)
- omp_target_ptr_members_accessed.empty ();
-}
-
-tree
-finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
-{
- tree last_inserted_clause = NULL_TREE;
-
- if (omp_target_this_expr)
- {
- /* See if explicit user-specified map(this[:]) clause already exists.
- If not, we create an implicit map(tofrom:this[:1]) clause. */
- tree *explicit_this_deref_map = NULL;
- for (tree *c = &clauses; *c; c = &OMP_CLAUSE_CHAIN (*c))
- if (OMP_CLAUSE_CODE (*c) == OMP_CLAUSE_MAP
- && TREE_CODE (OMP_CLAUSE_DECL (*c)) == INDIRECT_REF
- && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*c), 0),
- omp_target_this_expr))
- {
- explicit_this_deref_map = c;
- break;
- }
-
- if (DECL_LAMBDA_FUNCTION_P (current_function_decl))
- {
- /* For lambda functions, we need to first create a copy of the
- __closure object. */
- tree closure = DECL_ARGUMENTS (current_function_decl);
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
- OMP_CLAUSE_DECL (c) = build_simple_mem_ref (closure);
- OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure)));
-
- tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
- OMP_CLAUSE_DECL (c2) = closure;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
- OMP_CLAUSE_CHAIN (c2) = clauses;
- OMP_CLAUSE_CHAIN (c) = c2;
- last_inserted_clause = c2;
- clauses = c;
-
- STRIP_NOPS (omp_target_this_expr);
- gcc_assert (DECL_HAS_VALUE_EXPR_P (omp_target_this_expr));
- omp_target_this_expr = DECL_VALUE_EXPR (omp_target_this_expr);
-
- if (explicit_this_deref_map)
- {
- /* Transform *this into *__closure->this in maps. */
- tree this_map = *explicit_this_deref_map;
- OMP_CLAUSE_DECL (this_map)
- = build_simple_mem_ref (omp_target_this_expr);
- tree nc = OMP_CLAUSE_CHAIN (this_map);
- gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (nc)
- == GOMP_MAP_FIRSTPRIVATE_POINTER));
- OMP_CLAUSE_DECL (nc) = omp_target_this_expr;
- OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_POINTER);
-
- /* Move map(*__closure->this) map(always_pointer:__closure->this)
- sequence to right after __closure map. */
- *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc);
- OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c2);
- OMP_CLAUSE_CHAIN (c2) = this_map;
- last_inserted_clause = nc;
- }
- else
- {
- tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_TOFROM);
- OMP_CLAUSE_DECL (c3)
- = build_simple_mem_ref (omp_target_this_expr);
- OMP_CLAUSE_SIZE (c3)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
-
- tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER);
-
- OMP_CLAUSE_DECL (c4) = omp_target_this_expr;
- OMP_CLAUSE_SIZE (c4) = size_zero_node;
-
- OMP_CLAUSE_CHAIN (c3) = c4;
- OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (c2);
- OMP_CLAUSE_CHAIN (c2) = c3;
- last_inserted_clause = c4;
- }
- }
- else
- {
- /* For the non-lambda case, we only need to create map(this[:1]) when
- it's not present, no transforming needed. */
- if (!explicit_this_deref_map)
- {
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
- OMP_CLAUSE_DECL (c) = build_simple_mem_ref (omp_target_this_expr);
- OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
-
- tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
- STRIP_NOPS (omp_target_this_expr);
- OMP_CLAUSE_DECL (c2) = omp_target_this_expr;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
- OMP_CLAUSE_CHAIN (c2) = clauses;
- OMP_CLAUSE_CHAIN (c) = c2;
- clauses = c;
- last_inserted_clause = c2;
- }
- }
- omp_target_this_expr = NULL_TREE;
- }
-
- if (last_inserted_clause && !omp_target_ptr_members_accessed.is_empty ())
- for (hash_map<tree, tree>::iterator i
- = omp_target_ptr_members_accessed.begin ();
- i != omp_target_ptr_members_accessed.end (); ++i)
- {
- /* For each referenced member that is of pointer or reference-to-pointer
- type, create the equivalent of map(alloc:this->ptr[:0]). */
- tree field_decl = (*i).first;
- tree ptr_member = (*i).second;
-
- for (tree nc = OMP_CLAUSE_CHAIN (last_inserted_clause);
- nc != NULL_TREE; nc = OMP_CLAUSE_CHAIN (nc))
- {
- /* If map(this->ptr[:N] already exists, avoid creating another
- such map. */
- tree decl = OMP_CLAUSE_DECL (nc);
- if ((TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == MEM_REF)
- && operand_equal_p (TREE_OPERAND (decl, 0),
- ptr_member))
- goto next_ptr_member;
- }
-
- if (!cxx_mark_addressable (ptr_member))
- gcc_unreachable ();
-
- if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE)
- {
- /* For reference to pointers, we need to map the referenced pointer
- first for things to be correct. */
- tree ptr_member_type = TREE_TYPE (ptr_member);
-
- /* Map pointer target as zero-length array section. */
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
- OMP_CLAUSE_DECL (c)
- = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member);
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
-
- /* Map pointer to zero-length array section. */
- tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND
- (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION);
- OMP_CLAUSE_DECL (c2) = ptr_member;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
-
- /* Attach reference-to-pointer field to pointer. */
- tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH);
- OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0);
- OMP_CLAUSE_SIZE (c3) = size_zero_node;
-
- OMP_CLAUSE_CHAIN (c) = c2;
- OMP_CLAUSE_CHAIN (c2) = c3;
- OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (last_inserted_clause);
-
- OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
- last_inserted_clause = c3;
- }
- else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE)
- {
- /* Map pointer target as zero-length array section. */
- tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
- OMP_CLAUSE_DECL (c)
- = build2 (MEM_REF, char_type_node, ptr_member,
- build_int_cst (build_pointer_type (char_type_node), 0));
- OMP_CLAUSE_SIZE (c) = size_zero_node;
- OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
-
- /* Attach zero-length array section to pointer. */
- tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND
- (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
- OMP_CLAUSE_DECL (c2) = ptr_member;
- OMP_CLAUSE_SIZE (c2) = size_zero_node;
-
- OMP_CLAUSE_CHAIN (c) = c2;
- OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (last_inserted_clause);
- OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
- last_inserted_clause = c2;
- }
- else
- gcc_unreachable ();
-
- next_ptr_member:
- ;
- }
-
- tree stmt = make_node (OMP_TARGET);
- TREE_TYPE (stmt) = void_type_node;
- OMP_TARGET_CLAUSES (stmt) = clauses;
- OMP_TARGET_BODY (stmt) = body;
- OMP_TARGET_COMBINED (stmt) = combined_p;
- SET_EXPR_LOCATION (stmt, loc);
-
- tree c = clauses;
- while (c)
- {
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
- switch (OMP_CLAUSE_MAP_KIND (c))
- {
- case GOMP_MAP_TO:
- case GOMP_MAP_ALWAYS_TO:
- case GOMP_MAP_FROM:
- case GOMP_MAP_ALWAYS_FROM:
- case GOMP_MAP_TOFROM:
- case GOMP_MAP_ALWAYS_TOFROM:
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
- case GOMP_MAP_ALWAYS_POINTER:
- case GOMP_MAP_ATTACH_DETACH:
- case GOMP_MAP_ATTACH:
- case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
- case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
- break;
- default:
- error_at (OMP_CLAUSE_LOCATION (c),
- "%<#pragma omp target%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
- break;
- }
- c = OMP_CLAUSE_CHAIN (c);
- }
- return add_stmt (stmt);
-}
-
tree
finish_omp_parallel (tree clauses, tree body)
{
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH:
case GOMP_MAP_DETACH:
- case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
- case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
+++ /dev/null
-// { dg-do compile }
-// { dg-additional-options "-fdump-tree-gimple" }
-extern "C" void abort ();
-
-struct S
-{
- int a, b, c, d;
-
- int sum (void)
- {
- int val = 0;
- val += a + b + this->c + this->d;
- return val;
- }
-
- int sum_offload (void)
- {
- int val = 0;
- #pragma omp target map(val)
- val += a + b + this->c + this->d;
- return val;
- }
-};
-
-int main (void)
-{
- S s = { 1, 2, 3, 4 };
- if (s.sum () != s.sum_offload ())
- abort ();
- return 0;
-}
-
-/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+++ /dev/null
-// We use 'auto' without a function return type, so specify dialect here
-// { dg-do compile }
-// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
-
-extern "C" void abort ();
-
-struct T
-{
- int x, y;
-
- auto sum_func (int n)
- {
- auto fn = [=](int m) -> int
- {
- int v;
- v = (x + y) * n + m;
- return v;
- };
- return fn;
- }
-
- auto sum_func_offload (int n)
- {
- auto fn = [=](int m) -> int
- {
- int v;
- #pragma omp target map(from:v)
- v = (x + y) * n + m;
- return v;
- };
- return fn;
- }
-
-};
-
-int main (void)
-{
- T a = { 1, 2 };
-
- auto s1 = a.sum_func (3);
- auto s2 = a.sum_func_offload (3);
-
- if (s1 (1) != s2 (1))
- abort ();
-
- return 0;
-}
-
-/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */
+++ /dev/null
-// { dg-do compile }
-// { dg-additional-options "-fdump-tree-gimple" }
-#include <cstdlib>
-#include <cstring>
-extern "C" void abort ();
-
-struct S
-{
- int * ptr;
- int ptr_len;
-
- int *&refptr;
- int refptr_len;
-
- bool set_ptr (int n)
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (ptr != NULL)
- for (int i = 0; i < ptr_len; i++)
- ptr[i] = n;
- mapped = (ptr != NULL);
- }
- return mapped;
- }
-
- bool set_refptr (int n)
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (refptr != NULL)
- for (int i = 0; i < refptr_len; i++)
- refptr[i] = n;
- mapped = (refptr != NULL);
- }
- return mapped;
- }
-};
-
-int main (void)
-{
- #define N 10
- int *ptr1 = new int[N];
- int *ptr2 = new int[N];
-
- memset (ptr1, 0, sizeof (int) * N);
- memset (ptr2, 0, sizeof (int) * N);
-
- S s = { ptr1, N, ptr2, N };
-
- bool mapped;
- int val = 123;
-
- mapped = s.set_ptr (val);
- if (mapped)
- abort ();
- if (s.ptr != ptr1)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr1[i] != 0)
- abort ();
-
- mapped = s.set_refptr (val);
- if (mapped)
- abort ();
- if (s.refptr != ptr2)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr2[i] != 0)
- abort ();
-
- #pragma omp target data map(ptr1[:N])
- mapped = s.set_ptr (val);
-
- if (!mapped)
- abort ();
- if (s.set_refptr (0))
- abort ();
- if (s.ptr != ptr1 || s.refptr != ptr2)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr1[i] != val)
- abort ();
-
- #pragma omp target data map(ptr2[:N])
- mapped = s.set_refptr (val);
-
- if (!mapped)
- abort ();
- if (s.set_ptr (0))
- abort ();
- if (s.ptr != ptr1 || s.refptr != ptr2)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr2[i] != val)
- abort ();
-
- return 0;
-}
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*this->refptr \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
+++ /dev/null
-// We use 'auto' without a function return type, so specify dialect here
-// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
-#include <cstdlib>
-#include <cstring>
-
-struct T
-{
- int *ptr;
- int ptr_len;
-
- int *&refptr;
- int refptr_len;
-
- auto set_ptr_func (int n)
- {
- auto fn = [=](void) -> bool
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (ptr)
- for (int i = 0; i < ptr_len; i++)
- ptr[i] = n;
- mapped = (ptr != NULL);
- }
- return mapped;
- };
- return fn;
- }
-
- auto set_refptr_func (int n)
- {
- auto fn = [=](void) -> bool
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (refptr)
- for (int i = 0; i < refptr_len; i++)
- refptr[i] = n;
- mapped = (refptr != NULL);
- }
- return mapped;
- };
- return fn;
- }
-};
-
-int main (void)
-{
- #define N 10
- int *ptr1 = new int[N];
- int *ptr2 = new int[N];
-
- memset (ptr1, 0, sizeof (int) * N);
- memset (ptr2, 0, sizeof (int) * N);
-
- T a = { ptr1, N, ptr2, N };
-
- auto p1 = a.set_ptr_func (1);
- auto r2 = a.set_refptr_func (2);
-
- if (p1 ())
- abort ();
- if (r2 ())
- abort ();
-
- if (a.ptr != ptr1)
- abort ();
- if (a.refptr != ptr2)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr1[i] != 0)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr2[i] != 0)
- abort ();
-
- #pragma omp target data map(ptr1[:N], ptr2[:N])
- {
- if (!p1 ())
- abort ();
- if (!r2 ())
- abort ();
- }
-
- if (a.ptr != ptr1)
- abort ();
- if (a.refptr != ptr2)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr1[i] != 1)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr2[i] != 2)
- abort ();
-
- return 0;
-}
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+->refptr \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_3->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_POINTER:
- case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, "alloc");
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
pp_string (pp, "force_present,noncontig_array");
break;
- case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
- pp_string (pp, "attach_zero_length_array_section");
- break;
default:
gcc_unreachable ();
}
case GOMP_MAP_ALWAYS_POINTER:
pp_string (pp, " [pointer assign, bias: ");
break;
- case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
- pp_string (pp, " [pointer assign, zero-length array section, bias: ");
- break;
case GOMP_MAP_TO_PSET:
pp_string (pp, " [pointer set, len: ");
break;
case GOMP_MAP_DETACH:
case GOMP_MAP_FORCE_DETACH:
case GOMP_MAP_ATTACH_DETACH:
- case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, " [bias: ");
break;
default:
No refcount is bumped by this, and the store is done unconditionally. */
GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_FLAG_SPECIAL | 1),
- /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to
- NULL if target is not mapped. */
- GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION
- = (GOMP_MAP_FLAG_SPECIAL_2
- | GOMP_MAP_FLAG_SPECIAL | 2),
/* Forced deallocation of zero length array section. */
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
= (GOMP_MAP_FLAG_SPECIAL_2
GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT = (GOMP_MAP_NONCONTIG_ARRAY
| GOMP_MAP_FORCE_PRESENT),
- /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections
- (i.e. set to NULL when array section is not mapped) Currently only used
- by OpenMP. */
- GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
- = (GOMP_MAP_DEEP_COPY | 2),
-
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1),
((X) == GOMP_MAP_ALWAYS_POINTER)
#define GOMP_MAP_POINTER_P(X) \
- ((X) == GOMP_MAP_POINTER \
- || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
+ ((X) == GOMP_MAP_POINTER)
#define GOMP_MAP_ALWAYS_TO_P(X) \
(((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
extern void gomp_attach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree,
splay_tree_key, uintptr_t, size_t,
- struct gomp_coalesce_buf *, bool);
+ struct gomp_coalesce_buf *);
extern void gomp_detach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree_key,
uintptr_t, bool, struct gomp_coalesce_buf *);
}
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
- 0, NULL, false);
+ 0, NULL);
gomp_mutex_unlock (&acc_dev->lock);
}
if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
{
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
- (uintptr_t) h, s, NULL, false);
+ (uintptr_t) h, s, NULL);
/* OpenACC 'attach'/'detach' doesn't affect structured/dynamic
reference counts ('n->refcount', 'n->dynamic_refcount'). */
}
splay_tree_key m
= lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
- (uintptr_t) hostaddrs[j], sizes[j], NULL,
- false);
+ (uintptr_t) hostaddrs[j], sizes[j], NULL);
}
bool processed = false;
struct gomp_coalesce_buf *cbuf,
htab_t *refcount_set)
{
- assert (kind != GOMP_MAP_ATTACH
- || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
+ assert (kind != GOMP_MAP_ATTACH);
tgt_var->key = oldn;
tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
static void
gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias,
- struct gomp_coalesce_buf *cbuf,
- bool allow_zero_length_array_sections)
+ struct gomp_coalesce_buf *cbuf)
{
struct gomp_device_descr *devicep = tgt->device_descr;
struct splay_tree_s *mem_map = &devicep->mem_map;
splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
if (n == NULL)
{
- if (allow_zero_length_array_sections)
- cur_node.tgt_offset = 0;
- else
- {
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("Pointer target of array section wasn't mapped");
- }
- }
- else
- {
- cur_node.host_start -= n->host_start;
- cur_node.tgt_offset
- = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
- /* At this point tgt_offset is target address of the
- array section. Now subtract bias to get what we want
- to initialize the pointer with. */
- cur_node.tgt_offset -= bias;
- }
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("Pointer target of array section wasn't mapped");
+ }
+ cur_node.host_start -= n->host_start;
+ cur_node.tgt_offset
+ = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
+ /* At this point tgt_offset is target address of the
+ array section. Now subtract bias to get what we want
+ to initialize the pointer with. */
+ cur_node.tgt_offset -= bias;
gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset),
(void *) &cur_node.tgt_offset, sizeof (void *), true,
cbuf);
gomp_attach_pointer (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, splay_tree mem_map,
splay_tree_key n, uintptr_t attach_to, size_t bias,
- struct gomp_coalesce_buf *cbufp,
- bool allow_zero_length_array_sections)
+ struct gomp_coalesce_buf *cbufp)
{
struct splay_tree_key_s s;
size_t size, idx;
if (!tn)
{
- if (allow_zero_length_array_sections)
- {
- /* When allowing attachment to zero-length array sections, we
- allow attaching to NULL pointers when the target region is not
- mapped. */
- data = 0;
- }
- else
- {
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("pointer target not mapped for attach");
- }
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("pointer target not mapped for attach");
}
- else
- data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
+
+ data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
gomp_debug (1,
"%s: attaching host %p, target %p (struct base %p) to %p\n",
has_firstprivate = true;
continue;
}
- else if ((kind & typemask) == GOMP_MAP_ATTACH
- || ((kind & typemask)
- == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
+ else if ((kind & typemask) == GOMP_MAP_ATTACH)
{
tgt->list[i].key = NULL;
has_firstprivate = true;
(uintptr_t) *(void **) hostaddrs[j],
k->tgt_offset + ((uintptr_t) hostaddrs[j]
- k->host_start),
- sizes[j], cbufp, false);
+ sizes[j], cbufp);
}
}
i = j - 1;
++i;
continue;
case GOMP_MAP_ATTACH:
- case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
{
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start + sizeof (void *);
structured/dynamic reference counts ('n->refcount',
'n->dynamic_refcount'). */
- bool zlas
- = ((kind & typemask)
- == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
gomp_attach_pointer (devicep, aq, mem_map, n,
(uintptr_t) hostaddrs[i], sizes[i],
- cbufp, zlas);
+ cbufp);
}
else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
{
cbufp);
break;
case GOMP_MAP_POINTER:
- case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
- gomp_map_pointer
- (tgt, aq, (uintptr_t) *(void **) k->host_start,
- k->tgt_offset, sizes[i], cbufp,
- ((kind & typemask)
- == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION));
+ gomp_map_pointer (tgt, aq,
+ (uintptr_t) *(void **) k->host_start,
+ k->tgt_offset, sizes[i], cbufp);
break;
case GOMP_MAP_TO_PSET:
gomp_copy_host2dev (devicep, aq,
k->tgt_offset
+ ((uintptr_t) hostaddrs[j]
- k->host_start),
- sizes[j], cbufp, false);
+ sizes[j], cbufp);
}
}
i = j - 1;
+++ /dev/null
-extern "C" void abort ();
-
-struct S
-{
- int a, b, c, d;
-
- int sum (void)
- {
- int val = 0;
- val += a + b + this->c + this->d;
- return val;
- }
-
- int sum_offload (void)
- {
- int val = 0;
- #pragma omp target map(val)
- val += a + b + this->c + this->d;
- return val;
- }
-};
-
-int main (void)
-{
- S s = { 1, 2, 3, 4 };
- if (s.sum () != s.sum_offload ())
- abort ();
- return 0;
-}
+++ /dev/null
-
-// We use 'auto' without a function return type, so specify dialect here
-// { dg-additional-options "-std=c++14" }
-
-extern "C" void abort ();
-
-struct T
-{
- int x, y;
-
- auto sum_func (int n)
- {
- auto fn = [=](int m) -> int
- {
- int v;
- v = (x + y) * n + m;
- return v;
- };
- return fn;
- }
-
- auto sum_func_offload (int n)
- {
- auto fn = [=](int m) -> int
- {
- int v;
- #pragma omp target map(from:v)
- v = (x + y) * n + m;
- return v;
- };
- return fn;
- }
-
-};
-
-int main (void)
-{
- T a = { 1, 2 };
-
- auto s1 = a.sum_func (3);
- auto s2 = a.sum_func_offload (3);
-
- if (s1 (1) != s2 (1))
- abort ();
-
- return 0;
-}
+++ /dev/null
-#include <stdio.h>
-#include <string.h>
-extern "C" void abort ();
-
-struct S
-{
- int * ptr;
- int ptr_len;
-
- int *&refptr;
- int refptr_len;
-
- bool set_ptr (int n)
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (ptr != NULL)
- for (int i = 0; i < ptr_len; i++)
- ptr[i] = n;
- mapped = (ptr != NULL);
- }
- return mapped;
- }
-
- bool set_refptr (int n)
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (refptr != NULL)
- for (int i = 0; i < refptr_len; i++)
- refptr[i] = n;
- mapped = (refptr != NULL);
- }
- return mapped;
- }
-};
-
-int main (void)
-{
- #define N 10
- int *ptr1 = new int[N];
- int *ptr2 = new int[N];
-
- memset (ptr1, 0, sizeof (int) * N);
- memset (ptr2, 0, sizeof (int) * N);
-
- S s = { ptr1, N, ptr2, N };
-
- bool mapped;
- int val = 123;
-
- mapped = s.set_ptr (val);
- if (mapped)
- abort ();
- if (s.ptr != ptr1)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr1[i] != 0)
- abort ();
-
- mapped = s.set_refptr (val);
- if (mapped)
- abort ();
- if (s.refptr != ptr2)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr2[i] != 0)
- abort ();
-
- #pragma omp target data map(ptr1[:N])
- mapped = s.set_ptr (val);
-
- if (!mapped)
- abort ();
- if (s.set_refptr (0))
- abort ();
- if (s.ptr != ptr1 || s.refptr != ptr2)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr1[i] != val)
- abort ();
-
- #pragma omp target data map(ptr2[:N])
- mapped = s.set_refptr (val);
-
- if (!mapped)
- abort ();
- if (s.set_ptr (0))
- abort ();
- if (s.ptr != ptr1 || s.refptr != ptr2)
- abort ();
- for (int i = 0; i < N; i++)
- if (ptr2[i] != val)
- abort ();
-
- return 0;
-}
+++ /dev/null
-
-// We use 'auto' without a function return type, so specify dialect here
-// { dg-additional-options "-std=c++14" }
-#include <cstdlib>
-#include <cstring>
-
-struct T
-{
- int *ptr;
- int ptr_len;
-
- int *&refptr;
- int refptr_len;
-
- auto set_ptr_func (int n)
- {
- auto fn = [=](void) -> bool
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (ptr)
- for (int i = 0; i < ptr_len; i++)
- ptr[i] = n;
- mapped = (ptr != NULL);
- }
- return mapped;
- };
- return fn;
- }
-
- auto set_refptr_func (int n)
- {
- auto fn = [=](void) -> bool
- {
- bool mapped;
- #pragma omp target map(from:mapped)
- {
- if (refptr)
- for (int i = 0; i < refptr_len; i++)
- refptr[i] = n;
- mapped = (refptr != NULL);
- }
- return mapped;
- };
- return fn;
- }
-};
-
-int main (void)
-{
- #define N 10
- int *ptr1 = new int[N];
- int *ptr2 = new int[N];
-
- memset (ptr1, 0, sizeof (int) * N);
- memset (ptr2, 0, sizeof (int) * N);
-
- T a = { ptr1, N, ptr2, N };
-
- auto p1 = a.set_ptr_func (1);
- auto r2 = a.set_refptr_func (2);
-
- if (p1 ())
- abort ();
- if (r2 ())
- abort ();
-
- if (a.ptr != ptr1)
- abort ();
- if (a.refptr != ptr2)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr1[i] != 0)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr2[i] != 0)
- abort ();
-
- #pragma omp target data map(ptr1[:N], ptr2[:N])
- {
- if (!p1 ())
- abort ();
- if (!r2 ())
- abort ();
- }
-
- if (a.ptr != ptr1)
- abort ();
- if (a.refptr != ptr2)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr1[i] != 1)
- abort ();
-
- for (int i = 0; i < N; i++)
- if (ptr2[i] != 2)
- abort ();
-
- return 0;
-}
projects / thirdparty / gcc.git / commitdiff
