Loosely related to PR 66053 which is about explicit mapping. The original code boils down to: this->j = this->i[1] + this->i[2] This ends up in omp_notice_variable (via gimplify_var_or_parm_decl) as "this", but it is – in line with OpenMP 4.5 – only mapped as: map(alloc:MEM[(char *)this] [len: 0]) map(firstprivate:this [pointer assign, bias: 0]) However, for OpenMP 5 – and to be usable – one also needs the actual data, i.e. map(tofrom:this[:1]). OpenMP 4.5, "2.15.5 Data-mapping Attribute Rules and Clauses": "A variable that is of type reference to pointer is treated as if it had appeared in a map clause as a1zero-length array section." OpenMP 5.0: "If the target construct is within a class non-static member function, and a variable is an accessible data member of the object for which the non-static data member function is invoked, the variable is treated as if the this[:1] expression had appeared in a map clause with a map-type of tofrom. […]" #include <stdlib.h> // For abort. class foo { public: int i[2], j; void func() { int k[2]; #pragma omp target j = i[1]+i[2]; } }; int bar(int x) { foo cl; cl.i[1] = x; cl.i[2] = 42; cl.func(); return cl.j; } int main() { if (bar(3) != 45) abort(); return 0; }
[Tacking bug, just to not to forget it. – It is not completely clear to me whether it shall be handled in gimplify.c or in the front end.]
For OpenMP 5 there is quite a lot of mapping related changes that need to be implemented, first of all, the C/C++ array section support needs to be rewritten because the rules changed there too much and the current limited parsing is not usable, this[:1] needs to be handled (this PR), declare mapper support added, non-contiguous arrays in to/from, etc. This all is deferred for GCC 11.
The master branch has been updated by Chung-Lin Tang <cltang@gcc.gnu.org>: https://gcc.gnu.org/g:0ab29cf0bb68960c1f87405f14b4fb2109254e2f commit r12-5835-g0ab29cf0bb68960c1f87405f14b4fb2109254e2f Author: Chung-Lin Tang <cltang@codesourcery.com> Date: Wed Dec 8 22:28:03 2021 +0800 openmp: Improve OpenMP target support for C++ (PR92120) This patch implements several C++ specific mapping capabilities introduced for OpenMP 5.0, including implicit mapping of this[:1] for non-static member functions, zero-length array section mapping of pointer-typed members, lambda captured variable access in target regions, and use of lambda objects inside target regions. Several adjustments to the C/C++ front-ends to allow more member-access syntax as valid is also included. PR middle-end/92120 gcc/cp/ChangeLog: * cp-tree.h (finish_omp_target): New declaration. (finish_omp_target_clauses): Likewise. * parser.c (cp_parser_omp_clause_map): Adjust call to cp_parser_omp_var_list_no_open to set 'allow_deref' argument to true. (cp_parser_omp_target): Factor out code, adjust into calls to new function finish_omp_target. * pt.c (tsubst_expr): Add call to finish_omp_target_clauses for OMP_TARGET case. * semantics.c (handle_omp_array_sections_1): Add handling to create 'this->member' from 'member' FIELD_DECL. Remove case of rejecting 'this' when not in declare simd. (handle_omp_array_sections): Likewise. (finish_omp_clauses): Likewise. Adjust to allow 'this[]' in OpenMP map clauses. Handle 'A->member' case in map clauses. Remove case of rejecting 'this' when not in declare simd. (struct omp_target_walk_data): New struct for walking over target-directive tree body. (finish_omp_target_clauses_r): New function for tree walk. (finish_omp_target_clauses): New function. (finish_omp_target): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in call to c_parser_omp_variable_list to 'true'. * c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in array base handling. (c_finish_omp_clauses): Handle 'A->member' case in map clauses. gcc/ChangeLog: * gimplify.c ("tree-hash-traits.h"): Add include. (gimplify_scan_omp_clauses): Change struct_map_to_clause to type hash_map<tree_operand, tree> *. Adjust struct map handling to handle cases of *A and A->B expressions. Under !DECL_P case of GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to struct_deref_set for map(*ptr_to_struct) cases. Add MEM_REF case when handling component_ref_p case. Add unshare_expr and gimplification when created GOMP_MAP_STRUCT is not a DECL. Add code to add firstprivate pointer for *pointer-to-struct case. (gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for exit data directives code to earlier position. * omp-low.c (lower_omp_target): Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. * tree-pretty-print.c (dump_omp_clause): Likewise. gcc/testsuite/ChangeLog: * gcc.dg/gomp/target-3.c: New testcase. * g++.dg/gomp/target-3.C: New testcase. * g++.dg/gomp/target-lambda-1.C: New testcase. * g++.dg/gomp/target-lambda-2.C: New testcase. * g++.dg/gomp/target-this-1.C: New testcase. * g++.dg/gomp/target-this-2.C: New testcase. * g++.dg/gomp/target-this-3.C: New testcase. * g++.dg/gomp/target-this-4.C: New testcase. * g++.dg/gomp/target-this-5.C: New testcase. * g++.dg/gomp/this-2.C: Adjust testcase. include/ChangeLog: * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. (GOMP_MAP_POINTER_P): Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION. libgomp/ChangeLog: * libgomp.h (gomp_attach_pointer): Add bool parameter. * oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_existing): Update assert condition to include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION. (gomp_map_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for mapping a pointer with NULL target. (gomp_attach_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for attaching a pointer with NULL target. (gomp_map_vars_internal): Update calls to gomp_map_pointer and gomp_attach_pointer, add handling for GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases. * testsuite/libgomp.c++/target-23.C: New testcase. * testsuite/libgomp.c++/target-lambda-1.C: New testcase. * testsuite/libgomp.c++/target-lambda-2.C: New testcase. * testsuite/libgomp.c++/target-this-1.C: New testcase. * testsuite/libgomp.c++/target-this-2.C: New testcase. * testsuite/libgomp.c++/target-this-3.C: New testcase. * testsuite/libgomp.c++/target-this-4.C: New testcase. * testsuite/libgomp.c++/target-this-5.C: New testcase.
Without offloading, a few of the new test cases FAIL; please adjust: +PASS: libgomp.c++/target-lambda-1.C (test for excess errors) +FAIL: libgomp.c++/target-lambda-1.C execution test +PASS: libgomp.c++/target-this-3.C (test for excess errors) +FAIL: libgomp.c++/target-this-3.C execution test +PASS: libgomp.c++/target-this-4.C (test for excess errors) +FAIL: libgomp.c++/target-this-4.C execution test
Fixed.