398e3feb8a
* tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description. Add OMP_CLAUSE_USE_DEVICE_ADDR clause. * tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries for OMP_CLAUSE_USE_DEVICE_ADDR clause. (walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Likewise. * gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Likewise. * omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise. Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR clause with array or reference to array types, no matter what type except for reference it has. gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. Set PRAGMA_OACC_CLAUSE_USE_DEVICE equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate enumeration value. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause. (c_parser_omp_clause_use_device_addr): New function. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no map or use_device_* clauses. * c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR in OpenMP, require pointer type rather than pointer or array type. Handle OMP_CLAUSE_USE_DEVICE_ADDR. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause. (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. (cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no map or use_device_* clauses. * semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR in OpenMP, require pointer or reference to pointer type rather than pointer or array or reference to pointer or array type. Handle OMP_CLAUSE_USE_DEVICE_ADDR. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR. gcc/testsuite/ * c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause instead of use_device_ptr clause where required by OpenMP 5.0, add further tests for both use_device_ptr and use_device_addr clauses. libgomp/ * testsuite/libgomp.c/target-18.c (struct S): New type. (foo): Use use_device_addr clause instead of use_device_ptr clause where required by OpenMP 5.0, add further tests for both use_device_ptr and use_device_addr clauses. * testsuite/libgomp.c++/target-9.C (struct S): New type. (foo): Use use_device_addr clause instead of use_device_ptr clause where required by OpenMP 5.0, add further tests for both use_device_ptr and use_device_addr clauses. Add t and u arguments. (main): Adjust caller. From-SVN: r274159
77 lines
1.5 KiB
C
77 lines
1.5 KiB
C
extern void abort (void);
|
|
struct S { int e, f; };
|
|
|
|
void
|
|
foo (int n)
|
|
{
|
|
int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
|
|
struct S d = { 5, 6 };
|
|
int *p = a + 1, i, err;
|
|
for (i = 0; i < n; i++)
|
|
b[i] = 9 + i;
|
|
#pragma omp target data map(to:a)
|
|
#pragma omp target data use_device_ptr(p) map(from:err)
|
|
#pragma omp target is_device_ptr(p) private(i) map(from:err)
|
|
{
|
|
err = 0;
|
|
for (i = 0; i < 4; i++)
|
|
if (p[i - 1] != i)
|
|
err = 1;
|
|
}
|
|
if (err)
|
|
abort ();
|
|
for (i = 0; i < 4; i++)
|
|
a[i] = 23 + i;
|
|
#pragma omp target data map(to:a)
|
|
#pragma omp target data use_device_addr(a) map(from:err)
|
|
#pragma omp target is_device_ptr(a) private(i) map(from:err)
|
|
{
|
|
err = 0;
|
|
for (i = 0; i < 4; i++)
|
|
if (a[i] != 23 + i)
|
|
err = 1;
|
|
}
|
|
if (err)
|
|
abort ();
|
|
#pragma omp target data map(to:b)
|
|
#pragma omp target data use_device_addr(b) map(from:err)
|
|
#pragma omp target is_device_ptr(b) private(i) map(from:err)
|
|
{
|
|
err = 0;
|
|
for (i = 0; i < 4; i++)
|
|
if (b[i] != 9 + i)
|
|
err = 1;
|
|
}
|
|
if (err)
|
|
abort ();
|
|
#pragma omp target data map(to:c)
|
|
#pragma omp target data use_device_addr(c) map(from:err)
|
|
{
|
|
int *q = &c;
|
|
#pragma omp target is_device_ptr(q) map(from:err)
|
|
{
|
|
err = *q != 4;
|
|
}
|
|
}
|
|
if (err)
|
|
abort ();
|
|
#pragma omp target data map(to:d)
|
|
#pragma omp target data use_device_addr(d) map(from:err)
|
|
{
|
|
struct S *r = &d;
|
|
#pragma omp target is_device_ptr(r) map(from:err)
|
|
{
|
|
err = r->e != 5 || r->f != 6;
|
|
}
|
|
}
|
|
if (err)
|
|
abort ();
|
|
}
|
|
|
|
int
|
|
main ()
|
|
{
|
|
foo (9);
|
|
return 0;
|
|
}
|