From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 62515 invoked by alias); 22 Oct 2015 19:20:32 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 62506 invoked by uid 89); 22 Oct 2015 19:20:31 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_LOW,SPF_PASS,T_FROM_12LTRDOM autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 22 Oct 2015 19:20:30 +0000 Received: from svr-orw-fem-02x.mgc.mentorg.com ([147.34.96.206] helo=SVR-ORW-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZpLPX-0004W4-8a from James_Norris@mentor.com ; Thu, 22 Oct 2015 12:20:27 -0700 Received: from [172.30.80.51] (147.34.91.1) by svr-orw-fem-02.mgc.mentorg.com (147.34.96.168) with Microsoft SMTP Server id 14.3.224.2; Thu, 22 Oct 2015 12:20:26 -0700 Message-ID: <562936F8.8070302@codesourcery.com> Date: Thu, 22 Oct 2015 19:22:00 -0000 From: James Norris User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.7.0 MIME-Version: 1.0 To: GCC Patches CC: "Joseph S. Myers" , Nathan Sidwell , Jakub Jelinek Subject: [OpenACC 7/7] host_data construct (runtime tests) References: <56293476.5020801@codesourcery.com> In-Reply-To: <56293476.5020801@codesourcery.com> X-TagToolbar-Keys: D20151022142024378 Content-Type: multipart/mixed; boundary="------------030305070503090505030203" X-SW-Source: 2015-10/txt/msg02339.txt.bz2 --------------030305070503090505030203 Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 295 libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c - New runtime tests for host_data. --------------030305070503090505030203 Content-Type: text/x-patch; name="p7.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="p7.patch" Content-length: 4123 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c new file mode 100644 index 0000000..15ccb27 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c @@ -0,0 +1,125 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-lcuda -lcublas -lcudart" } */ + +#include +#include +#include +#include +#include +#include + +void +saxpy_host (int n, float a, float *x, float *y) +{ + int i; + + for (i = 0; i < n; i++) + y[i] = y[i] + a * x[i]; +} + +#pragma acc routine +void +saxpy_target (int n, float a, float *x, float *y) +{ + int i; + + for (i = 0; i < n; i++) + y[i] = y[i] + a * x[i]; +} + +int +main(int argc, char **argv) +{ + const int N = 8; + int i; + float *x_ref, *y_ref; + float *x, *y; + cublasHandle_t h; + float a = 2.0; + + x_ref = (float*) malloc (N * sizeof(float)); + y_ref = (float*) malloc (N * sizeof(float)); + + x = (float*) malloc (N * sizeof(float)); + y = (float*) malloc (N * sizeof(float)); + +#pragma acc data copyin (x[0:N]) copy (y[0:N]) + { + float *xp, *yp; +#pragma acc host_data use_device (x, y) + { +#pragma acc parallel pcopy (xp, yp) present (x, y) + { + xp = x; + yp = y; + } + } + + if (xp != acc_deviceptr (x) || yp != acc_deviceptr (y)) + abort (); + } + + for (i = 0; i < N; i++) + { + x[i] = x_ref[i] = 4.0 + i; + y[i] = y_ref[i] = 3.0; + } + + saxpy_host (N, a, x_ref, y_ref); + + cublasCreate (&h); + +#pragma acc data copyin (x[0:N]) copy (y[0:N]) + { +#pragma acc host_data use_device (x, y) + { + cublasSaxpy (h, N, &a, x, 1, y, 1); + } + } + + for (i = 0; i < N; i++) + { + if (y[i] != y_ref[i]) + abort (); + } + +#pragma acc data create (x[0:N]) copyout (y[0:N]) + { +#pragma acc kernels + for (i = 0; i < N; i++) + y[i] = 3.0; + +#pragma acc host_data use_device (x, y) + { + cublasSaxpy (h, N, &a, x, 1, y, 1); + } + } + + cublasDestroy (h); + + for (i = 0; i < N; i++) + { + if (y[i] != y_ref[i]) + abort (); + } + + for (i = 0; i < N; i++) + y[i] = 3.0; + +#pragma acc data copyin (x[0:N]) copyin (a, N) copy (y[0:N]) + { +#pragma acc host_data use_device (x, y) + { +#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a, N) + saxpy_target (N, a, x, y); + } + } + + for (i = 0; i < N; i++) + { + if (y[i] != y_ref[i]) + abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c new file mode 100644 index 0000000..511ec64 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c @@ -0,0 +1,50 @@ +/* { dg-do run } */ + +#include + +struct by_lightning { + int a; + int b; + int c; +}; + +int main (int argc, char* argv[]) +{ + int x; + void *q = NULL, *r = NULL, *p = NULL, *s = NULL, *t = NULL; + long u; + struct by_lightning on_the_head = {1, 2, 3}; + int arr[10], *f = NULL; + _Complex float cf; + #pragma acc enter data copyin (x, arr, on_the_head, cf) + #pragma acc host_data use_device (x, arr, on_the_head, cf) + { + q = &x; + { + f = &arr[5]; + r = f; + s = &__real__ cf; + t = &on_the_head.c; + u = (long) &__imag__ cf; + #pragma acc parallel copyout(p) present (x, arr, on_the_head, cf) + { + /* This will not (and must not) call GOACC_deviceptr, but '&x' will be + the address on the device (if appropriate) regardless. */ + p = &x; + } + } + } + #pragma acc exit data delete (x) + +#if ACC_MEM_SHARED + if (q != &x || f != &arr[5] || r != f || s != &(__real__ cf) + || t != &on_the_head.c || u != (long) &(__imag__ cf) || p != &x) + abort (); +#else + if (q == &x || f == &arr[5] || r != f || s == &(__real__ cf) + || t == &on_the_head.c || u == (long) &(__imag__ cf) || p == &x) + abort (); +#endif + + return 0; +} --------------030305070503090505030203--