a6163563f2
The clause makes any device code use the local memory address for each
of the variables specified unless the given variable is already present
on the current device.
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
gcc/
* omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC.
* tree-pretty-print.c (dump_omp_clause): Likewise.
gcc/c-family/
* c-pragma.h (pragma_omp_clause): Add
PRAGMA_OACC_CLAUSE_NO_CREATE.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Support no_create.
(c_parser_oacc_data_clause): Likewise.
(c_parser_oacc_all_clauses): Likewise.
(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
(OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add
PRAGMA_OACC_CLAUSE_NO_CREATE.
* c-typeck.c (handle_omp_array_sections): Support
GOMP_MAP_NO_ALLOC.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Support no_create.
(cp_parser_oacc_data_clause): Likewise.
(cp_parser_oacc_all_clauses): Likewise.
(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
(OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE.
* semantics.c (handle_omp_array_sections): Support no_create.
gcc/fortran/
* gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC.
* openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE.
(gfc_match_omp_clauses): Support no_create.
(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES)
(OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
* trans-openmp.c (gfc_trans_omp_clauses_1): Support
OMP_MAP_NO_ALLOC.
gcc/testsuite/
* gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
* gfortran.dg/goacc/common-block-1.f90: Likewise.
* gfortran.dg/goacc/data-clauses.f95: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
include/
* gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC.
libgomp/
* target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC.
* testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test.
* testsuite/libgomp.oacc-fortran/no_create-1.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-2.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-3.F90: New test.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Maciej W. Rozycki <macro@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
From-SVN: r279551
50 lines
1.0 KiB
C
50 lines
1.0 KiB
C
/* Test 'no_create' clause on 'data' construct and nested compute construct,
|
|
with data not present on the device. */
|
|
|
|
#include <stdlib.h>
|
|
#include <stdio.h>
|
|
#include <openacc.h>
|
|
|
|
#define N 128
|
|
|
|
int
|
|
main (int argc, char *argv[])
|
|
{
|
|
int var;
|
|
int *arr = (int *) malloc (N * sizeof (*arr));
|
|
int *devptr[2];
|
|
|
|
#pragma acc data no_create(var, arr[0:N])
|
|
{
|
|
devptr[0] = (int *) acc_deviceptr (&var);
|
|
devptr[1] = (int *) acc_deviceptr (&arr[2]);
|
|
|
|
#if ACC_MEM_SHARED
|
|
if (devptr[0] == NULL)
|
|
__builtin_abort ();
|
|
if (devptr[1] == NULL)
|
|
__builtin_abort ();
|
|
#else
|
|
if (devptr[0] != NULL)
|
|
__builtin_abort ();
|
|
if (devptr[1] != NULL)
|
|
__builtin_abort ();
|
|
#endif
|
|
|
|
#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
|
|
{
|
|
devptr[0] = &var;
|
|
devptr[1] = &arr[2];
|
|
}
|
|
|
|
if (devptr[0] != &var)
|
|
__builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } }
|
|
if (devptr[1] != &arr[2])
|
|
__builtin_abort ();
|
|
}
|
|
|
|
free (arr);
|
|
|
|
return 0;
|
|
}
|