From: Julian Brown Date: Tue, 9 Jun 2020 13:21:34 +0000 (-0700) Subject: openacc: Set bias to zero for explicit attach/detach clauses in C and C++ X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=0d00fe404c162ad0cf922ca8455aa23a74042b63;p=gcc.git 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 Thomas Schwinge 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. --- diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index b28c2c5ff62..fb5c288b549 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -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_) { diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index d63cea96e23..4a3ef3d2839 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -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) { diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c index fb5841a709d..337c1f7cc77 100644 --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -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" } } */ diff --git a/include/gomp-constants.h b/include/gomp-constants.h index b42b41403aa..7e44238ae03 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -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), diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-1.c new file mode 100644 index 00000000000..0457c232bc9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-1.c @@ -0,0 +1,46 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +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; +} + diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-2.c new file mode 100644 index 00000000000..0575e726738 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr95270-2.c @@ -0,0 +1,48 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include +#include +#include + +#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; +} +