Mercurial > hg > CbC > CbC_gcc
view libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c @ 144:8f4e72ab4e11
fix segmentation fault caused by nothing next cur_op to end
author | Takahiro SHIMIZU <anatofuz@cr.ie.u-ryukyu.ac.jp> |
---|---|
date | Sun, 23 Dec 2018 21:23:56 +0900 |
parents | 84e7813d76e9 |
children |
line wrap: on
line source
/* { dg-do run } */ /* { dg-additional-options "-lcuda" { target openacc_nvidia_accel_selected } } */ #include <openacc.h> #include <stdlib.h> #if defined ACC_DEVICE_TYPE_nvidia #include "cuda.h" #endif #include <stdio.h> #include <sys/time.h> int main (int argc, char **argv) { #if defined ACC_DEVICE_TYPE_nvidia CUresult r; CUstream stream1; #endif int N = 128; //1024 * 1024; float *a, *b, *c, *d, *e; int i; int nbytes; #if defined ACC_DEVICE_TYPE_nvidia acc_init (acc_device_nvidia); #endif 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 (); } #if defined ACC_DEVICE_TYPE_nvidia 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); #endif 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 (); } #if defined ACC_DEVICE_TYPE_nvidia 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); #endif 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 (); } #if defined ACC_DEVICE_TYPE_nvidia acc_shutdown (acc_device_nvidia); #endif return 0; }