On Thu, Jun 16, 2016 at 09:05:40PM +0200, Jakub Jelinek wrote: > Fixed thusly, bootstrapped/regtested on x86_64-linux and i686-linux, and > tested with x86_64-intelmicemul-linux offloading on x86_64-linux, committed > to trunk.
And the following testcase shows similar issues in the array section handling path. Tested on x86_64-linux and i686-linux, plus with intelmicemul offloading, committed to trunk. 2016-06-17 Jakub Jelinek <ja...@redhat.com> * semantics.c (handle_omp_array_sections_1): Don't ICE when processing_template_decl when checking for bitfields and unions. Look through REFERENCE_REF_P as base of COMPONENT_REF. (finish_omp_clauses): Look through REFERENCE_REF_P even for array sections with COMPONENT_REF bases. * testsuite/libgomp.c++/target-21.C: New test. --- gcc/cp/semantics.c.jj 2016-06-16 17:29:53.000000000 +0200 +++ gcc/cp/semantics.c 2016-06-17 14:34:21.929203440 +0200 @@ -4487,7 +4487,8 @@ handle_omp_array_sections_1 (tree c, tre || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM) && !type_dependent_expression_p (t)) { - if (DECL_BIT_FIELD (TREE_OPERAND (t, 1))) + if (TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL + && DECL_BIT_FIELD (TREE_OPERAND (t, 1))) { error_at (OMP_CLAUSE_LOCATION (c), "bit-field %qE in %qs clause", @@ -4496,7 +4497,8 @@ handle_omp_array_sections_1 (tree c, tre } while (TREE_CODE (t) == COMPONENT_REF) { - if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == UNION_TYPE) + if (TREE_TYPE (TREE_OPERAND (t, 0)) + && TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == UNION_TYPE) { error_at (OMP_CLAUSE_LOCATION (c), "%qE is a member of a union", t); @@ -4504,6 +4506,8 @@ handle_omp_array_sections_1 (tree c, tre } t = TREE_OPERAND (t, 0); } + if (REFERENCE_REF_P (t)) + t = TREE_OPERAND (t, 0); } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { @@ -6623,6 +6627,8 @@ finish_omp_clauses (tree clauses, enum c { while (TREE_CODE (t) == COMPONENT_REF) t = TREE_OPERAND (t, 0); + if (REFERENCE_REF_P (t)) + t = TREE_OPERAND (t, 0); if (bitmap_bit_p (&map_field_head, DECL_UID (t))) break; if (bitmap_bit_p (&map_head, DECL_UID (t))) --- libgomp/testsuite/libgomp.c++/target-21.C.jj 2016-06-17 13:18:59.684314656 +0200 +++ libgomp/testsuite/libgomp.c++/target-21.C 2016-06-17 15:10:21.860516966 +0200 @@ -0,0 +1,173 @@ +extern "C" void abort (); +struct T { char t[270]; }; +struct S { int (&x)[10]; int *&y; T t; int &z; S (); ~S (); }; + +template <int N> +void +foo (S s) +{ + int err; + #pragma omp target map (s.x[0:N], s.y[0:N]) map (s.t.t[16:3]) map (from: err) + { + err = s.x[2] != 28 || s.y[2] != 37 || s.t.t[17] != 81; + s.x[2]++; + s.y[2]++; + s.t.t[17]++; + } + if (err || s.x[2] != 29 || s.y[2] != 38 || s.t.t[17] != 82) + abort (); +} + +template <int N> +void +bar (S s) +{ + int err; + #pragma omp target map (s.x, s.z)map(from:err) + { + err = s.x[2] != 29 || s.z != 6; + s.x[2]++; + s.z++; + } + if (err || s.x[2] != 30 || s.z != 7) + abort (); +} + +template <int N> +void +foo2 (S &s) +{ + int err; + #pragma omp target map (s.x[N:10], s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3]) + { + err = s.x[2] != 30 || s.y[2] != 38 || s.t.t[17] != 81; + s.x[2]++; + s.y[2]++; + s.t.t[17]++; + } + if (err || s.x[2] != 31 || s.y[2] != 39 || s.t.t[17] != 82) + abort (); +} + +template <int N> +void +bar2 (S &s) +{ + int err; + #pragma omp target map (s.x, s.z)map(from:err) + { + err = s.x[2] != 31 || s.z != 7; + s.x[2]++; + s.z++; + } + if (err || s.x[2] != 32 || s.z != 8) + abort (); +} + +template <typename U> +void +foo3 (U s) +{ + int err; + #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3]) + { + err = s.x[2] != 32 || s.y[2] != 39 || s.t.t[17] != 82; + s.x[2]++; + s.y[2]++; + s.t.t[17]++; + } + if (err || s.x[2] != 33 || s.y[2] != 40 || s.t.t[17] != 83) + abort (); +} + +template <typename U> +void +bar3 (U s) +{ + int err; + #pragma omp target map (s.x, s.z)map(from:err) + { + err = s.x[2] != 33 || s.z != 8; + s.x[2]++; + s.z++; + } + if (err || s.x[2] != 34 || s.z != 9) + abort (); +} + +template <typename U> +void +foo4 (U &s) +{ + int err; + #pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3]) + { + err = s.x[2] != 34 || s.y[2] != 40 || s.t.t[17] != 82; + s.x[2]++; + s.y[2]++; + s.t.t[17]++; + } + if (err || s.x[2] != 35 || s.y[2] != 41 || s.t.t[17] != 83) + abort (); +} + +template <typename U> +void +bar4 (U &s) +{ + int err; + #pragma omp target map (s.x, s.z)map(from:err) + { + err = s.x[2] != 35 || s.z != 9; + s.x[2]++; + s.z++; + } + if (err || s.x[2] != 36 || s.z != 10) + abort (); +} + +int xt[10] = { 1, 2, 28, 3, 4, 5, 6, 7, 8, 9 }; +int yt[10] = { 1, 2, 37, 3, 4, 5, 6, 7, 8, 9 }; +int *yp = yt; +int zt = 6; + +S::S () : x (xt), y (yp), z (zt) +{ +} + +S::~S () +{ +} + +int +main () +{ + S s; + s.t.t[16] = 5; + s.t.t[17] = 81; + s.t.t[18] = 9; + foo <10> (s); + if (s.t.t[17] != 81) + abort (); + bar <7> (s); + foo2 <0> (s); + if (s.t.t[17] != 82) + abort (); + bar2 <21> (s); + foo3 <S> (s); + if (s.t.t[17] != 82) + abort (); + bar3 <S> (s); + foo4 <S> (s); + if (s.t.t[17] != 83) + abort (); + bar4 <S> (s); + s.x[2] -= 4; + s.y[2] -= 2; + s.z -= 2; + s.t.t[17]--; + foo3 <S &> (s); + if (s.t.t[17] != 83) + abort (); + bar3 <S &> (s); +} Jakub