From patchwork Fri Mar 30 13:31:40 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 893335 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-475644-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="PDAYY8Q2"; 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 40CMwZ2hw6z9s1r for ; Sat, 31 Mar 2018 00:31:56 +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:to:cc :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=nA8QHrP17ArR01z/3iGWwmOFpvrCKRD3HDV41uS+bjZ9eZUHYM SLHwfCwp9GFY0uv6t2g62GqFCFpK9A9QGbJAplwfDCZ/TpKZA2PpkGuc1ctkozlj kTBB+2zsLyjgYZf2x6mJE6W/RTUzI62NO2sgKPCK30j9+/ebbbIphWzdI= 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:to:cc :from:subject:message-id:date:mime-version:content-type; s= default; bh=XEF667a0WmeqxpPLsG/mzIMPZZ0=; b=PDAYY8Q2hUd/1L/qyDMj Lurcvvr2HJPowB5bSnRRJJJhsnmFNbshA9C18uWrJn2fIHp0UUL/ixfVc9K2kiny jNWKJoTZ3ai3Pd7a7HSnIRvp0IwnU7L4GQro77ZodE4GM9MgiGOa5teyN9D0+fcp IaQqDPnNMGh3eAh4zVlcg64= Received: (qmail 67222 invoked by alias); 30 Mar 2018 13:31:49 -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 67201 invoked by uid 89); 30 Mar 2018 13:31:48 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.8 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=sk:fdumpt, sk:fdump-t 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; Fri, 30 Mar 2018 13:31:47 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1f1u89-0000q1-Bp from Tom_deVries@mentor.com for gcc-patches@gcc.gnu.org; Fri, 30 Mar 2018 06:31:45 -0700 Received: from [172.30.72.232] (137.202.0.87) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Fri, 30 Mar 2018 14:31:41 +0100 To: GCC Patches , Thomas Schwinge CC: Cesar Philippidis From: Tom de Vries Subject: [og7, openacc, committed] Add vector-length-128-{1, 2, 3}.c test-cases Message-ID: <6e629c43-0e76-908b-0552-8f7651eb9ba0@mentor.com> Date: Fri, 30 Mar 2018 15:31:40 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.6.0 MIME-Version: 1.0 X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) Hi, this patch adds three testcases, setting vector length to 128 in three different ways: 1. the vector_length clause 2. the -fopenacc-dim option 3. the GOMP_OPENACC_DIM variable The tests contains: - a check of the dimensions that the compiler decides upon - a check of the dimensions used at runtime by libgomp [ The first check is made possible by the "Add scan-offload-tree-dump" patch I've just committed. The second used setting environment variable GOMP_DEBUG to 1. We cannot do this in main using setenv, as we can for GOMP_OPENACC_DIM, so we have to use dg-set-target-env-var. This means the dg-output scans will fail in a remote test setup where dg-set-target-env-var is broken. This is annoying, but not as annoying as not having tests that DTRT in a local test setup. Also annoying is the fact that GOMP_DEBUG=1 prints a lot of lines into libgomp.log. We could improve on that by not emitting the PTX code for GOMP_DEBUG=1, and maybe emit that depending on the value of another variable, say GOMP_OPENACC_DEBUG=n. ] Currently, in all three test cases we don't use vector length 128, but fall back to warp_size, so checks test for 32, not 128 vector_length. Tested libgomp on x86_64 build with nvptx accelerator. Committed. Thanks, - Tom [openacc] Add vector-length-128-{1,2,3}.c test-cases 2018-03-30 Tom de Vries * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.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-3.c: New test. --- .../vector-length-128-1.c | 39 ++++++++++++++++++++ .../vector-length-128-2.c | 40 +++++++++++++++++++++ .../vector-length-128-3.c | 42 ++++++++++++++++++++++ 3 files changed, 121 insertions(+) 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 new file mode 100644 index 0000000..fab5b0d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c @@ -0,0 +1,39 @@ +/* { 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 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-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" } */ 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 0000000..cc6fd55 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c @@ -0,0 +1,40 @@ +/* { 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-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" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c new file mode 100644 index 0000000..c403e74 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c @@ -0,0 +1,42 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */ +/* We default to warp size 32 for the vector length, so the GOMP_OPENACC_DIM has + no effect. */ +/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "-:-:128" } */ +/* { 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, 32\\)" "oaccdevlow" } } */ +/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */