From patchwork Sat Jan 12 22:21:23 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024005 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493946-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="hcoMItbS"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ2h6lV1z9sCh for ; Sun, 13 Jan 2019 09:21:32 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=w1EpKL4HCemxuHzx5vn/3vYPRWEPwGPpUFadIZTNJtjLcOVZGCGo1 RbS03IyoaLOeYvKLQATtzs3LXiB4kTJFdczoO6jVGyO3TsTQBJpOVbWeDn9Ui8kx lrFW67tl5Fyad4wnrZ1IQcbzw5QgCwNJXfjsM/DP0MPwGT8sxDkip0= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=2IkhLMDWJw0Pn0/jFbGBPDYfD5A=; b=hcoMItbSbD1n6Vdy3sC4 6TvMpNV0pzhYvIU0RkZ+Agup+Cj+c1Eke1HyGbR95eUITIXD+5dvPy2P5s0Xt97L fKKEjCvkq3WgU11Vtk9mEjoTfCNSqbWfrn5y0SrqBo8lnL0yyUfYwcbTx1dk1TjS N8e41yMX8R4RtpBBa3inc7o= Received: (qmail 103920 invoked by alias); 12 Jan 2019 22:21:13 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 103826 invoked by uid 89); 12 Jan 2019 22:21:13 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy=Hx-languages-length:3293 X-HELO: mx1.suse.de Received: from mx2.suse.de (HELO mx1.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:11 +0000 Received: from relay1.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id 25692AC15; Sat, 12 Jan 2019 22:21:09 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 1/9] [nvptx] Enable large vectors Date: Sat, 12 Jan 2019 23:21:23 +0100 Message-Id: <20190112222131.29519-2-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-IsSubscribed: yes Allow vector_length clauses to accept values larger than warp size. Note that this does not enable setting vector_length to values larger than warp size using -fopenacc-dim. 2018-12-17 Tom de Vries * config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector lengths into account. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect vector length to be 128. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector length 2097152 to be reduced to 1024 instead of 32. --- gcc/config/nvptx/nvptx.c | 2 +- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c | 4 ++-- libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c | 5 ++--- 3 files changed, 5 insertions(+), 6 deletions(-) diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 1d9704543d9..8d2740cd50f 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -96,7 +96,7 @@ #define PTX_NUM_PER_WORKER_BARRIERS (PTX_CTA_NUM_BARRIERS - PTX_NUM_PER_CTA_BARRIERS) #define PTX_DEFAULT_VECTOR_LENGTH PTX_WARP_SIZE -#define PTX_MAX_VECTOR_LENGTH PTX_WARP_SIZE +#define PTX_MAX_VECTOR_LENGTH PTX_CTA_SIZE #define PTX_WORKER_LENGTH 32 #define PTX_DEFAULT_RUNTIME_DIM 0 /* Defer to runtime. */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c index 4a9854662cc..d7cd0461b53 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c @@ -350,7 +350,7 @@ int main () int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max; gangs_min = workers_min = vectors_min = INT_MAX; gangs_max = workers_max = vectors_max = INT_MIN; -#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 2097152" "" { target openacc_nvidia_accel_configured } } */ \ +#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(1024\\), ignoring 2097152" "" { target openacc_nvidia_accel_configured } } */ \ vector_length (VECTORS) { if (acc_on_device (acc_device_host)) @@ -361,7 +361,7 @@ int main () else if (acc_on_device (acc_device_nvidia)) { /* The GCC nvptx back end enforces vector_length (32). */ - vectors_actual = 32; + vectors_actual = 1024; } else __builtin_abort (); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c index fab5b0d25d1..18d77cc5ecb 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c @@ -33,7 +33,6 @@ main (void) return 0; } -/* { dg-prune-output "using vector_length \\(32\\), ignoring 128" } */ -/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */ -/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */ +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */ From patchwork Sat Jan 12 22:21:24 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024004 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493945-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="qbC/qVlP"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ2W4qtCz9s7T for ; Sun, 13 Jan 2019 09:21:20 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=q5ImsYtynbQ27+DQITKrHQrc4QdglvbcjYxf3x0cAfnMMGMkJV2lr lzfXt9d2VKqftiV2UPj18JM8xl0HXvd0Ya6jumgshwP6pdxDRIyRO718y9sbWjSU h9t7BslV0AKwD5p5oiuq6pHeGil3nAnX93NPx59cFJ+1XGqDtn+siY= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=xfxliSUocoNWahKj88SdKuWnmPU=; b=qbC/qVlPdjO6auD+fkll hBOlkPESHaZ2AjRhALtIC9VZszJ1lrlZBAKSm7YwL8/3ehlrBYj7oAcUNiSJk3XF zmaCR3QLlv+tAp7FDQQoL9NUNUvL/YG6+pEWxMf/nBjP9KBiPLL6z8oC1IkQGCQi 8r0TpcsaS7xme28lw5FHNw0= Received: (qmail 103824 invoked by alias); 12 Jan 2019 22:21:12 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 103814 invoked by uid 89); 12 Jan 2019 22:21:12 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy= X-HELO: mx1.suse.de Received: from mx2.suse.de (HELO mx1.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:11 +0000 Received: from relay2.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id 5CAF8AC66; Sat, 12 Jan 2019 22:21:09 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 2/9] [nvptx] Update insufficient launch message for variable vector_length Date: Sat, 12 Jan 2019 23:21:24 +0100 Message-Id: <20190112222131.29519-3-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-IsSubscribed: yes Update message in nvptx libgomp plugin about insufficient resources to launch kernel, to accommodate for the fact the vector_length can now be variable. 19-01-08 Tom de Vries * plugin/plugin-nvptx.c (nvptx_exec): Update insufficient hardware resources diagnostic. --- libgomp/plugin/plugin-nvptx.c | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index c80da64c422..8912660966a 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -1297,14 +1297,16 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs, if (dims[GOMP_DIM_WORKER] * dims[GOMP_DIM_VECTOR] > targ_fn->max_threads_per_block) { - int suggest_workers - = targ_fn->max_threads_per_block / dims[GOMP_DIM_VECTOR]; - GOMP_PLUGIN_fatal ("The Nvidia accelerator has insufficient resources to" - " launch '%s' with num_workers = %d; recompile the" - " program with 'num_workers = %d' on that offloaded" - " region or '-fopenacc-dim=:%d'", - targ_fn->launch->fn, dims[GOMP_DIM_WORKER], - suggest_workers, suggest_workers); + const char *msg + = ("The Nvidia accelerator has insufficient resources to launch '%s'" + " with num_workers = %d and vector_length = %d" + "; " + "recompile the program with 'num_workers = x and vector_length = y'" + " on that offloaded region or '-fopenacc-dim=:x:y' where" + " x * y <= %d" + ".\n"); + GOMP_PLUGIN_fatal (msg, targ_fn->launch->fn, dims[GOMP_DIM_WORKER], + dims[GOMP_DIM_VECTOR], targ_fn->max_threads_per_block); } /* Check if the accelerator has sufficient barrier resources to From patchwork Sat Jan 12 22:21:25 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024006 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493947-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="wGZGZ1r7"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ2x4M2fz9s7T for ; Sun, 13 Jan 2019 09:21:45 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=vmK+whUNoJOR8unkiNUA0Xvv9qF7YzEhqo/uAVXxAF46MTEvslXjF nLau4tgd5P7v6ye7VCtR3Y/ZjgirTD8u5RgbAqmcO9RpiKYwCWXOHaqZVfn2/B2i jbW/iO98rxU0B2C3Q3Nj2KnXNaR/scHN5jtgkGi94SqGJ4bPbNILxU= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=sHGmed2fYxt9h9aovYgdDTXaSDQ=; b=wGZGZ1r7yZ+OEX5xhppW 5W0gNrzA9F5d3TttUsNfYv0XckMRt9pHpy91BBi8EIvDKVYI69z3WdiMY866UX2+ vF5/s6YQZFCpPjfCBR8Hmu9BKc1RitRgaYnWfWJw+ic/9/DnNockH3H2eX5MVguO wteYIpOuVTfgkrVL3vTXBG8= Received: (qmail 104147 invoked by alias); 12 Jan 2019 22:21:14 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 104062 invoked by uid 89); 12 Jan 2019 22:21:14 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy=foffload X-HELO: mx1.suse.de Received: from mx2.suse.de (HELO mx1.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:11 +0000 Received: from relay1.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id 9C7EAACEF; Sat, 12 Jan 2019 22:21:09 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 3/9] [nvptx] Enable large vectors -- test-cases Date: Sat, 12 Jan 2019 23:21:25 +0100 Message-Id: <20190112222131.29519-4-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-IsSubscribed: yes Add various test-cases with vector length 128. 2018-12-17 Tom de Vries * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test. --- .../vector-length-128-4.c | 40 +++++++++++++++++++++ .../vector-length-128-6.c | 41 ++++++++++++++++++++++ .../vector-length-128-7.c | 40 +++++++++++++++++++++ 3 files changed, 121 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c new file mode 100644 index 00000000000..e5d1df09b8a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c @@ -0,0 +1,40 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ + +#include + +#define N 1024 + +unsigned int a[N]; +unsigned int b[N]; +unsigned int c[N]; +unsigned int n = N; + +int +main (void) +{ + for (unsigned int i = 0; i < n; ++i) + { + a[i] = i % 3; + b[i] = i % 5; + } + +#pragma acc parallel num_workers (2) vector_length (128) copyin (a,b) copyout (c) + { +#pragma acc loop worker + for (unsigned int i = 0; i < 4; i++) +#pragma acc loop vector + for (unsigned int j = 0; j < n / 4; j++) + c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j]; + } + + for (unsigned int i = 0; i < n; ++i) + if (c[i] != (i % 3) + (i % 5)) + abort (); + + return 0; +} + +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c new file mode 100644 index 00000000000..a1f67622f84 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c @@ -0,0 +1,41 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-set-target-env-var "GOMP_OPENACC_DIM" ":2:" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ + +#include + +#define N 1024 + +unsigned int a[N]; +unsigned int b[N]; +unsigned int c[N]; +unsigned int n = N; + +int +main (void) +{ + for (unsigned int i = 0; i < n; ++i) + { + a[i] = i % 3; + b[i] = i % 5; + } + +#pragma acc parallel vector_length (128) copyin (a,b) copyout (c) + { +#pragma acc loop worker + for (unsigned int i = 0; i < 4; i++) +#pragma acc loop vector + for (unsigned int j = 0; j < n / 4; j++) + c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j]; + } + + for (unsigned int i = 0; i < n; ++i) + if (c[i] != (i % 3) + (i % 5)) + abort (); + + return 0; +} + +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c new file mode 100644 index 00000000000..c419f6499b5 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c @@ -0,0 +1,40 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ + +#include + +#define N 1024 + +unsigned int a[N]; +unsigned int b[N]; +unsigned int c[N]; +unsigned int n = N; + +int +main (void) +{ + for (unsigned int i = 0; i < n; ++i) + { + a[i] = i % 3; + b[i] = i % 5; + } + +#pragma acc parallel vector_length (128) copyin (a,b) copyout (c) + { +#pragma acc loop worker + for (unsigned int i = 0; i < 4; i++) +#pragma acc loop vector + for (unsigned int j = 0; j < n / 4; j++) + c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j]; + } + + for (unsigned int i = 0; i < n; ++i) + if (c[i] != (i % 3) + (i % 5)) + abort (); + + return 0; +} + +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=8, vectors=128" } */ From patchwork Sat Jan 12 22:21:26 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024012 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493953-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="OhJbe4iv"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ4L48XZz9s7T for ; Sun, 13 Jan 2019 09:22:58 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=P7FJght4OtnUtBb8XtBldpwK6Lxd0JHqvgh2ogVj4v5kl7e18mvwL 93v8TaZSJDxh1hYgrQGQPVn8vYMzhBvGeu/XsnYgTXxOFD2MaGQ21pMzstriCvGJ AlMw3pgMTRJasf1Pw8qmiv+Y3XbNGknnCxQxh5ahzp0PzVMILQSfDo= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=CIYetiG2L4OOaZ3/MHYeF7iG0sg=; b=OhJbe4ivXqPC1ew+dBhB JC/BobullP6GhOpicJ8MMsYVwKmVFAo66BKDjGcaTlGxczqlPykktKckE+xjWRLt rN4x7Pwi5yMCAM38IdnIF/Sv6/D9XRsEvJe57m5o/hRc6TKQKUo8qxhQPJ49cSvs mZjy3KSdz8ZsCdYU2lvGDP0= Received: (qmail 104958 invoked by alias); 12 Jan 2019 22:21:19 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 104882 invoked by uid 89); 12 Jan 2019 22:21:19 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.2 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_SOFTFAIL autolearn=ham version=3.3.2 spammy=HX-detected-operating-system:timestamps, m*, exercise, ij X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (209.51.188.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:15 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1giReS-0007E1-U8 for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:13 -0500 Received: from mx2.suse.de ([195.135.220.15]:52480 helo=mx1.suse.de) by eggs.gnu.org with esmtps (TLS1.0:DHE_RSA_AES_256_CBC_SHA1:32) (Exim 4.71) (envelope-from ) id 1giReS-0007Bq-LD for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:12 -0500 Received: from relay2.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id E2D64AD48; Sat, 12 Jan 2019 22:21:09 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 4/9] [nvptx] Enable large vectors -- reduction testcases Date: Sat, 12 Jan 2019 23:21:26 +0100 Message-Id: <20190112222131.29519-5-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-detected-operating-system: by eggs.gnu.org: GNU/Linux 2.2.x-3.x (no timestamps) [generic] X-Received-From: 195.135.220.15 X-IsSubscribed: yes Add various reduction test-cases with vector length 128. 2018-12-17 Tom de Vries * testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: New test. * testsuite/libgomp.oacc-fortran/gemm.f90: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c: New test. --- .../vector-length-128-10.c | 39 +++++++++++ .../libgomp.oacc-c-c++-common/vred2d-128.c | 55 +++++++++++++++ libgomp/testsuite/libgomp.oacc-fortran/gemm.f90 | 79 ++++++++++++++++++++++ 3 files changed, 173 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vred2d-128.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/gemm.f90 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c new file mode 100644 index 00000000000..0658cfde7ad --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c @@ -0,0 +1,39 @@ +/* { dg-do run } */ + +#include + +#define N 1024 + +unsigned int a[N]; +unsigned int b[N]; +unsigned int c[N]; +unsigned int n = N; + +int +main (void) +{ + for (unsigned int i = 0; i < n; ++i) + { + a[i] = i % 3; + b[i] = i % 5; + } + + unsigned int res = 1; + unsigned long long res2 = 1; +#pragma acc parallel vector_length (128) copyin (a,b) reduction (+:res, res2) copy (res, res2) + { +#pragma acc loop vector reduction (+:res, res2) + for (unsigned int i = 0; i < n; i++) + { + res += ((a[i] + b[i]) % 2); + res2 += ((a[i] + b[i]) % 2); + } + } + + if (res != 478) + abort (); + if (res2 != 478) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vred2d-128.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vred2d-128.c new file mode 100644 index 00000000000..86171d456e0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vred2d-128.c @@ -0,0 +1,55 @@ +/* Test large vector lengths. */ + +#include + +#define n 10000 +int a1[n], a2[n]; + +#define gentest(name, outer, inner) \ + void name () \ + { \ + long i, j, t1, t2, t3; \ + _Pragma(outer) \ + for (i = 0; i < n; i++) \ + { \ + t1 = 0; \ + t2 = 0; \ + _Pragma(inner) \ + for (j = i; j < n; j++) \ + { \ + t1++; \ + t2--; \ + } \ + a1[i] = t1; \ + a2[i] = t2; \ + } \ + for (i = 0; i < n; i++) \ + { \ + assert (a1[i] == n-i); \ + assert (a2[i] == -(n-i)); \ + } \ + } \ + +gentest (test1, "acc parallel loop gang vector_length (128) firstprivate (t1, t2)", + "acc loop vector reduction(+:t1) reduction(-:t2)") + +gentest (test2, "acc parallel loop gang vector_length (128) firstprivate (t1, t2)", + "acc loop worker vector reduction(+:t1) reduction(-:t2)") + +gentest (test3, "acc parallel loop gang worker vector_length (128) firstprivate (t1, t2)", + "acc loop vector reduction(+:t1) reduction(-:t2)") + +gentest (test4, "acc parallel loop firstprivate (t1, t2)", + "acc loop reduction(+:t1) reduction(-:t2)") + + +int +main () +{ + test1 (); + test2 (); + test3 (); + test4 (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/gemm.f90 b/libgomp/testsuite/libgomp.oacc-fortran/gemm.f90 new file mode 100644 index 00000000000..de78148c7b3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/gemm.f90 @@ -0,0 +1,79 @@ +! Exercise three levels of parallelism using SGEMM from BLAS. + +! { dg-do run } + +! Explicitly set vector_length to 128 using a vector_length clause. +subroutine openacc_sgemm_128 (m, n, k, alpha, a, b, beta, c) + integer :: m, n, k + real :: alpha, beta + real :: a(k,*), b(k,*), c(m,*) + + integer :: i, j, l + real :: temp + + !$acc parallel loop copy(c(1:m,1:n)) copyin(a(1:k,1:m),b(1:k,1:n)) vector_length (128) firstprivate (temp) + do j = 1, n + !$acc loop + do i = 1, m + temp = 0.0 + !$acc loop reduction(+:temp) + do l = 1, k + temp = temp + a(l,i)*b(l,j) + end do + if(beta == 0.0) then + c(i,j) = alpha*temp + else + c(i,j) = alpha*temp + beta*c(i,j) + end if + end do + end do +end subroutine openacc_sgemm_128 + +subroutine host_sgemm (m, n, k, alpha, a, b, beta, c) + integer :: m, n, k + real :: alpha, beta + real :: a(k,*), b(k,*), c(m,*) + + integer :: i, j, l + real :: temp + + do j = 1, n + do i = 1, m + temp = 0.0 + do l = 1, k + temp = temp + a(l,i)*b(l,j) + end do + if(beta == 0.0) then + c(i,j) = alpha*temp + else + c(i,j) = alpha*temp + beta*c(i,j) + end if + end do + end do +end subroutine host_sgemm + +program main + integer, parameter :: M = 100, N = 50, K = 2000 + real :: a(K, M), b(K, N), c(M, N), d (M, N), e (M, N) + real alpha, beta + integer i, j + + a(:,:) = 1.0 + b(:,:) = 0.25 + + c(:,:) = 0.0 + d(:,:) = 0.0 + e(:,:) = 0.0 + + alpha = 1.05 + beta = 1.25 + + call openacc_sgemm_128 (M, N, K, alpha, a, b, beta, d) + call host_sgemm (M, N, K, alpha, a, b, beta, e) + + do i = 1, m + do j = 1, n + if (d(i,j) /= e(i,j)) call abort + end do + end do +end program main From patchwork Sat Jan 12 22:21:27 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024008 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493949-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="c0aGtwar"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ3Q2j41z9s7T for ; Sun, 13 Jan 2019 09:22:10 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=LXy72wP7b7XyOhZUjW+6w8Ge2/5Ve75QRRKEGiqjYhJFaoV22vv// 5kiWzMw7g9OcJ984hdEddyNM8XajuBc8cIoeIW4OD8y8LB0Yh22miedIWb+M7qzj BlpFMymmorwlhLPbEP0X9k/onJIvuodvwdMS73mxccTstWjIPtMNgg= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=B3f8/BN+IrDrKXCPlBWCnr8cKag=; b=c0aGtwar2SWbgwalnxVk gBBg7SCuOBgHzxiVpkM3EQkGTtv67tExBpvk03eLgGuurBHWSQjCBfIwJpBYxFTQ pnG6r25KZDEx/u6zic3+eNRHJISREZtvdN0hnh2KjE9nY2aIeUAtHBHgH2P1l6b6 YVPoKepwDYuOIyDpwLQ7fmM= Received: (qmail 104483 invoked by alias); 12 Jan 2019 22:21:17 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 104434 invoked by uid 89); 12 Jan 2019 22:21:16 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy= X-HELO: mx1.suse.de Received: from mx2.suse.de (HELO mx1.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:14 +0000 Received: from relay1.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id 27B73ADF2; Sat, 12 Jan 2019 22:21:10 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 5/9] [nvptx] Don't emit barriers for empty loops -- test-cases Date: Sat, 12 Jan 2019 23:21:27 +0100 Message-Id: <20190112222131.29519-6-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-IsSubscribed: yes Add test-cases for PR85381. 2018-12-17 Tom de Vries PR target/85381 * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85381.c: New test. --- .../libgomp.oacc-c-c++-common/pr85381-5.c | 24 ++++++++++++++++++++++ .../testsuite/libgomp.oacc-c-c++-common/pr85381.c | 18 ++++++++++++++++ 2 files changed, 42 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c new file mode 100644 index 00000000000..61e7e48f0c9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-5.c @@ -0,0 +1,24 @@ +/* { dg-additional-options "-save-temps" } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } + { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ + +#define n 1024 + +int +main (void) +{ + #pragma acc parallel vector_length(128) + { + #pragma acc loop vector + for (int i = 0; i < n; i++) + ; + + #pragma acc loop vector + for (int i = 0; i < n; i++) + ; + } + + return 0; +} + +/* { dg-final { scan-assembler-not "bar.sync" } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c new file mode 100644 index 00000000000..2864dfcf3cb --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381.c @@ -0,0 +1,18 @@ +/* { dg-additional-options "-save-temps" } */ +/* { dg-do run { target openacc_nvidia_accel_selected } } + { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */ + +int +main (void) +{ + int v1; + + #pragma acc parallel vector_length (128) + #pragma acc loop vector + for (v1 = 0; v1 < 20; v1 += 2) + ; + + return 0; +} + +/* { dg-final { scan-assembler-not "bar.sync" } } */ From patchwork Sat Jan 12 22:21:28 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024011 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493952-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="WKLQlnRC"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ4602T7z9s7T for ; Sun, 13 Jan 2019 09:22:45 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=HntN1rHGrHCGm+xTOS6SlDnYvyTkEaaWFs2del0hEtLVx9zSjhD/4 aR+kvaDkoDaZ1ZHYn8JEDDbrBTH2qmTKmYA3oVeJKqLaCq+Qk/wT61NvvQnf2QFe guz1gi89XEnNRZqULJjgQtVa1e2c5dcQ1x9VrCBZHSAssQTEqMsDxk= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=ElTgjC2DHc4m1GGTj294dUD20Ng=; b=WKLQlnRCOqJHjqFO6Ukp YyGMmO4FiCMxBhFPnyoVd1DOIrXCYAuYPjT0kJhfg9DSZu1oGAzLlawUJ4UuvlMG tiJEDfigzTUtm2rmI6ELyGCrhQrcUgcMmE3+iqW8/mKBkxq1PM9DBeZPnwl4DD/p Vg2vbvQwp1QWAxHOpY/n4KY= Received: (qmail 104947 invoked by alias); 12 Jan 2019 22:21:19 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 104877 invoked by uid 89); 12 Jan 2019 22:21:19 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.2 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_SOFTFAIL autolearn=ham version=3.3.2 spammy=HX-detected-operating-system:timestamps, err X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (209.51.188.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:16 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1giReU-0007FJ-8U for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:15 -0500 Received: from mx2.suse.de ([195.135.220.15]:52520 helo=mx1.suse.de) by eggs.gnu.org with esmtps (TLS1.0:DHE_RSA_AES_256_CBC_SHA1:32) (Exim 4.71) (envelope-from ) id 1giReT-0007EF-VZ for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:14 -0500 Received: from relay2.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id 68798AE54; Sat, 12 Jan 2019 22:21:10 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 6/9] [nvptx] Force vl32 if calling vector-partitionable routines -- test-cases Date: Sat, 12 Jan 2019 23:21:28 +0100 Message-Id: <20190112222131.29519-7-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-detected-operating-system: by eggs.gnu.org: GNU/Linux 2.2.x-3.x (no timestamps) [generic] X-Received-From: 195.135.220.15 X-IsSubscribed: yes Add test-cases for "[nvptx] Force vl32 if calling vector-partitionable routines". 2018-12-17 Tom de Vries PR target/85486 * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test. --- .../libgomp.oacc-c-c++-common/pr85486-3.c | 54 ++++++++++++++++++++++ .../testsuite/libgomp.oacc-c-c++-common/pr85486.c | 51 ++++++++++++++++++++ 2 files changed, 105 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c new file mode 100644 index 00000000000..a959b90c29a --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c @@ -0,0 +1,54 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */ + +/* Minimized from ref-1.C. */ + +#include + +#pragma acc routine vector +void __attribute__((noinline, noclone)) +Vector (int *ptr, int n, const int inc) +{ + #pragma acc loop vector + for (unsigned ix = 0; ix < n; ix++) + ptr[ix] += inc; +} + +int +main (void) +{ + const int n = 32, m=32; + + int ary[m][n]; + unsigned ix, iy; + + for (ix = m; ix--;) + for (iy = n; iy--;) + ary[ix][iy] = (1 << 16) + (ix << 8) + iy; + + int err = 0; + +#pragma acc parallel copy (ary) + { + Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16)); + } + + for (ix = m; ix--;) + for (iy = n; iy--;) + if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy)) + { + printf ("ary[%u][%u] = %x expected %x\n", + ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy)); + err++; + } + + if (err) + { + printf ("%d failed\n", err); + return 1; + } + + return 0; +} + +/* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c new file mode 100644 index 00000000000..99c08059d37 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c @@ -0,0 +1,51 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ + +/* Minimized from ref-1.C. */ + +#include + +#pragma acc routine vector +void __attribute__((noinline, noclone)) +Vector (int *ptr, int n, const int inc) +{ + #pragma acc loop vector + for (unsigned ix = 0; ix < n; ix++) + ptr[ix] += inc; +} + +int +main (void) +{ + const int n = 32, m=32; + + int ary[m][n]; + unsigned ix, iy; + + for (ix = m; ix--;) + for (iy = n; iy--;) + ary[ix][iy] = (1 << 16) + (ix << 8) + iy; + + int err = 0; + +#pragma acc parallel copy (ary) vector_length (128) /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */ + { + Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16)); + } + + for (ix = m; ix--;) + for (iy = n; iy--;) + if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy)) + { + printf ("ary[%u][%u] = %x expected %x\n", + ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy)); + err++; + } + + if (err) + { + printf ("%d failed\n", err); + return 1; + } + + return 0; +} From patchwork Sat Jan 12 22:21:29 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024009 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493950-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="Q6Ga73j3"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ3f412Cz9s7T for ; Sun, 13 Jan 2019 09:22:22 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=h9AvegNUk7J/fYwaZmfJ4IQLF3274wy2o8TYCh8xpca+TqmnlIlD9 f7TKFixYnBQnKanVwYoNEDED6E/wmh1T8gU/aK5No0YqUo21F3hA8I+VN/ipomHZ SPL0mI7n/8ms3tccdf/vzLgcb+vYjnrVRHc/qvwuHdTvGybvn036r8= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=SbF7NQ1Soid5Kk1dm6hBmzrogM4=; b=Q6Ga73j33KCjPdM2sBRX 9hXYZ4boKkA9ZbwGhH4BBm5v7nsYLvpqCWm04ABUBv+i1gz0hziVp4Satj1tcpZH 1Q95VGW+eCsYRaXUniyxIQR34VPCSJKRadmZ07B+AFTQ2C/4HXXxc15WMx/6I22j /EjRUOP+dcD+k330kPeTz5E= Received: (qmail 104585 invoked by alias); 12 Jan 2019 22:21:17 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 104499 invoked by uid 89); 12 Jan 2019 22:21:17 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy= X-HELO: mx1.suse.de Received: from mx2.suse.de (HELO mx1.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:14 +0000 Received: from relay1.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id A13BAAE60; Sat, 12 Jan 2019 22:21:10 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 7/9] [nvptx] Add vector_length 64 test-cases Date: Sat, 12 Jan 2019 23:21:29 +0100 Message-Id: <20190112222131.29519-8-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-IsSubscribed: yes Add some test-cases using vector_length 64. 2019-01-10 Tom de Vries * testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c: New test. --- .../libgomp.oacc-c-c++-common/vector-length-64-1.c | 17 +++++++++++++++++ .../libgomp.oacc-c-c++-common/vector-length-64-2.c | 21 +++++++++++++++++++++ .../libgomp.oacc-c-c++-common/vector-length-64-3.c | 17 +++++++++++++++++ 3 files changed, 55 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c new file mode 100644 index 00000000000..b6ee732f863 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c @@ -0,0 +1,17 @@ +#include +#include + +int +main (void) +{ +#pragma acc parallel vector_length (64) num_workers (16) /* { dg-warning "using num_workers \\(15\\), ignoring 16" "" { target openacc_nvidia_accel_configured } } */ + { +#pragma acc loop worker + for (unsigned int i = 0; i < 32; i++) +#pragma acc loop vector + for (unsigned int j = 0; j < 64; j++) + ; + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c new file mode 100644 index 00000000000..4dfbae8de91 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c @@ -0,0 +1,21 @@ +/* { dg-set-target-env-var "GOMP_OPENACC_DIM" ":16:" } */ +/* { dg-shouldfail "" { openacc_nvidia_accel_selected } } */ + +#include +#include + +int +main (void) +{ +#pragma acc parallel vector_length (64) + { +#pragma acc loop worker + for (unsigned int i = 0; i < 32; i++) +#pragma acc loop vector + for (unsigned int j = 0; j < 64; j++) + ; + } + + return 0; +} +/* { dg-output "The Nvidia accelerator has insufficient barrier resources" { target openacc_nvidia_accel_selected } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c new file mode 100644 index 00000000000..1acb40e8357 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c @@ -0,0 +1,17 @@ +#include +#include + +int +main (void) +{ +#pragma acc parallel vector_length (64) + { +#pragma acc loop worker + for (unsigned int i = 0; i < 32; i++) +#pragma acc loop vector + for (unsigned int j = 0; j < 64; j++) + ; + } + + return 0; +} From patchwork Sat Jan 12 22:21:30 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024010 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493951-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="UaIvCtrj"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ3t0w4Bz9sCh for ; Sun, 13 Jan 2019 09:22:33 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=oeq7DIKdzJEI1AYuRXrsTisitGkOK5oPNaoQjZo2b/eszhllP6Rm8 NhwHwlRHEUeanXP+wS237gyIBkqDivCPPsC7jKBlWdvNke/n9VoMe/2daox6029R TXP/fyrakph1FBJxzZQvCKtH1l6Atk2FxO/XlTAcwbI7nQOUDYSHco= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=qB/Bb4LimY+qB4KklegHMHWkU3U=; b=UaIvCtrjtZVyNOBwQRpk ZDc2K0v9bZWVzVpVUYuiuOBgJRlmbFmju6w2DwZIcVwr3xIJkpdd8gYLQUPO/8w+ aBuisA0TF9+sb7j+0k02DdIzC8rzqYDARrmtOT6n1c7/p99fUNqCHumAjxatBslj Luiu1Xo5CxUlCPGHm4tFHLU= Received: (qmail 104767 invoked by alias); 12 Jan 2019 22:21:18 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 104717 invoked by uid 89); 12 Jan 2019 22:21:18 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.2 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_SOFTFAIL autolearn=ham version=3.3.2 spammy=HX-detected-operating-system:timestamps X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (209.51.188.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:16 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1giReU-0007FC-6w for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:14 -0500 Received: from mx2.suse.de ([195.135.220.15]:52524 helo=mx1.suse.de) by eggs.gnu.org with esmtps (TLS1.0:DHE_RSA_AES_256_CBC_SHA1:32) (Exim 4.71) (envelope-from ) id 1giReU-0007EH-0o for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:14 -0500 Received: from relay2.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id D880DAE63; Sat, 12 Jan 2019 22:21:10 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 8/9] [nvptx] Enable setting vector length using -fopenacc-dim Date: Sat, 12 Jan 2019 23:21:30 +0100 Message-Id: <20190112222131.29519-9-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-detected-operating-system: by eggs.gnu.org: GNU/Linux 2.2.x-3.x (no timestamps) [generic] X-Received-From: 195.135.220.15 X-IsSubscribed: yes Enable setting vector length using -fopenacc-dim, f.i. -fopenacc-dim=::128. 2019-01-12 Tom de Vries * config/nvptx/nvptx.c (nvptx_goacc_validate_dims_1): Alow setting vector length using -fopenacc-dim. * plugin/plugin-nvptx.c (nvptx_exec): Update error message. --- gcc/config/nvptx/nvptx.c | 3 ++- libgomp/plugin/plugin-nvptx.c | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 8d2740cd50f..03c0f82f4a2 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -5705,7 +5705,8 @@ nvptx_goacc_validate_dims_1 (tree decl, int dims[], int fn_level, unsigned used) if (oacc_default_dims_p) { - dims[GOMP_DIM_VECTOR] = default_vector_length; + if (dims[GOMP_DIM_VECTOR] < 0) + dims[GOMP_DIM_VECTOR] = default_vector_length; if (dims[GOMP_DIM_WORKER] < 0) dims[GOMP_DIM_WORKER] = PTX_DEFAULT_RUNTIME_DIM; if (dims[GOMP_DIM_GANG] < 0) diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index 8912660966a..dd2bcf3083f 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -1321,7 +1321,7 @@ nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs, " region or '-fopenacc-dim=:x:' where x <= 15" "; " "or, recompile the program with 'vector_length = 32' on that" - " offloaded region" + " offloaded region or '-fopenacc-dim=::32'" ".\n"); GOMP_PLUGIN_fatal (msg, targ_fn->launch->fn, dims[GOMP_DIM_WORKER], dims[GOMP_DIM_VECTOR]); From patchwork Sat Jan 12 22:21:31 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1024013 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493954-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="hGlI8SpX"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 43cZ4f74P6z9s7T for ; Sun, 13 Jan 2019 09:23:14 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=Kt/GwtcU6GgW9mVrzADFxOufqQ+mtL8ksayMylGtd0fHwn1uRK+t0 bpWtXJXsjqlpf8NjQ5X6wqRJagot+1yCO4fBtUjeUwjQdNcF0A1+wLTmDxM+Wdqz 2l9DiJn3qCiLbEY90WwYIguTGdva7UuTVoUI5lNsZ09Z7qKjIUJ19s= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=AwPar2rOh3EGzZFzgbnCLYGf73c=; b=hGlI8SpX095dlB11xy1W rJ141Bj7ToFHuH40El0stfvcKb/i9McODh+Y7KlNN/w0uVgCoNe/bGjPku07kT7a cB0Fee91KpS2/TRONcke67PcihcRTv8TQO2xRztanhCVZpa4c70S+oaojYW/5xTD g5rNBuHOt0b9TwUNNW0ddHg= Received: (qmail 105334 invoked by alias); 12 Jan 2019 22:21:22 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 105242 invoked by uid 89); 12 Jan 2019 22:21:21 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.2 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_SOFTFAIL autolearn=ham version=3.3.2 spammy=HX-detected-operating-system:timestamps, m*, exercise, ij X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (209.51.188.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Sat, 12 Jan 2019 22:21:17 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1giReU-0007Fb-L4 for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:15 -0500 Received: from mx2.suse.de ([195.135.220.15]:52526 helo=mx1.suse.de) by eggs.gnu.org with esmtps (TLS1.0:DHE_RSA_AES_256_CBC_SHA1:32) (Exim 4.71) (envelope-from ) id 1giReU-0007EZ-9q for gcc-patches@gcc.gnu.org; Sat, 12 Jan 2019 17:21:14 -0500 Received: from relay1.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id 33406AE7A; Sat, 12 Jan 2019 22:21:11 +0000 (UTC) From: Tom de Vries To: gcc-patches@gcc.gnu.org Cc: Thomas Schwinge Subject: [PATCH 9/9] [nvptx] Enable setting vector length using -fopenacc-dim -- testcases Date: Sat, 12 Jan 2019 23:21:31 +0100 Message-Id: <20190112222131.29519-10-tdevries@suse.de> In-Reply-To: <20190112222131.29519-1-tdevries@suse.de> References: <20190112222131.29519-1-tdevries@suse.de> X-detected-operating-system: by eggs.gnu.org: GNU/Linux 2.2.x-3.x (no timestamps) [generic] X-Received-From: 195.135.220.15 X-IsSubscribed: yes Add some test-cases that set vector length using -fopenacc-dim. 2019-01-12 Tom de Vries * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: New test. * testsuite/libgomp.oacc-fortran/gemm-2.f90: New test. --- .../libgomp.oacc-c-c++-common/pr85486-2.c | 52 ++++++++++++++ .../vector-length-128-2.c | 39 +++++++++++ .../vector-length-128-5.c | 41 +++++++++++ libgomp/testsuite/libgomp.oacc-fortran/gemm-2.f90 | 80 ++++++++++++++++++++++ 4 files changed, 212 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/gemm-2.f90 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c new file mode 100644 index 00000000000..f6ca263166d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c @@ -0,0 +1,52 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-fopenacc-dim=::128" } */ + +/* Minimized from ref-1.C. */ + +#include + +#pragma acc routine vector +void __attribute__((noinline, noclone)) +Vector (int *ptr, int n, const int inc) +{ + #pragma acc loop vector + for (unsigned ix = 0; ix < n; ix++) + ptr[ix] += inc; +} + +int +main (void) +{ + const int n = 32, m=32; + + int ary[m][n]; + unsigned ix, iy; + + for (ix = m; ix--;) + for (iy = n; iy--;) + ary[ix][iy] = (1 << 16) + (ix << 8) + iy; + + int err = 0; + +#pragma acc parallel copy (ary) + { + Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16)); + } + + for (ix = m; ix--;) + for (iy = n; iy--;) + if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy)) + { + printf ("ary[%u][%u] = %x expected %x\n", + ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy)); + err++; + } + + if (err) + { + printf ("%d failed\n", err); + return 1; + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c new file mode 100644 index 00000000000..8b5b2a4a92d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c @@ -0,0 +1,39 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-fopenacc-dim=::128" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ + +#include + +#define N 1024 + +unsigned int a[N]; +unsigned int b[N]; +unsigned int c[N]; +unsigned int n = N; + +int +main (void) +{ + for (unsigned int i = 0; i < n; ++i) + { + a[i] = i % 3; + b[i] = i % 5; + } + +#pragma acc parallel copyin (a,b) copyout (c) + { +#pragma acc loop vector + for (unsigned int i = 0; i < n; i++) + c[i] = a[i] + b[i]; + } + + for (unsigned int i = 0; i < n; ++i) + if (c[i] != (i % 3) + (i % 5)) + abort (); + + return 0; +} + +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c new file mode 100644 index 00000000000..e60f1c28db4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c @@ -0,0 +1,41 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-fopenacc-dim=:2:128" } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */ + +#include + +#define N 1024 + +unsigned int a[N]; +unsigned int b[N]; +unsigned int c[N]; +unsigned int n = N; + +int +main (void) +{ + for (unsigned int i = 0; i < n; ++i) + { + a[i] = i % 3; + b[i] = i % 5; + } + +#pragma acc parallel copyin (a,b) copyout (c) + { +#pragma acc loop worker + for (unsigned int i = 0; i < 4; i++) +#pragma acc loop vector + for (unsigned int j = 0; j < n / 4; j++) + c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j]; + } + + for (unsigned int i = 0; i < n; ++i) + if (c[i] != (i % 3) + (i % 5)) + abort (); + + return 0; +} + +/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/gemm-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/gemm-2.f90 new file mode 100644 index 00000000000..fe108732a5f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/gemm-2.f90 @@ -0,0 +1,80 @@ +! Exercise three levels of parallelism using SGEMM from BLAS. + +! { dg-do run } +! { dg-additional-options "-fopenacc-dim=::128" } + +! Implicitly set vector_length to 128 using -fopenacc-dim. +subroutine openacc_sgemm (m, n, k, alpha, a, b, beta, c) + integer :: m, n, k + real :: alpha, beta + real :: a(k,*), b(k,*), c(m,*) + + integer :: i, j, l + real :: temp + + !$acc parallel loop copy(c(1:m,1:n)) copyin(a(1:k,1:m),b(1:k,1:n)) firstprivate (temp) + do j = 1, n + !$acc loop + do i = 1, m + temp = 0.0 + !$acc loop reduction(+:temp) + do l = 1, k + temp = temp + a(l,i)*b(l,j) + end do + if(beta == 0.0) then + c(i,j) = alpha*temp + else + c(i,j) = alpha*temp + beta*c(i,j) + end if + end do + end do +end subroutine openacc_sgemm + +subroutine host_sgemm (m, n, k, alpha, a, b, beta, c) + integer :: m, n, k + real :: alpha, beta + real :: a(k,*), b(k,*), c(m,*) + + integer :: i, j, l + real :: temp + + do j = 1, n + do i = 1, m + temp = 0.0 + do l = 1, k + temp = temp + a(l,i)*b(l,j) + end do + if(beta == 0.0) then + c(i,j) = alpha*temp + else + c(i,j) = alpha*temp + beta*c(i,j) + end if + end do + end do +end subroutine host_sgemm + +program main + integer, parameter :: M = 100, N = 50, K = 2000 + real :: a(K, M), b(K, N), c(M, N), d (M, N), e (M, N) + real alpha, beta + integer i, j + + a(:,:) = 1.0 + b(:,:) = 0.25 + + c(:,:) = 0.0 + d(:,:) = 0.0 + e(:,:) = 0.0 + + alpha = 1.05 + beta = 1.25 + + call openacc_sgemm (M, N, K, alpha, a, b, beta, c) + call host_sgemm (M, N, K, alpha, a, b, beta, e) + + do i = 1, m + do j = 1, n + if (c(i,j) /= e(i,j)) call abort + end do + end do +end program main