openacc: Set bias to zero for explicit attach/detach clauses in C and C++
This is a fix for the pointer (or array) size inadvertently being used for the bias with attach and detach mapping kinds, for both C and C++. 2020-07-09 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> gcc/c/ PR middle-end/95270 * c-typeck.c (c_finish_omp_clauses): Set OMP_CLAUSE_SIZE (bias) to zero for standalone attach/detach clauses. gcc/cp/ PR middle-end/95270 * semantics.c (finish_omp_clauses): Likewise. include/ PR middle-end/95270 * gomp-constants.h (gomp_map_kind): Expand comment for attach/detach mapping kinds. gcc/testsuite/ PR middle-end/95270 * c-c++-common/goacc/mdc-1.c: Update expected dump output for zero bias. libgomp/ PR middle-end/95270 * testsuite/libgomp.oacc-c-c++-common/pr95270-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr95270-2.c: New test.
This commit is contained in:
parent
8d2e5026d2
commit
0d00fe404c
|
@ -14579,6 +14579,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
}
|
||||
if (c_oacc_check_attachments (c))
|
||||
remove = true;
|
||||
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|
||||
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|
||||
/* In this case, we have a single array element which is a
|
||||
pointer, and we already set OMP_CLAUSE_SIZE in
|
||||
handle_omp_array_sections above. For attach/detach clauses,
|
||||
reset the OMP_CLAUSE_SIZE (representing a bias) to zero
|
||||
here. */
|
||||
OMP_CLAUSE_SIZE (c) = size_zero_node;
|
||||
break;
|
||||
}
|
||||
if (t == error_mark_node)
|
||||
|
@ -14592,6 +14601,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
remove = true;
|
||||
break;
|
||||
}
|
||||
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|
||||
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|
||||
/* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
|
||||
bias) to zero here, so it is not set erroneously to the pointer
|
||||
size later on in gimplify.c. */
|
||||
OMP_CLAUSE_SIZE (c) = size_zero_node;
|
||||
if (TREE_CODE (t) == COMPONENT_REF
|
||||
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
|
||||
{
|
||||
|
|
|
@ -7362,6 +7362,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
}
|
||||
if (cp_oacc_check_attachments (c))
|
||||
remove = true;
|
||||
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|
||||
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|
||||
/* In this case, we have a single array element which is a
|
||||
pointer, and we already set OMP_CLAUSE_SIZE in
|
||||
handle_omp_array_sections above. For attach/detach clauses,
|
||||
reset the OMP_CLAUSE_SIZE (representing a bias) to zero
|
||||
here. */
|
||||
OMP_CLAUSE_SIZE (c) = size_zero_node;
|
||||
break;
|
||||
}
|
||||
if (t == error_mark_node)
|
||||
|
@ -7375,6 +7384,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
remove = true;
|
||||
break;
|
||||
}
|
||||
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|
||||
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|
||||
/* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
|
||||
bias) to zero here, so it is not set erroneously to the pointer
|
||||
size later on in gimplify.c. */
|
||||
OMP_CLAUSE_SIZE (c) = size_zero_node;
|
||||
if (REFERENCE_REF_P (t)
|
||||
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
|
||||
{
|
||||
|
|
|
@ -45,12 +45,12 @@ t1 ()
|
|||
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */
|
||||
|
|
|
@ -139,7 +139,12 @@ enum gomp_map_kind
|
|||
/* Decrement usage count and deallocate if zero. */
|
||||
GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_SPECIAL_2
|
||||
| GOMP_MAP_DELETE),
|
||||
/* In OpenACC, attach a pointer to a mapped struct field. */
|
||||
/* The attach/detach mappings below use the OMP_CLAUSE_SIZE field as a
|
||||
bias. This will typically be zero, except when mapping an array slice
|
||||
with a non-zero base. In that case the bias will indicate the
|
||||
(positive) difference between the start of the actual mapped data and
|
||||
the "virtual" origin of the array.
|
||||
In OpenACC, attach a pointer to a mapped struct field. */
|
||||
GOMP_MAP_ATTACH = (GOMP_MAP_DEEP_COPY | 0),
|
||||
/* In OpenACC, detach a pointer to a mapped struct field. */
|
||||
GOMP_MAP_DETACH = (GOMP_MAP_DEEP_COPY | 1),
|
||||
|
|
|
@ -0,0 +1,46 @@
|
|||
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
|
||||
|
||||
#include <assert.h>
|
||||
#include <openacc.h>
|
||||
#include <stdint.h>
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
int data;
|
||||
int *data_p_dev = (int *) acc_create (&data, sizeof data);
|
||||
int *data_p = &data;
|
||||
uintptr_t ptrbits;
|
||||
|
||||
acc_copyin (&data_p, sizeof data_p);
|
||||
|
||||
/* Test attach/detach directives. */
|
||||
#pragma acc enter data attach(data_p)
|
||||
#pragma acc serial copyout(ptrbits) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
|
||||
{
|
||||
ptrbits = (uintptr_t) data_p;
|
||||
}
|
||||
#pragma acc exit data detach(data_p)
|
||||
assert ((void *) ptrbits == data_p_dev);
|
||||
|
||||
acc_update_self (&data_p, sizeof data_p);
|
||||
assert (data_p == &data);
|
||||
|
||||
/* Test attach/detach API call. */
|
||||
acc_attach ((void **) &data_p);
|
||||
#pragma acc serial copyout(ptrbits) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
|
||||
{
|
||||
ptrbits = (uintptr_t) data_p;
|
||||
}
|
||||
acc_detach ((void **) &data_p);
|
||||
|
||||
assert ((void *) ptrbits == data_p_dev);
|
||||
acc_update_self (&data_p, sizeof data_p);
|
||||
assert (data_p == &data);
|
||||
|
||||
acc_delete (&data_p, sizeof data_p);
|
||||
acc_delete (&data, sizeof data);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,48 @@
|
|||
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
|
||||
|
||||
#include <assert.h>
|
||||
#include <openacc.h>
|
||||
#include <stdint.h>
|
||||
|
||||
#define N 128
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
int *ptrarr[N];
|
||||
int otherarr[N];
|
||||
int sum = 0, hostsum = 0;
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
otherarr[i] = i * 2 + 1;
|
||||
ptrarr[i] = &otherarr[N - 1 - i];
|
||||
hostsum += otherarr[i];
|
||||
}
|
||||
|
||||
acc_copyin (otherarr, sizeof otherarr);
|
||||
acc_copyin (ptrarr, sizeof ptrarr);
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
#pragma acc enter data attach(ptrarr[i])
|
||||
}
|
||||
|
||||
#pragma acc parallel loop copyin(ptrarr[0:N], otherarr[0:N]) \
|
||||
reduction(+:sum)
|
||||
for (int i = 0; i < N; i++)
|
||||
sum += *ptrarr[i];
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
#pragma acc exit data detach(ptrarr[i])
|
||||
}
|
||||
|
||||
assert (sum == hostsum);
|
||||
|
||||
acc_delete (ptrarr, sizeof ptrarr);
|
||||
acc_delete (otherarr, sizeof otherarr);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
Loading…
Reference in New Issue