2016-05-10 Cesar Philippidis gcc/c/ * c-parser.c (c_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTER. * c-typeck.c (c_finish_omp_clauses): Add specific errors and warning messages for OpenACC. Use firsrtprivate pointers for OpenACC subarrays. gcc/cp/ * parser.c (cp_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTERS. * semantics.c (finish_omp_clauses): Add specific errors and warning messages for OpenACC. Use firsrtprivate pointers for OpenACC subarrays. Add some checking for duplicate reference-typed data mappings. gcc/ * gimplify.c (omp_notice_variable): Use zero-length arrays for data pointers inside OACC_DATA regions. (gimplify_scan_omp_clauses): Prune firstprivate clause associated with OACC_DATA, OACC_ENTER_DATA and OACC_EXIT data regions. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: Adjust test. * c-c++-common/goacc/deviceptr-1.c: Likewise. * c-c++-common/goacc/kernels-alias-3.c: Likewise. * c-c++-common/goacc/kernels-alias-4.c: Likewise. * c-c++-common/goacc/kernels-alias-5.c: Likewise. * c-c++-common/goacc/kernels-alias-8.c: Likewise. * c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise. * c-c++-common/goacc/pcopy.c: Likewise. * c-c++-common/goacc/pcopyin.c: Likewise. * c-c++-common/goacc/pcopyout.c: Likewise. * c-c++-common/goacc/pcreate.c: Likewise. * c-c++-common/goacc/pr70688.c: New test. * c-c++-common/goacc/present-1.c: Adjust test. * c-c++-common/goacc/reduction-5.c: Likewise. * g++.dg/goacc/data-1.C: New test. * g++.dg/goacc/data-2.C: New test. * g++.dg/gomp/template-data.C: New test. libgomp/ * oacc-mem.c (acc_malloc): Update handling of shared-memory targets. (acc_free): Likewise. (acc_memcpy_to_device): Likewise. (acc_memcpy_from_device): Likewise. (acc_deviceptr): Likewise. (acc_hostptr): Likewise. (acc_is_present): Likewise. (acc_map_data): Likewise. (acc_unmap_data): Likewise. (present_create_copy): Likewise. (delete_copyout): Likewise. (update_dev_host): Likewise. * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Remove xfail. * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: New test. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Adjust test. * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c: New test. * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Adjust test so that it only runs on nvptx targets. * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 6523c08..5794c68 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -13602,6 +13602,7 @@ c_parser_oacc_declare (c_parser *parser) switch (OMP_CLAUSE_MAP_KIND (t)) { + case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_DEVICEPTR: diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 861aa12..1df18e2 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12502,7 +12502,7 @@ 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; + bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head; tree c, t, type, *pc; tree simdlen = NULL_TREE, safelen = NULL_TREE; bool branch_seen = false; @@ -12519,6 +12519,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -12542,7 +12543,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ort & C_ORT_OMP)) + if (handle_omp_array_sections (c, ort & (C_ORT_OMP | C_ORT_ACC))) { remove = true; break; @@ -12856,6 +12857,17 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } + else if (ort == C_ORT_ACC + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) + { + if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error ("%qD appears more than once in reduction clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_reduction_head, DECL_UID (t)); + } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) @@ -12867,7 +12879,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE && bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clasues", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -12893,7 +12908,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clasues", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -12986,7 +13004,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ort & C_ORT_OMP)) + if (handle_omp_array_sections (c, ort & (C_ORT_OMP | C_ORT_ACC))) remove = true; break; } @@ -13009,7 +13027,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ort & C_ORT_OMP)) + if (handle_omp_array_sections (c, ort & (C_ORT_OMP | C_ORT_ACC))) remove = true; else { @@ -13036,6 +13054,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion" " clauses", t); + else if (ort == C_ORT_ACC) + error ("%qD appears more than once in data" + " clauses", t); else error ("%qD appears more than once in map" " clauses", t); @@ -13137,7 +13158,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -13147,6 +13171,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion clauses", t); + else if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); else error ("%qD appears more than once in map clauses", t); remove = true; @@ -13154,7 +13180,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index f4c6f74..b87ccef 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -35225,6 +35225,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok) gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP); switch (OMP_CLAUSE_MAP_KIND (t)) { + case GOMP_MAP_FIRSTPRIVATE_POINTER: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_DEVICEPTR: diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index fed7e88..aef82b6 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5796,12 +5796,14 @@ 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; + bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head; tree c, t, *pc; tree safelen = NULL_TREE; bool branch_seen = false; bool copyprivate_seen = false; bool ordered_seen = false; + bool allow_fields = (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP + || ort == C_ORT_ACC; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -5810,6 +5812,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -5829,8 +5832,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD) - == C_ORT_OMP))) + if (handle_omp_array_sections (c, allow_fields)) { remove = true; break; @@ -6040,6 +6042,17 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } + else if (ort == C_ORT_ACC + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) + { + if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + { + error ("%qD appears more than once in reduction clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_reduction_head, DECL_UID (t)); + } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) @@ -6050,7 +6063,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE && bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6076,7 +6092,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_note_field_privatization (t, OMP_CLAUSE_DECL (c)); else t = OMP_CLAUSE_DECL (c); - if (t == current_class_ptr) + if (ort != C_ORT_ACC && t == current_class_ptr) { error ("% allowed in OpenMP only in %" " clauses"); @@ -6103,7 +6119,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6551,8 +6570,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD) - == C_ORT_OMP))) + if (handle_omp_array_sections (c, allow_fields)) remove = true; break; } @@ -6586,8 +6604,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD) - == C_ORT_OMP))) + if (handle_omp_array_sections (c, allow_fields)) remove = true; else { @@ -6616,6 +6633,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion" " clauses", t); + else if (ort == C_ORT_ACC) + error ("%qD appears more than once in data" + " clauses", t); else error ("%qD appears more than once in map" " clauses", t); @@ -6627,6 +6647,27 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&map_field_head, DECL_UID (t)); } } + else if (TREE_CODE (t) == TREE_LIST) + { + while (TREE_CODE (t = TREE_CHAIN (t)) == TREE_LIST) + ; + + if (DECL_P (t)) + { + if (bitmap_bit_p (&map_head, DECL_UID (t))) + { + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data " + "clauses", t); + else + error ("%qD appears more than once in map " + "clauses", t); + remove = true; + } + else + bitmap_set_bit (&map_head, DECL_UID (t)); + } + } } break; } @@ -6703,7 +6744,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (t == current_class_ptr) + else if (ort != C_ORT_ACC && t == current_class_ptr) { error ("% allowed in OpenMP only in %" " clauses"); @@ -6752,7 +6793,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else @@ -6762,6 +6806,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); else error ("%qD appears more than once in map clauses", t); remove = true; @@ -6769,7 +6815,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) { - error ("%qD appears both in data and map clauses", t); + if (ort == C_ORT_ACC) + error ("%qD appears more than once in data clauses", t); + else + error ("%qD appears both in data and map clauses", t); remove = true; } else diff --git a/gcc/gimplify.c b/gcc/gimplify.c index f13980d..512b3dd 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6255,6 +6255,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) error ("variable %qE declared in enclosing " "% region", DECL_NAME (decl)); nflags |= GOVD_MAP; + if (octx->region_type == ORT_ACC_DATA + && (n2->value & GOVD_MAP_0LEN_ARRAY)) + nflags |= GOVD_MAP_0LEN_ARRAY; goto found_outer; } } @@ -6830,9 +6833,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { case OMP_TARGET: break; + case OACC_DATA: + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) + break; case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: case OACC_HOST_DATA: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -7286,6 +7294,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, omp_notice_variable (outer_ctx, t, true); } } + if (code == OACC_DATA && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + flags |= GOVD_MAP_0LEN_ARRAY; omp_add_variable (ctx, decl, flags); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c)) @@ -7544,6 +7555,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gcc_unreachable (); } + if (code == OACC_DATA && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + remove = true; if (remove) *list_p = OMP_CLAUSE_CHAIN (c); else diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c index 7a1cf68..6245beb 100644 --- a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c +++ b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -2,12 +2,12 @@ void fun (void) { float *fp; -#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; } diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c index 08ddb10..3aa0e8a 100644 --- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c +++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -47,7 +47,7 @@ fun2 (void) /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */ /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */ /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */ - /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */ + /* { dg-error "'fp' appears more than once in data clauses" "fp more than once" { target *-*-* } 46 } */ ; } @@ -55,11 +55,11 @@ void fun3 (void) { float *fp; -#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; -#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ ; } diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c index 6989c1c..2934f12 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c @@ -17,5 +17,5 @@ foo (void) /* Only the omp_data_i related loads should be annotated with non-base 0 cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c index d41802c..f6ee5b5 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c @@ -19,5 +19,5 @@ foo (void) /* Only the omp_data_i related loads should be annotated with non-base 0 cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c index 6fefe183..74425fb 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c @@ -15,5 +15,5 @@ foo (int *a) /* Only the omp_data_i related loads should be annotated with cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 4 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c index 3b91acd..69200cc 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c @@ -7,7 +7,7 @@ extern void *acc_copyin (void *, size_t); void foo (int *a, size_t n) { - int *p = (int *)acc_copyin (&a, n); + int *p = (int *)acc_copyin (a, n); #pragma acc kernels deviceptr (p) pcopy(a[0:n]) { @@ -18,5 +18,5 @@ foo (int *a, size_t n) /* Only the omp_data_i related loads should be annotated with cliques. */ /* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */ -/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */ +/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c index 1eb56eb..1ea0e73 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c @@ -31,6 +31,5 @@ foo (void) free (c); } -/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" } } */ -/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */ -/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" { target c } } } */ +/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" { target c } } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopy.c b/gcc/testsuite/c-c++-common/goacc/pcopy.c index 02c4383..0e0aad5 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopy.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopy.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyin.c b/gcc/testsuite/c-c++-common/goacc/pcopyin.c index 10911fc..3085251 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopyin.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopyin.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyout.c b/gcc/testsuite/c-c++-common/goacc/pcopyout.c index 703ac2f..47c454c 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcopyout.c +++ b/gcc/testsuite/c-c++-common/goacc/pcopyout.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pcreate.c b/gcc/testsuite/c-c++-common/goacc/pcreate.c index 00bf155..a403e5a 100644 --- a/gcc/testsuite/c-c++-common/goacc/pcreate.c +++ b/gcc/testsuite/c-c++-common/goacc/pcreate.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c new file mode 100644 index 0000000..5a23665 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c @@ -0,0 +1,48 @@ +const int n = 100; + +int +private_reduction () +{ + int i, r; + + #pragma acc parallel + #pragma acc loop private (r) reduction (+:r) + for (i = 0; i < 100; i++) + r += 10; + + return r; +} + +int +parallel_reduction () +{ + int sum = 0; + int dummy = 0; + +#pragma acc data copy (dummy) + { +#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) + { + int v = 5; + sum += 10 + v; + } + } + + return sum; +} + +int +main () +{ + int i, s = 0; + +#pragma acc parallel num_gangs (10) copy (s) reduction (+:s) + for (i = 0; i < n; i++) + s += i+1; + +#pragma acc parallel num_gangs (10) reduction (+:s) copy (s) + for (i = 0; i < n; i++) + s += i+1; + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/present-1.c b/gcc/testsuite/c-c++-common/goacc/present-1.c index 7537948..51362b2 100644 --- a/gcc/testsuite/c-c++-common/goacc/present-1.c +++ b/gcc/testsuite/c-c++-common/goacc/present-1.c @@ -7,4 +7,4 @@ f (char *cp) ; } -/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-5.c b/gcc/testsuite/c-c++-common/goacc/reduction-5.c index 74daad3..dfdbab9 100644 --- a/gcc/testsuite/c-c++-common/goacc/reduction-5.c +++ b/gcc/testsuite/c-c++-common/goacc/reduction-5.c @@ -7,9 +7,9 @@ main(void) { int v1; -#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "appears more than once in data clauses" } */ +#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "invalid private reduction" } */ ; -#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "appears more than once in data clauses" } */ +#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "invalid private reduction" } */ ; return 0; diff --git a/gcc/testsuite/g++.dg/goacc/data-1.C b/gcc/testsuite/g++.dg/goacc/data-1.C new file mode 100644 index 0000000..54676dc --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/data-1.C @@ -0,0 +1,39 @@ +void +foo (int &a, int (&b)[100], int &n) +{ +#pragma acc enter data copyin (a, b) async wait +#pragma acc enter data create (b[20:30]) async wait +#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */ +#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */ +#pragma acc exit data delete (a) if (0) +#pragma acc exit data copyout (b) if (a) +#pragma acc exit data delete (b) +#pragma acc enter /* { dg-error "expected 'data' in" } */ +#pragma acc exit /* { dg-error "expected 'data' in" } */ +#pragma acc enter data /* { dg-error "has no data movement clause" } */ +#pragma acc exit data /* { dg-error "has no data movement clause" } */ +#pragma acc enter Data /* { dg-error "invalid pragma before" } */ +#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */ +} + +template +void +foo (T &a, T (&b)[100], T &n) +{ +#pragma acc enter data copyin (a, b) async wait +#pragma acc enter data create (b[20:30]) async wait +#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */ +#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */ +#pragma acc exit data delete (a) if (0) +#pragma acc exit data copyout (b) if (a) +#pragma acc exit data delete (b) +#pragma acc enter /* { dg-error "expected 'data' in" } */ +#pragma acc exit /* { dg-error "expected 'data' in" } */ +#pragma acc enter data /* { dg-error "has no data movement clause" } */ +#pragma acc exit data /* { dg-error "has no data movement clause" } */ +#pragma acc enter Data /* { dg-error "invalid pragma before" } */ +#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */ +} + +/* { dg-error "has no data movement clause" "" { target *-*-* } 6 } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } 25 } */ diff --git a/gcc/testsuite/g++.dg/goacc/data-2.C b/gcc/testsuite/g++.dg/goacc/data-2.C new file mode 100644 index 0000000..efa002d --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/data-2.C @@ -0,0 +1,30 @@ +void +fun (float (&fp)[100]) +{ + float *dptr = &fp[50]; + +#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +#pragma acc data create(fp[:10]) deviceptr(dptr) + ; +#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +} + +template +void +fun (T (&fp)[100]) +{ + T *dptr = &fp[50]; + +#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +#pragma acc data create(fp[:10]) deviceptr(dptr) + ; +#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +} diff --git a/gcc/testsuite/g++.dg/gomp/template-data.C b/gcc/testsuite/g++.dg/gomp/template-data.C new file mode 100644 index 0000000..0be14d4 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/template-data.C @@ -0,0 +1,18 @@ +void +fun (float (&fp)[100]) +{ + float *dptr = &fp[50]; + +#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */ + ; +} + +template +void +fun (T (&fp)[100]) +{ + T *dptr = &fp[50]; + +#pragma omp target data map(tofrom:fp[0:2], fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +} diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index ce1905c..665e208 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -32,6 +32,7 @@ #include "gomp-constants.h" #include "oacc-int.h" #include +#include #include /* Return block containing [H->S), or NULL if not contained. The device lock @@ -104,6 +105,9 @@ acc_malloc (size_t s) assert (thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return malloc (s); + return thr->dev->alloc_func (thr->dev->target_id, s); } @@ -124,6 +128,9 @@ acc_free (void *d) struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return free (d); + gomp_mutex_lock (&acc_dev->lock); /* We don't have to call lazy open here, as the ptr value must have @@ -154,6 +161,12 @@ acc_memcpy_to_device (void *d, void *h, size_t s) assert (thr && thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + { + memmove (d, h, s); + return; + } + thr->dev->host2dev_func (thr->dev->target_id, d, h, s); } @@ -166,6 +179,12 @@ acc_memcpy_from_device (void *h, void *d, size_t s) assert (thr && thr->dev); + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + { + memmove (h, d, s); + return; + } + thr->dev->dev2host_func (thr->dev->target_id, h, d, s); } @@ -184,6 +203,9 @@ acc_deviceptr (void *h) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h; + gomp_mutex_lock (&dev->lock); n = lookup_host (dev, h, 1); @@ -218,6 +240,9 @@ acc_hostptr (void *d) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return d; + gomp_mutex_lock (&acc_dev->lock); n = lookup_dev (acc_dev->openacc.data_environ, d, 1); @@ -252,6 +277,9 @@ acc_is_present (void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h != NULL; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -271,7 +299,7 @@ acc_is_present (void *h, size_t s) void acc_map_data (void *h, void *d, size_t s) { - struct target_mem_desc *tgt; + struct target_mem_desc *tgt = NULL; size_t mapnum = 1; void *hostaddrs = h; void *devaddrs = d; @@ -287,9 +315,6 @@ acc_map_data (void *h, void *d, size_t s) { if (d != h) gomp_fatal ("cannot map data on shared-memory system"); - - tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, - GOMP_MAP_VARS_OPENACC); } else { @@ -335,6 +360,10 @@ acc_unmap_data (void *h) /* No need to call lazy open, as the address must have been mapped. */ + /* This is a no-op on shared-memory targets. */ + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + size_t host_size; gomp_mutex_lock (&acc_dev->lock); @@ -405,6 +434,9 @@ present_create_copy (unsigned f, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return h; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -496,6 +528,9 @@ delete_copyout (unsigned f, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); @@ -553,6 +588,9 @@ update_dev_host (int is_dev, void *h, size_t s) struct goacc_thread *thr = goacc_thread (); struct gomp_device_descr *acc_dev = thr->dev; + if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return; + gomp_mutex_lock (&acc_dev->lock); n = lookup_host (acc_dev, h, s); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c index f3b490a..d478ce2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c @@ -1,6 +1,4 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ -/* . - { dg-xfail-run-if "TODO" { *-*-* } } */ /* { dg-additional-options "-lcuda" } */ #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c new file mode 100644 index 0000000..e1aa2c9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c @@ -0,0 +1,185 @@ +/* This test is similar to data-2.c, but it uses acc_* library functions + to move data. */ + +/* { dg-do run } */ + +#include +#include +#include + +int +main (int argc, char **argv) +{ + int N = 128; //1024 * 1024; + float *a, *b, *c, *d, *e; + void *d_a, *d_b, *d_c, *d_d; + int i; + int nbytes; + + nbytes = N * sizeof (float); + + a = (float *) malloc (nbytes); + b = (float *) malloc (nbytes); + c = (float *) malloc (nbytes); + d = (float *) malloc (nbytes); + e = (float *) malloc (nbytes); + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + + d_a = acc_copyin (a, nbytes); + d_b = acc_copyin (b, nbytes); + acc_copyin (&N, sizeof (int)); + +#pragma acc parallel present (a[0:N], b[0:N], N) async wait +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + + acc_wait_all (); + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + + for (i = 0; i < N; i++) + { + assert (a[i] == 3.0); + assert (b[i] == 3.0); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = a[i]; + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + + for (i = 0; i < N; i++) + { + assert (a[i] == 2.0); + assert (b[i] == 2.0); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + d_c = acc_copyin (c, nbytes); + d_d = acc_copyin (d, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) +#pragma acc loop + for (i = 0; i < N; i++) + b[i] = (a[i] * a[i] * a[i]) / a[i]; + +#pragma acc parallel present (a[0:N], c[0:N], N) async (2) +#pragma acc loop + for (i = 0; i < N; i++) + c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; + +#pragma acc parallel present (a[0:N], d[0:N], N) async (3) +#pragma acc loop + for (i = 0; i < N; i++) + d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; + + acc_wait_all (); + + acc_memcpy_from_device (a, d_a, nbytes); + acc_memcpy_from_device (b, d_b, nbytes); + acc_memcpy_from_device (c, d_c, nbytes); + acc_memcpy_from_device (d, d_d, nbytes); + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 9.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + e[i] = 0.0; + } + + acc_update_device (a, nbytes); + acc_update_device (b, nbytes); + acc_update_device (c, nbytes); + acc_update_device (d, nbytes); + acc_copyin (e, nbytes); + +#pragma acc parallel present (a[0:N], b[0:N], N) async (1) + for (int ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + +#pragma acc parallel present (a[0:N], c[0:N], N) async (2) + for (int ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + +#pragma acc parallel present (a[0:N], d[0:N], N) async (3) + for (int ii = 0; ii < N; ii++) + d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) \ + async (4) + for (int ii = 0; ii < N; ii++) + e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; + + acc_wait_all (); + acc_copyout (a, nbytes); + acc_copyout (b, nbytes); + acc_copyout (c, nbytes); + acc_copyout (d, nbytes); + acc_copyout (e, nbytes); + acc_delete (&N, sizeof (int)); + + for (i = 0; i < N; i++) + { + if (a[i] != 2.0) + abort (); + + if (b[i] != 4.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + + if (e[i] != 11.0) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c index f867a66..c1c0825 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c @@ -1,3 +1,5 @@ +/* Test 'acc enter/exit data' regions. */ + /* { dg-do run } */ #include @@ -25,7 +27,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async -#pragma acc parallel async wait +#pragma acc parallel present (a[0:N], b[0:N]) async wait #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -49,7 +51,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -76,17 +78,17 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = (a[i] * a[i] * a[i]) / a[i]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) #pragma acc loop for (i = 0; i < N; i++) c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) #pragma acc loop for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; @@ -120,26 +122,27 @@ main (int argc, char **argv) #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) for (int ii = 0; ii < N; ii++) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) for (int ii = 0; ii < N; ii++) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) for (int ii = 0; ii < N; ii++) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; -#pragma acc parallel wait (1) async (4) +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \ + wait (1) async (4) for (int ii = 0; ii < N; ii++) e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ + copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) #pragma acc wait (1) - for (i = 0; i < N; i++) { if (a[i] != 2.0) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c index 747109f..0bf706a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c @@ -1,3 +1,5 @@ +/* Test 'acc enter/exit data' regions with 'acc update'. */ + /* { dg-do run } */ #include @@ -25,7 +27,7 @@ main (int argc, char **argv) } #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async -#pragma acc parallel async wait +#pragma acc parallel present (a[0:N], b[0:N]) async wait #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -49,7 +51,7 @@ main (int argc, char **argv) } #pragma acc update device (a[0:N], b[0:N]) async (1) -#pragma acc parallel async (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) #pragma acc loop for (i = 0; i < N; i++) b[i] = a[i]; @@ -78,17 +80,17 @@ main (int argc, char **argv) #pragma acc update device (b[0:N]) async (2) #pragma acc enter data copyin (c[0:N], d[0:N]) async (3) -#pragma acc parallel async (1) wait (1,2) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1,2) #pragma acc loop for (i = 0; i < N; i++) b[i] = (a[i] * a[i] * a[i]) / a[i]; -#pragma acc parallel async (2) wait (1,3) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1,3) #pragma acc loop for (i = 0; i < N; i++) c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i]; -#pragma acc parallel async (3) wait (1,3) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1,3) #pragma acc loop for (i = 0; i < N; i++) d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i]; @@ -123,27 +125,28 @@ main (int argc, char **argv) #pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) #pragma acc enter data copyin (e[0:N]) async (5) -#pragma acc parallel async (1) wait (1) +#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1) for (int ii = 0; ii < N; ii++) b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; -#pragma acc parallel async (2) wait (1) +#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1) for (int ii = 0; ii < N; ii++) c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; -#pragma acc parallel async (3) wait (1) +#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1) for (int ii = 0; ii < N; ii++) d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; -#pragma acc parallel wait (1,5) async (4) +#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \ + wait (1,5) async (4) for (int ii = 0; ii < N; ii++) e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) +#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \ + copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1) #pragma acc exit data delete (N) #pragma acc wait (1) - for (i = 0; i < N; i++) { if (a[i] != 2.0) @@ -162,5 +165,11 @@ main (int argc, char **argv) abort (); } + free (a); + free (b); + free (c); + free (d); + free (e); + return 0; } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c new file mode 100644 index 0000000..b5b37b2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c @@ -0,0 +1,70 @@ +/* Verify enter/exit data interoperablilty between pragmas and + acc library calls. */ + +/* { dg-do run } */ + +#include +#include +#include + +int +main () +{ + int *p = (int *)malloc (sizeof (int)); + + /* Test 1: pragma input, library output. */ + +#pragma acc enter data copyin (p[0:1]) + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 1; + } + + acc_copyout (p, sizeof (int)); + + assert (p[0] == 1); + + /* Test 2: library input, pragma output. */ + + acc_copyin (p, sizeof (int)); + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 2; + } + +#pragma acc exit data copyout (p[0:1]) + + assert (p[0] == 2); + + /* Test 3: library input, library output. */ + + acc_copyin (p, sizeof (int)); + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 3; + } + + acc_copyout (p, sizeof (int)); + + assert (p[0] == 3); + + /* Test 4: pragma input, pragma output. */ + +#pragma acc enter data copyin (p[0:1]) + +#pragma acc parallel present (p[0:1]) num_gangs (1) + { + p[0] = 3; + } + +#pragma acc exit data copyout (p[0:1]) + + assert (p[0] == 3); + + free (p); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c index 7098ef3..d665533 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present and acc_delete. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c index a9632f7..ee21257 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c index 4f6a731..50c1701 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Check acc_is_present and acc_copyout. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c index 28e4e5c..c81a78d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Test if duplicate data mappings with acc_copy_in. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c index 7d1767e..a3487e8 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c @@ -1,4 +1,7 @@ -/* { dg-do run } */ +/* Check acc_copyout failure with acc_device_nvidia. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c index 160b33c..b686cc9 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Verify that acc_delete unregisters data mappings on the device. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c index 4f8e14c..25ceb3a 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c index d908700..b170f81 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c index a6c0197..65ff440 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c index 2339dd6..fd4dc59 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_copyin and acc_copyout on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c index d7de8e3..09e2817 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_create, acc_is_present and acc_delete. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c index bb709d3..5f00ccb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_create and acc_delete on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c index 9304daa..7a96ab2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_delete with a NULL address on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c index 92e3858..318a060 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_delete with size zero on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c index e81627d..9bc9ecc 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise an invalid partial acc_delete on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c index 031c731..a24916d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise an invalid acc_present_or_create on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c index de5d1c1..30b90d4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device on unmapped data on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c index 0d593f0..5db2912 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device with a NULL data address on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c index e98ecc4..8bbf016 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_device with size zero data on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c index f26fc33..c214042 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_self with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c index 253ce59..afa137f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_update_self with a size zero data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c index cfbb077..25c70c2 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c index 5de376d..a8ee7df 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c index 3e621c3..fc221f4 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c @@ -1,4 +1,6 @@ -/* { dg-do run } */ +/* Exercise acc_map_data with data size of zero on nvidia targets. */ + +/* { dg-do run { target openacc_nvidia_accel_selected } } */ #include #include