https://gcc.gnu.org/g:9d1f31a22fa5d9714c7e03bfdf716e3c8e38c3e5
commit r17-5499-g9d1f31a22fa5d9714c7e03bfdf716e3c8e38c3e5 Author: Thomas Schwinge <[email protected]> Date: Thu Oct 22 12:51:45 2026 +0200 openmp c++: Temporary fix for some lambda mapping problems: More testing, XFAILed regressions [PR127559] Since commit 84d5fecd2694ad70f391a20451d43d82fde37eb5 "openmp c++: Temporary fix for some lambda mapping problems [PR127559]", the C++ front end no longer maps lambda objects that are subobjects of aggregates, such as 'c.x' or 'd[y].x' in 'libgomp.c++/target-lambda-6.C', or their by-reference captures. 'gcc/cp/semantics.cc:finish_omp_target_clauses_r' looks through 'COMPONENT_REF's, 'ARRAY_REF's for a 'VAR_DECL', 'PARM_DECL', 'RESULT_DECL' base, but the base is wrapped in a 'VIEW_CONVERT_EXPR' (location wrapper), and therefore not recognized. Without the mapping for the by-reference capture of 'x', code offloading execution dereferences a host stack address, which faults in case that the device can't access host memory. Thus, XFAILing the 'libgomp.c++/target-lambda-6.C' execution test accordingly, and adding XFAILed 'scan-tree-dump's to it. The new 'g++.dg/gomp/target-lambda-3.C' also covers a nested aggregate, and a constant array index. With the commit mentioned above reverted, these all PASS, so this is a regression. Also adding a 'scan-tree-dump' to 'libgomp.c++/target-lambda-5.C' to verify that locally declared lambdas don't get mapped. Assisted-by: Anthropic Claude Opus 5.5 PR c++/127559 gcc/testsuite/ * g++.dg/gomp/target-lambda-3.C: New. libgomp/ * testsuite/libgomp.c++/target-lambda-5.C: Enhance testing. * testsuite/libgomp.c++/target-lambda-6.C: Likewise. Diff: --- gcc/testsuite/g++.dg/gomp/target-lambda-3.C | 63 +++++++++++++++++++++++++ libgomp/testsuite/libgomp.c++/target-lambda-5.C | 4 ++ libgomp/testsuite/libgomp.c++/target-lambda-6.C | 10 ++++ 3 files changed, 77 insertions(+) diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-3.C b/gcc/testsuite/g++.dg/gomp/target-lambda-3.C new file mode 100644 index 000000000000..163cc0e05095 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-3.C @@ -0,0 +1,63 @@ +/* PR c++/127559 + { dg-do compile { target c++11 } } + { dg-additional-options -fdump-tree-original } */ + +/* Lambda objects that are subobjects of aggregates must be mapped, including + their by-reference captures. */ + +template <typename T> struct S { T x; int y; }; +template <typename T> int baz (T &x) { return x.x (); } + +void +f1 (int x) +{ + auto b = [&] { return x; }; + S <decltype (b)> c1 = { b, 42 }; + #pragma omp target + if (c1.x () != 42) + __builtin_abort (); +} + +/* { dg-final { scan-tree-dump {map\(to:c1\.x \[len: [0-9]+\]\) map\(tofrom:\*c1\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:c1\.x\.__x \[pointer assign, bias: 0\]\)} original { xfail *-*-* } } } */ + +void +f2 (int x, int y) +{ + auto b = [&] { return x; }; + S <decltype (b)> d2[2] = { { b, 42 }, { b, 42 } }; + #pragma omp target + if (d2[1].x () != 42 || d2[y].x () != 42) + __builtin_abort (); +} + +/* { dg-final { scan-tree-dump {map\(to:d2\[1\]\.x \[len: [0-9]+\]\) map\(tofrom:\*d2\[1\]\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:d2\[1\]\.x\.__x \[pointer assign, bias: 0\]\)} original { xfail *-*-* } } } + { dg-final { scan-tree-dump {map\(to:d2\[y\]\.x \[len: [0-9]+\]\) map\(tofrom:\*d2\[y\]\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:d2\[y\]\.x\.__x \[pointer assign, bias: 0\]\)} original { xfail *-*-* } } } */ + +void +f3 (int x) +{ + auto b = [&] { return x; }; + S <S <decltype (b)> > e3 = { { b, 42 }, 42 }; + #pragma omp target + if (e3.x.x () != 42) + __builtin_abort (); +} + +/* { dg-final { scan-tree-dump {map\(to:e3\.x\.x \[len: [0-9]+\]\) map\(tofrom:\*e3\.x\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:e3\.x\.x\.__x \[pointer assign, bias: 0\]\)} original { xfail *-*-* } } } */ + +void +f4 (int x) +{ + auto b = [&] { return x; }; + S <decltype (b)> c4 = { b, 42 }; + #pragma omp target + if (baz (c4) != 42) + __builtin_abort (); +} + +/* No lambda-typed expression appears in the target region's body; see + commit 84d5fecd2694ad70f391a20451d43d82fde37eb5 + "openmp c++: Temporary fix for some lambda mapping problems [PR127559]". + { dg-final { scan-tree-dump {(?n)map\([^)]*c4} original { xfail *-*-* } } } + Eventually might look as follows: + { dg-final { scan-tree-dump {map\(to:c4\.x \[len: [0-9]+\]\) map\(tofrom:\*c4\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:c4\.x\.__x \[pointer assign, bias: 0\]\)} original { xfail *-*-* } } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-5.C b/libgomp/testsuite/libgomp.c++/target-lambda-5.C index 6fbc65e2de82..fb5b61736716 100644 --- a/libgomp/testsuite/libgomp.c++/target-lambda-5.C +++ b/libgomp/testsuite/libgomp.c++/target-lambda-5.C @@ -1,5 +1,6 @@ // PR c++/127559 // { dg-do run } +// { dg-additional-options -fdump-tree-gimple } void foo (int *a) @@ -19,3 +20,6 @@ main () int a = 42; foo (&a); } + +/* Locally declared lambdas 'b', 'c' must not be mapped. + { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(to:\*a \[len: [0-9]+\]\) map\(firstprivate:a \[pointer assign, bias: 0\]\)$} gimple } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-6.C b/libgomp/testsuite/libgomp.c++/target-lambda-6.C index 3385c7ec187c..a997f2118f82 100644 --- a/libgomp/testsuite/libgomp.c++/target-lambda-6.C +++ b/libgomp/testsuite/libgomp.c++/target-lambda-6.C @@ -1,5 +1,7 @@ // PR c++/127559 // { dg-do run } +// { dg-xfail-run-if PR127559 { offload_device_nonshared_as } } +// { dg-additional-options {-fdump-tree-original -fdump-tree-gimple} } template <typename T> struct S { @@ -31,3 +33,11 @@ main () { foo (42, 0); } + +/* Lambda objects 'c.x', 'd[y].x' must be mapped, including their by-reference capture of 'x'. + { dg-final { scan-tree-dump {map\(to:c\.x \[len: [0-9]+\]\) map\(tofrom:\*c\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:c\.x\.__x \[pointer assign, bias: 0\]\)} original { xfail *-*-* } } } + { dg-final { scan-tree-dump {map\(to:d\[y\]\.x \[len: [0-9]+\]\) map\(tofrom:\*d\[y\]\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:d\[y\]\.x\.__x \[pointer assign, bias: 0\]\)} original { xfail *-*-* } } } + { dg-final { scan-tree-dump {map\(struct:c \[len: 1\]\) map\(to:c\.x \[len: [0-9]+\]\) map\(tofrom:\*c\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:c\.x\.__x \[pointer assign, bias: 0\]\)} gimple { xfail *-*-* } } } + { dg-final { scan-tree-dump {map\(struct_unord:d \[len: 1\]\) map\(to:d\[y\]\.x \[len: [0-9]+\]\) map\(tofrom:\*d\[y\]\.x\.__x \[len: [0-9]+\]\) map\(always_pointer:d\[y\]\.x\.__x \[pointer assign, bias: 0\]\)} gimple { xfail *-*-* } } } + ..., and in particular not as a whole, implicitly: + { dg-final { scan-tree-dump-not {runtime_implicit} gimple { xfail *-*-* } } } */
