diff options
| -rw-r--r-- | gcc/c-family/c-omp.c | 10 | ||||
| -rw-r--r-- | gcc/c/c-typeck.c | 63 | ||||
| -rw-r--r-- | gcc/cp/semantics.c | 65 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/gomp/pr99928-10.c | 48 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/gomp/pr99928-16.c | 16 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/gomp/pr99928-8.c | 44 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/gomp/pr99928-9.c | 22 | ||||
| -rw-r--r-- | gcc/tree.h | 5 |
8 files changed, 190 insertions, 83 deletions
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index 1b4a112..28fbb1d 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -1989,6 +1989,16 @@ c_omp_split_clauses (location_t loc, enum tree_code code, "%<parallel for%>, %<parallel for simd%>"); OMP_CLAUSE_REDUCTION_INSCAN (clauses) = 0; } + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)) != 0) + { + c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM); + OMP_CLAUSE_MAP_IMPLICIT (c) = 1; + OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; + cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = c; + } if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SCHEDULE)) != 0) { if (code == OMP_SIMD) diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index bee01df..5f32287 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13657,6 +13657,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !c_mark_addressable (t)) return false; @@ -13927,7 +13928,8 @@ tree c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head; + bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head; + bitmap_head oacc_reduction_head; tree c, t, type, *pc; tree simdlen = NULL_TREE, safelen = NULL_TREE; bool branch_seen = false; @@ -13947,7 +13949,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) has been seen, -2 if mixed inscan/normal reduction diagnosed. */ int reduction_seen = 0; bool allocate_seen = false; - bool firstprivate_implicit_moved = false; + bool implicit_moved = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -13957,6 +13959,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */ bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); + bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack); /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head instead. */ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); @@ -14389,28 +14392,36 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) - && !firstprivate_implicit_moved) + if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) && !implicit_moved) { - firstprivate_implicit_moved = true; - /* Move firstprivate clauses with - OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT set to the end of + move_implicit: + implicit_moved = true; + /* Move firstprivate and map clauses with + OMP_CLAUSE_{FIRSTPRIVATE,MAP}_IMPLICIT set to the end of clauses chain. */ - tree cl = NULL, *pc1 = pc, *pc2 = &cl; + tree cl1 = NULL_TREE, cl2 = NULL_TREE; + tree *pc1 = pc, *pc2 = &cl1, *pc3 = &cl2; while (*pc1) if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_FIRSTPRIVATE && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (*pc1)) { + *pc3 = *pc1; + pc3 = &OMP_CLAUSE_CHAIN (*pc3); + *pc1 = OMP_CLAUSE_CHAIN (*pc1); + } + else if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IMPLICIT (*pc1)) + { *pc2 = *pc1; pc2 = &OMP_CLAUSE_CHAIN (*pc2); *pc1 = OMP_CLAUSE_CHAIN (*pc1); } else pc1 = &OMP_CLAUSE_CHAIN (*pc1); - *pc2 = NULL; - *pc1 = cl; - if (pc1 != pc) - continue; + *pc3 = NULL; + *pc2 = cl2; + *pc1 = cl1; + continue; } t = OMP_CLAUSE_DECL (c); need_complete = true; @@ -14421,6 +14432,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) "%qE is not a variable in clause %<firstprivate%>", t); remove = true; } + else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) + && !OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT_TARGET (c) + && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) + remove = true; else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { @@ -14673,6 +14688,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; case OMP_CLAUSE_MAP: + if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) + goto move_implicit; + /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: @@ -14706,6 +14724,16 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { while (TREE_CODE (t) == COMPONENT_REF) t = TREE_OPERAND (t, 0); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IMPLICIT (c) + && (bitmap_bit_p (&map_head, DECL_UID (t)) + || bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, + DECL_UID (t)))) + { + remove = true; + break; + } if (bitmap_bit_p (&map_field_head, DECL_UID (t))) break; if (bitmap_bit_p (&map_head, DECL_UID (t))) @@ -14860,6 +14888,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IMPLICIT (c) + && (bitmap_bit_p (&map_head, DECL_UID (t)) + || bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))) + remove = true; + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) { if (bitmap_bit_p (&generic_head, DECL_UID (t)) @@ -14880,7 +14914,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else - bitmap_set_bit (&generic_head, DECL_UID (t)); + { + bitmap_set_bit (&generic_head, DECL_UID (t)); + bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); + } } else if (bitmap_bit_p (&map_head, DECL_UID (t)) && (ort != C_ORT_OMP diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 6fafd06..fe370a2 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5548,6 +5548,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c); if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && !cxx_mark_addressable (t)) return false; @@ -5574,6 +5575,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2)); + OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c); OMP_CLAUSE_DECL (c3) = ptr; if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH) @@ -6510,7 +6512,8 @@ tree finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head; + bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head; + bitmap_head oacc_reduction_head; tree c, t, *pc; tree safelen = NULL_TREE; bool branch_seen = false; @@ -6527,7 +6530,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bool allocate_seen = false; tree detach_seen = NULL_TREE; bool mergeable_seen = false; - bool firstprivate_implicit_moved = false; + bool implicit_moved = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -6537,6 +6540,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */ bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); + bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack); /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head instead. */ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); @@ -6852,28 +6856,36 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) - && !firstprivate_implicit_moved) + if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) && !implicit_moved) { - firstprivate_implicit_moved = true; - /* Move firstprivate clauses with - OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT set to the end of + move_implicit: + implicit_moved = true; + /* Move firstprivate and map clauses with + OMP_CLAUSE_{FIRSTPRIVATE,MAP}_IMPLICIT set to the end of clauses chain. */ - tree cl = NULL, *pc1 = pc, *pc2 = &cl; + tree cl1 = NULL_TREE, cl2 = NULL_TREE; + tree *pc1 = pc, *pc2 = &cl1, *pc3 = &cl2; while (*pc1) if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_FIRSTPRIVATE && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (*pc1)) { + *pc3 = *pc1; + pc3 = &OMP_CLAUSE_CHAIN (*pc3); + *pc1 = OMP_CLAUSE_CHAIN (*pc1); + } + else if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IMPLICIT (*pc1)) + { *pc2 = *pc1; pc2 = &OMP_CLAUSE_CHAIN (*pc2); *pc1 = OMP_CLAUSE_CHAIN (*pc1); } else pc1 = &OMP_CLAUSE_CHAIN (*pc1); - *pc2 = NULL; - *pc1 = cl; - if (pc1 != pc) - continue; + *pc3 = NULL; + *pc2 = cl2; + *pc1 = cl1; + continue; } t = omp_clause_decl_field (OMP_CLAUSE_DECL (c)); if (t) @@ -6904,6 +6916,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t); remove = true; } + else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) + && !OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT_TARGET (c) + && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) + remove = true; else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { @@ -7620,6 +7636,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; case OMP_CLAUSE_MAP: + if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) + goto move_implicit; + /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: @@ -7651,6 +7670,16 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IMPLICIT (c) + && (bitmap_bit_p (&map_head, DECL_UID (t)) + || bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, + DECL_UID (t)))) + { + remove = true; + break; + } if (bitmap_bit_p (&map_field_head, DECL_UID (t))) break; if (bitmap_bit_p (&map_head, DECL_UID (t))) @@ -7832,6 +7861,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IMPLICIT (c) + && (bitmap_bit_p (&map_head, DECL_UID (t)) + || bitmap_bit_p (&map_field_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, + DECL_UID (t)))) + remove = true; + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) { if (bitmap_bit_p (&generic_head, DECL_UID (t)) @@ -7852,7 +7888,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; } else - bitmap_set_bit (&generic_head, DECL_UID (t)); + { + bitmap_set_bit (&generic_head, DECL_UID (t)); + bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); + } } else if (bitmap_bit_p (&map_head, DECL_UID (t)) && !bitmap_bit_p (&map_field_head, DECL_UID (t))) diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-10.c b/gcc/testsuite/c-c++-common/gomp/pr99928-10.c index d2987a1..6c44600 100644 --- a/gcc/testsuite/c-c++-common/gomp/pr99928-10.c +++ b/gcc/testsuite/c-c++-common/gomp/pr99928-10.c @@ -95,22 +95,22 @@ bar (void) #pragma omp section r12[1]++; } - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 60\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r13 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 60\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r13 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r13 \\+ 4" "gimple" } } */ #pragma omp target parallel reduction(+:r13[1:15]) r13[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 64\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r14 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 64\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r14 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r14 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */ /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r14 \\+ 4" "gimple" } } *//* FIXME. */ #pragma omp target parallel for reduction(+:r14[1:16]) for (int i = 0; i < 64; i++) r14[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 68\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r15 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 68\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r15 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r15 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */ /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r15 \\+ 4" "gimple" } } *//* FIXME. */ @@ -118,31 +118,31 @@ bar (void) #pragma omp target parallel for simd reduction(+:r15[1:17]) for (int i = 0; i < 64; i++) r15[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 72\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r16 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 72\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r16 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } *//* FIXME: Should be shared, but firstprivate is an optimization. */ /* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */ /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */ #pragma omp target parallel loop reduction(+:r16[1:18]) for (int i = 0; i < 64; i++) r16[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 76\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r17 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 76\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r17 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r17 \\+ 4" "gimple" } } */ #pragma omp target teams reduction(+:r17[1:19]) r17[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 80\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r18 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 80\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r18 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r18 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r18 \\+ 4" "gimple" } } */ #pragma omp target teams distribute reduction(+:r18[1:20]) for (int i = 0; i < 64; i++) r18[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 84\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r19 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 84\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r19 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r19 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r19 \\+ 4" "gimple" } } */ @@ -151,8 +151,8 @@ bar (void) #pragma omp target teams distribute parallel for reduction(+:r19[1:21]) for (int i = 0; i < 64; i++) r19[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 88\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r20 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 88\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r20 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r20 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r20 \\+ 4" "gimple" } } */ @@ -162,8 +162,8 @@ bar (void) #pragma omp target teams distribute parallel for simd reduction(+:r20[1:22]) for (int i = 0; i < 64; i++) r20[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 92\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r21 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 92\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r21 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r21 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r21 \\+ 4" "gimple" } } */ @@ -171,9 +171,9 @@ bar (void) #pragma omp target teams distribute simd reduction(+:r21[1:23]) for (int i = 0; i < 64; i++) r21[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 96\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r22 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 96\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r22 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } *//* FIXME: Should be shared, but firstprivate is an optimization. */ /* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r22 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } *//* NOTE: This is implementation detail. */ @@ -182,8 +182,8 @@ bar (void) #pragma omp target teams loop reduction(+:r22[1:24]) for (int i = 0; i < 64; i++) r22[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 100\\\]" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r23 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 100\\\]" "gimple" } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r23 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r23 \\+ 4" "gimple" } } */ #pragma omp target simd reduction(+:r23[1:25]) diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-16.c b/gcc/testsuite/c-c++-common/gomp/pr99928-16.c new file mode 100644 index 0000000..84cd85d --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/pr99928-16.c @@ -0,0 +1,16 @@ +/* PR middle-end/99928 */ + +void +foo (void) +{ + int a[6] = {}; + #pragma omp target simd reduction(+:a[:3]) + for (int i = 0; i < 6; i++) + a[0]++; + #pragma omp target simd reduction(+:a[:3]) map(always, tofrom: a) + for (int i = 0; i < 6; i++) + a[0]++; + #pragma omp target simd reduction(+:a[:3]) map(always, tofrom: a[:6]) + for (int i = 0; i < 6; i++) + a[0]++; +} diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-8.c b/gcc/testsuite/c-c++-common/gomp/pr99928-8.c index fc57573..27e6ad1 100644 --- a/gcc/testsuite/c-c++-common/gomp/pr99928-8.c +++ b/gcc/testsuite/c-c++-common/gomp/pr99928-8.c @@ -94,48 +94,48 @@ bar (void) #pragma omp section r12++; } - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r13\\)" "gimple" } } */ #pragma omp target parallel reduction(+:r13) r13++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r14\\)" "gimple" } } *//* FIXME: This should be on for instead. */ /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:r14\\)" "gimple" } } *//* FIXME. */ #pragma omp target parallel for reduction(+:r14) for (int i = 0; i < 64; i++) r14++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } *//* FIXME: This should be on for instead. */ /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } *//* FIXME. */ /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } */ #pragma omp target parallel for simd reduction(+:r15) for (int i = 0; i < 64; i++) r15++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r16\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:r16\\)" "gimple" } } *//* NOTE: This is implementation detail. */ /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r16\\)" "gimple" } } *//* NOTE: This is implementation detail. */ #pragma omp target parallel loop reduction(+:r16) for (int i = 0; i < 64; i++) r16++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r17\\)" "gimple" } } */ #pragma omp target teams reduction(+:r17) r17++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r18\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r18\\)" "gimple" } } */ #pragma omp target teams distribute reduction(+:r18) for (int i = 0; i < 64; i++) r18++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } *//* FIXME: This should be on for instead. */ @@ -143,8 +143,8 @@ bar (void) #pragma omp target teams distribute parallel for reduction(+:r19) for (int i = 0; i < 64; i++) r19++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } *//* FIXME: This should be on for instead. */ @@ -153,16 +153,16 @@ bar (void) #pragma omp target teams distribute parallel for simd reduction(+:r20) for (int i = 0; i < 64; i++) r20++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */ #pragma omp target teams distribute simd reduction(+:r21) for (int i = 0; i < 64; i++) r21++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*shared\\(r22\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:r22\\)" "gimple" } } *//* NOTE: This is implementation detail. */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r22\\)" "gimple" } } *//* NOTE: This is implementation detail. */ @@ -171,8 +171,8 @@ bar (void) #pragma omp target teams loop reduction(+:r22) for (int i = 0; i < 64; i++) r22++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23" "gimple" { xfail *-*-* } } } */ - /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23" "gimple" } } */ + /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r23\\)" "gimple" } } */ #pragma omp target simd reduction(+:r23) for (int i = 0; i < 64; i++) diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-9.c b/gcc/testsuite/c-c++-common/gomp/pr99928-9.c index c049bb0..8623527 100644 --- a/gcc/testsuite/c-c++-common/gomp/pr99928-9.c +++ b/gcc/testsuite/c-c++-common/gomp/pr99928-9.c @@ -94,19 +94,19 @@ bar (void) #pragma omp section r12[1]++; } - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r13 \\+ 4" "gimple" } } */ #pragma omp target parallel reduction(+:r13[1:2]) r13[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r14 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */ /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r14 \\+ 4" "gimple" } } *//* FIXME. */ #pragma omp target parallel for reduction(+:r14[1:2]) for (int i = 0; i < 64; i++) r14[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r15 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead. */ /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r15 \\+ 4" "gimple" } } *//* FIXME. */ @@ -114,7 +114,7 @@ bar (void) #pragma omp target parallel for simd reduction(+:r15[1:2]) for (int i = 0; i < 64; i++) r15[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r16\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */ @@ -122,19 +122,19 @@ bar (void) #pragma omp target parallel loop reduction(+:r16[1:2]) for (int i = 0; i < 64; i++) r16[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r17 \\+ 4" "gimple" } } */ #pragma omp target teams reduction(+:r17[1:2]) r17[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r18 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r18 \\+ 4" "gimple" } } */ #pragma omp target teams distribute reduction(+:r18[1:2]) for (int i = 0; i < 64; i++) r18[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r19 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r19 \\+ 4" "gimple" } } */ @@ -143,7 +143,7 @@ bar (void) #pragma omp target teams distribute parallel for reduction(+:r19[1:2]) for (int i = 0; i < 64; i++) r19[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r20 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r20 \\+ 4" "gimple" } } */ @@ -153,7 +153,7 @@ bar (void) #pragma omp target teams distribute parallel for simd reduction(+:r20[1:2]) for (int i = 0; i < 64; i++) r20[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r21 \\+ 4" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r21 \\+ 4" "gimple" } } */ @@ -161,7 +161,7 @@ bar (void) #pragma omp target teams distribute simd reduction(+:r21[1:2]) for (int i = 0; i < 64; i++) r21[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*shared\\(r22\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r22 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail. */ @@ -171,7 +171,7 @@ bar (void) #pragma omp target teams loop reduction(+:r22[1:2]) for (int i = 0; i < 64; i++) r22[1]++; - /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */ + /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */ /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */ /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r23 \\+ 4" "gimple" } } */ #pragma omp target simd reduction(+:r23[1:2]) @@ -1654,6 +1654,11 @@ class auto_suppress_location_wrappers variable. */ #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \ TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) +/* Nonzero on map clauses added implicitly for reduction clauses on combined + or composite constructs. They shall be removed if there is an explicit + map clause. */ +#define OMP_CLAUSE_MAP_IMPLICIT(NODE) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.default_def_flag) /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' clause. */ |