From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 437163857340; Mon, 4 Jul 2022 20:11:06 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 437163857340 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] OpenMP, C++: Add template support for the has_device_addr clause. X-Act-Checkin: gcc X-Git-Author: Marcel Vollweiler X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: edd96cddde899770296a6e9ffc2f19a1c447a2ea X-Git-Newrev: b584b32b1f76a9011d5a2473acc2781c78a0802c Message-Id: <20220704201106.437163857340@sourceware.org> Date: Mon, 4 Jul 2022 20:11:06 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 04 Jul 2022 20:11:06 -0000 https://gcc.gnu.org/g:b584b32b1f76a9011d5a2473acc2781c78a0802c commit b584b32b1f76a9011d5a2473acc2781c78a0802c Author: Marcel Vollweiler Date: Mon Jul 4 21:13:39 2022 +0200 OpenMP, C++: Add template support for the has_device_addr clause. This patch adds support for list items in the has_device_addr clause which type is given by C++ template parameters. gcc/cp/ChangeLog: * pt.cc (tsubst_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR. * semantics.cc (finish_omp_clauses): Added template decl processing. libgomp/ChangeLog: * testsuite/libgomp.c++/target-has-device-addr-7.C: New test. * testsuite/libgomp.c++/target-has-device-addr-8.C: New test. * testsuite/libgomp.c++/target-has-device-addr-9.C: New test. (cherry picked from commit b4fb9f4f9a10d825302cfb5a0ecefa796570d8bc) Diff: --- gcc/cp/ChangeLog.omp | 8 ++++ gcc/cp/pt.cc | 2 + gcc/cp/semantics.cc | 10 ++++- libgomp/ChangeLog.omp | 9 +++++ .../libgomp.c++/target-has-device-addr-7.C | 36 +++++++++++++++++ .../libgomp.c++/target-has-device-addr-8.C | 47 ++++++++++++++++++++++ .../libgomp.c++/target-has-device-addr-9.C | 30 ++++++++++++++ 7 files changed, 140 insertions(+), 2 deletions(-) diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 17d8625e009..5704076ab70 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,11 @@ +2022-07-04 Tobias Burnus + + Backport from mainline: + 2022-05-16 Marcel Vollweiler + + * pt.cc (tsubst_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR. + * semantics.cc (finish_omp_clauses): Added template decl processing. + 2022-06-30 Tobias Burnus * parser.cc (cp_parser_omp_requires): Add missing %<...%> in error. diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index 3aa7c1e0211..f1df95efbc9 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17740,6 +17740,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: OMP_CLAUSE_DECL (nc) @@ -17885,6 +17886,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: + case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: case OMP_CLAUSE_ALLOCATE: diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 8a79def9be7..51918d5091c 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -8759,14 +8759,20 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else { t = OMP_CLAUSE_DECL (c); + while (TREE_CODE (t) == TREE_LIST) + t = TREE_CHAIN (t); while (TREE_CODE (t) == INDIRECT_REF || TREE_CODE (t) == ARRAY_REF) t = TREE_OPERAND (t, 0); } } - bitmap_set_bit (&is_on_device_head, DECL_UID (t)); if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) - cxx_mark_addressable (t); + { + bitmap_set_bit (&is_on_device_head, DECL_UID (t)); + if (!processing_template_decl + && !cxx_mark_addressable (t)) + remove = true; + } goto check_dup_generic_t; case OMP_CLAUSE_USE_DEVICE_ADDR: diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 43b749417c8..d10eb5174df 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,12 @@ +2022-07-04 Tobias Burnus + + Backport from mainline: + 2022-05-16 Marcel Vollweiler + + * testsuite/libgomp.c++/target-has-device-addr-7.C: New test. + * testsuite/libgomp.c++/target-has-device-addr-8.C: New test. + * testsuite/libgomp.c++/target-has-device-addr-9.C: New test. + 2022-07-04 Tobias Burnus Backport from mainline: diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C new file mode 100644 index 00000000000..2c4571be455 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C @@ -0,0 +1,36 @@ +/* Testing 'has_device_addr' clause on the target construct with template. */ + +template +void +foo (T x) +{ + x = 24; + #pragma omp target data map(x) use_device_addr(x) + #pragma omp target has_device_addr(x) + x = 42; + + if (x != 42) + __builtin_abort (); +} + +template +void +bar (T (&x)[]) +{ + x[0] = 24; + #pragma omp target data map(x[:2]) use_device_addr(x) + #pragma omp target has_device_addr(x[:2]) + x[0] = 42; + + if (x[0] != 42) + __builtin_abort (); +} + +int +main () +{ + int a[] = { 24, 42}; + foo (42); + bar (a); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C new file mode 100644 index 00000000000..2adfd30296f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C @@ -0,0 +1,47 @@ +/* Testing 'has_device_addr' clause on the target construct with template. */ + +#include + +template +void +foo (T (&x)[]) +{ + #pragma omp target has_device_addr(x) + for (int i = 0; i < 15; i++) + x[i] = 2 * i; + + #pragma omp target has_device_addr(x[15:15]) + for (int i = 15; i < 30; i++) + x[i] = 3 * i; +} + +int +main () +{ + int *dp = (int*)omp_target_alloc (30*sizeof(int), 0); + + #pragma omp target is_device_ptr(dp) + for (int i = 0; i < 30; i++) + dp[i] = i; + + int (&x)[30] = *static_cast(static_cast(dp)); + + foo (x); + + int y[30]; + for (int i = 0; i < 30; ++i) + y[i] = 0; + int h = omp_get_initial_device (); + int t = omp_get_default_device (); + omp_target_memcpy (&y, dp, 30 * sizeof(int), 0, 0, h, t); + for (int i = 0; i < 15; ++i) + if (y[i] != 2 * i) + __builtin_abort (); + for (int i = 15; i < 30; ++i) + if (y[i] != 3 * i) + __builtin_abort (); + + omp_target_free (dp, 0); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C new file mode 100644 index 00000000000..0c34cab56d2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C @@ -0,0 +1,30 @@ +/* Testing 'has_device_addr' clause on the target construct with template. */ + +#include + +template +void +foo (T (&x)) +{ + #pragma omp target has_device_addr(x) + x = 24; +} + +int +main () +{ + int *dp = (int*)omp_target_alloc (sizeof(int), 0); + int &x = *dp; + + foo (x); + + int y = 42; + int h = omp_get_initial_device (); + int t = omp_get_default_device (); + omp_target_memcpy (&y, dp, sizeof(int), 0, 0, h, t); + if (y != 24) + __builtin_abort (); + + omp_target_free (dp, 0); + return 0; +}