@@ -12090,7 +12090,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
/* Adding the decl for a struct access: we haven't created
GOMP_MAP_STRUCT nodes yet, so this statement needs to predict
- whether they will be created in gimplify_adjust_omp_clauses. */
+ whether they will be created in gimplify_adjust_omp_clauses.
+ NOTE: Technically we should probably look through DECL_VALUE_EXPR
+ here because something that looks like a DECL_P may actually be a
+ struct access, e.g. variables in a lambda closure
+ (__closure->__foo) or class members (this->foo). Currently in both
+ those cases we map the whole of the containing object (directly in
+ the C++ FE) though, so struct nodes are not created. */
if (c == grp_end
&& addr_tokens[0]->type == STRUCTURE_BASE
&& addr_tokens[0]->u.structure_base_kind == BASE_DECL
@@ -13895,6 +13901,18 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
remove = true;
break;
}
+ /* If we have a DECL_VALUE_EXPR (e.g. this is a class member and/or
+ a variable captured in a lambda closure), look through that now
+ before the DECL_P check below. (A code other than COMPONENT_REF,
+ i.e. INDIRECT_REF, will be a VLA/variable-length array
+ section. A global var may be a variable in a common block. We
+ don't want to do this here for either of those.) */
+ if ((ctx->region_type & ORT_ACC) == 0
+ && DECL_P (decl)
+ && !is_global_var (decl)
+ && DECL_HAS_VALUE_EXPR_P (decl)
+ && TREE_CODE (DECL_VALUE_EXPR (decl)) == COMPONENT_REF)
+ decl = OMP_CLAUSE_DECL (c) = DECL_VALUE_EXPR (decl);
if (TREE_CODE (decl) == TARGET_EXPR)
{
if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL,
@@ -855,7 +855,7 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
if (n == NULL)
{
if (allow_zero_length_array_sections)
- cur_node.tgt_offset = 0;
+ cur_node.tgt_offset = cur_node.host_start;
else if (devicep->is_usm_ptr_func
&& devicep->is_usm_ptr_func ((void*)cur_node.host_start))
cur_node.tgt_offset = cur_node.host_start;
@@ -1023,9 +1023,8 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
{
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;
+ copy the host pointer when the target region is not mapped. */
+ data = target;
else
{
gomp_mutex_unlock (&devicep->lock);
@@ -2,6 +2,7 @@
#include <cstdlib>
#include <cstring>
+#include <cstdint>
template <typename L>
void
@@ -22,9 +23,11 @@ struct S
auto fn = [=](void) -> bool
{
bool mapped;
+ uintptr_t hostptr = (uintptr_t) ptr;
+ uintptr_t hostiptr = (uintptr_t) iptr;
#pragma omp target map(from:mapped)
{
- mapped = (ptr != NULL && iptr != NULL);
+ mapped = (ptr != (int*) hostptr && iptr != (int*) hostiptr);
if (mapped)
{
for (int i = 0; i < len; i++)
@@ -2,6 +2,7 @@
#include <stdio.h>
#include <string.h>
+#include <stdint.h>
extern "C" void abort ();
struct S
@@ -15,12 +16,13 @@ struct S
bool set_ptr (int n)
{
bool mapped;
+ uintptr_t hostptr = (uintptr_t) ptr;
#pragma omp target map(from:mapped)
{
- if (ptr != NULL)
+ if (ptr != (int *) hostptr)
for (int i = 0; i < ptr_len; i++)
ptr[i] = n;
- mapped = (ptr != NULL);
+ mapped = (ptr != (int *) hostptr);
}
return mapped;
}
@@ -28,12 +30,13 @@ struct S
bool set_refptr (int n)
{
bool mapped;
+ uintptr_t hostrefptr = (uintptr_t) refptr;
#pragma omp target map(from:mapped)
{
- if (refptr != NULL)
+ if (refptr != (int *) hostrefptr)
for (int i = 0; i < refptr_len; i++)
refptr[i] = n;
- mapped = (refptr != NULL);
+ mapped = (refptr != (int *) hostrefptr);
}
return mapped;
}
@@ -4,6 +4,7 @@
#include <cstdlib>
#include <cstring>
+#include <cstdint>
struct T
{
@@ -18,12 +19,13 @@ struct T
auto fn = [=](void) -> bool
{
bool mapped;
+ uintptr_t hostptr = (uintptr_t) ptr;
#pragma omp target map(from:mapped)
{
- if (ptr)
+ if (ptr != (int *) hostptr)
for (int i = 0; i < ptr_len; i++)
ptr[i] = n;
- mapped = (ptr != NULL);
+ mapped = (ptr != (int *) hostptr);
}
return mapped;
};
@@ -35,12 +37,13 @@ struct T
auto fn = [=](void) -> bool
{
bool mapped;
+ uintptr_t hostrefptr = (uintptr_t) refptr;
#pragma omp target map(from:mapped)
{
- if (refptr)
+ if (refptr != (int *) hostrefptr)
for (int i = 0; i < refptr_len; i++)
refptr[i] = n;
- mapped = (refptr != NULL);
+ mapped = (refptr != (int *) hostrefptr);
}
return mapped;
};