e46c777050
gcc/c/ * c-parser.c (c_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTER. * c-typeck.c (handle_omp_array_sections_1): Replace bool is_omp argument with enum c_omp_region_type ort. (handle_omp_array_sections): Likewise. Update call to handle_omp_array_sections_1. (c_finish_omp_clauses): Add specific errors and warning messages for OpenACC. Use firsrtprivate pointers for OpenACC subarrays. Update call to handle_omp_array_sections. gcc/cp/ * parser.c (cp_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTER. * semantics.c (handle_omp_array_sections_1): Replace bool is_omp argument with enum c_omp_region_type ort. Don't privatize OpenACC non-static members. (handle_omp_array_sections): Replace bool is_omp argument with enum c_omp_region_type ort. Update call to handle_omp_array_sections_1. (finish_omp_clauses): Add specific errors and warning messages for OpenACC. Use firsrtprivate pointers for OpenACC subarrays. Update call to handle_omp_array_sections. 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. (gimplify_adjust_omp_clauses): Fix typo in comment. 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. 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. From-SVN: r236678
899 lines
15 KiB
C
899 lines
15 KiB
C
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
|
/* { dg-additional-options "-lcuda" } */
|
|
|
|
#include <openacc.h>
|
|
#include <stdlib.h>
|
|
#include "cuda.h"
|
|
|
|
#include <stdio.h>
|
|
#include <sys/time.h>
|
|
|
|
int
|
|
main (int argc, char **argv)
|
|
{
|
|
CUresult r;
|
|
CUstream stream1;
|
|
int N = 128; //1024 * 1024;
|
|
float *a, *b, *c, *d, *e;
|
|
int i;
|
|
int nbytes;
|
|
|
|
acc_init (acc_device_nvidia);
|
|
|
|
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;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = a[ii];
|
|
}
|
|
|
|
#pragma acc wait
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 3.0)
|
|
abort ();
|
|
|
|
if (b[i] != 3.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 2.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 2.0)
|
|
abort ();
|
|
|
|
if (b[i] != 2.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 3.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
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;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc parallel wait (1) async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
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 ();
|
|
}
|
|
|
|
|
|
r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING);
|
|
if (r != CUDA_SUCCESS)
|
|
{
|
|
fprintf (stderr, "cuStreamCreate failed: %d\n", r);
|
|
abort ();
|
|
}
|
|
|
|
acc_set_cuda_stream (1, stream1);
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 5.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N], b[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 5.0)
|
|
abort ();
|
|
|
|
if (b[i] != 5.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 7.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 7.0)
|
|
abort ();
|
|
|
|
if (b[i] != 49.0)
|
|
abort ();
|
|
|
|
if (c[i] != 4.0)
|
|
abort ();
|
|
|
|
if (d[i] != 1.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 3.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
e[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc parallel wait (1) async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
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 ();
|
|
|
|
if (e[i] != 17.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 4.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
e[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc update host (a[0:N], b[0:N], c[0:N]) wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 4.0)
|
|
abort ();
|
|
|
|
if (b[i] != 16.0)
|
|
abort ();
|
|
|
|
if (c[i] != 4.0)
|
|
abort ();
|
|
}
|
|
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 5.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
e[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc parallel async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc update host (a[0:N], b[0:N], c[0:N]) async (1)
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 5.0)
|
|
abort ();
|
|
|
|
if (b[i] != 25.0)
|
|
abort ();
|
|
|
|
if (c[i] != 4.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 3.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = a[ii];
|
|
}
|
|
|
|
#pragma acc wait
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 3.0)
|
|
abort ();
|
|
|
|
if (b[i] != 3.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 2.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 2.0)
|
|
abort ();
|
|
|
|
if (b[i] != 2.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 3.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
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;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc kernels wait (1) async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
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 ();
|
|
}
|
|
|
|
|
|
r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING);
|
|
if (r != CUDA_SUCCESS)
|
|
{
|
|
fprintf (stderr, "cuStreamCreate failed: %d\n", r);
|
|
abort ();
|
|
}
|
|
|
|
acc_set_cuda_stream (1, stream1);
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 5.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N], b[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 5.0)
|
|
abort ();
|
|
|
|
if (b[i] != 5.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 7.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 7.0)
|
|
abort ();
|
|
|
|
if (b[i] != 49.0)
|
|
abort ();
|
|
|
|
if (c[i] != 4.0)
|
|
abort ();
|
|
|
|
if (d[i] != 1.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 3.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
e[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
|
|
}
|
|
|
|
#pragma acc kernels wait (1) async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
|
|
}
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
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 ();
|
|
|
|
if (e[i] != 17.0)
|
|
abort ();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 4.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
e[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc update host (a[0:N], b[0:N], c[0:N]) wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 4.0)
|
|
abort ();
|
|
|
|
if (b[i] != 16.0)
|
|
abort ();
|
|
|
|
if (c[i] != 4.0)
|
|
abort ();
|
|
}
|
|
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 5.0;
|
|
b[i] = 0.0;
|
|
c[i] = 0.0;
|
|
d[i] = 0.0;
|
|
e[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N)
|
|
{
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc kernels async (1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
|
|
}
|
|
|
|
#pragma acc update host (a[0:N], b[0:N], c[0:N]) async (1)
|
|
|
|
#pragma acc wait (1)
|
|
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 5.0)
|
|
abort ();
|
|
|
|
if (b[i] != 25.0)
|
|
abort ();
|
|
|
|
if (c[i] != 4.0)
|
|
abort ();
|
|
}
|
|
|
|
acc_shutdown (acc_device_nvidia);
|
|
|
|
return 0;
|
|
}
|