6e232ba424
2015-11-12 James Norris <jnorris@codesourcery.com> Joseph Myers <joseph@codesourcery.com> gcc/c-family/ * c-pragma.c (oacc_pragmas): Add entry for declare directive. * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE. (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OACC_CLAUSE_LINK. gcc/c/ * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_DECLARE. (c_parser_omp_clause_name): Handle 'device_resident' clause. (c_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OMP_CLAUSE_LINK. (c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OACC_CLAUSE_LINK. (OACC_DECLARE_CLAUSE_MASK): New definition. (c_parser_oacc_declare): New function. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Handle 'device_resident' clause. (cp_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OMP_CLAUSE_LINK. (cp_paser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and PRAGMA_OMP_CLAUSE_LINK. (OACC_DECLARE_CLAUSE_MASK): New definition. (cp_parser_oacc_declare): New function. (cp_parser_pragma): Handle PRAGMA_OACC_DECLARE. * pt.c (tsubst_expr): Handle OACC_DECLARE. gcc/ * gimple-pretty-print.c (dump_gimple_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DECLARE. (is_gomple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. * gimplify.c (oacc_declare_returns): New. (gimplify_bind_expr): Prepend 'exit' stmt to cleanup. (device_resident_p): New function. (oacc_default_clause): Handle device_resident clause. (gimplify_oacc_declare_1, gimplify_oacc_declare): New functions. (gimplify_expr): Handle OACC_DECLARE. * omp-builtins.def (BUILT_IN_GOACC_DECLARE): New builtin. * omp-low.c (expand_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE and BUILTIN_GOACC_DECLARE. (build_omp_regions_1): Handlde GF_OMP_TARGET_KIND_OACC_DECLARE. (lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE, GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK. (make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_LINK and GOMP_MAP_DEVICE_RESIDENT. gcc/testsuite * c-c++-common/goacc/declare-1.c: New test. * c-c++-common/goacc/declare-2.c: Likewise. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK. libgomp/ * libgomp.map (GOACC_2.0.1): Export GOACC_declare. * oacc-parallel.c (GOACC_declare): New function. * testsuite/libgomp.oacc-c-c++-common/declare-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-5.c: Likewise. * testsuite/libgomp.oacc-c++/declare-1.C: Likewise. Co-Authored-By: Joseph Myers <joseph@codesourcery.com> From-SVN: r230275
32 lines
329 B
C
32 lines
329 B
C
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
|
|
|
#include <stdlib.h>
|
|
|
|
template<class T>
|
|
T foo()
|
|
{
|
|
T a, b;
|
|
#pragma acc declare create (a)
|
|
|
|
#pragma acc parallel copyout (b)
|
|
{
|
|
a = 5;
|
|
b = a;
|
|
}
|
|
|
|
return b;
|
|
}
|
|
|
|
int
|
|
main (void)
|
|
{
|
|
int rc;
|
|
|
|
rc = foo<int>();
|
|
|
|
if (rc != 5)
|
|
abort ();
|
|
|
|
return 0;
|
|
}
|