From patchwork Thu Nov 10 10:47:07 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 81627 Delivered-To: patch@linaro.org Received: by 10.140.97.165 with SMTP id m34csp655580qge; Thu, 10 Nov 2016 02:48:00 -0800 (PST) X-Received: by 10.98.137.21 with SMTP id v21mr9072159pfd.48.1478774880556; Thu, 10 Nov 2016 02:48:00 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id yg8si3500932pab.187.2016.11.10.02.48.00 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Thu, 10 Nov 2016 02:48:00 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-440938-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) client-ip=209.132.180.131; Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org; spf=pass (google.com: domain of gcc-patches-return-440938-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-440938-patch=linaro.org@gcc.gnu.org DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :subject:to:cc:message-id:date:mime-version:content-type; q=dns; s=default; b=Q9llBCgAhPim8d5uroD0XwveImUShareq6hWi7Ctft76E7ytWv kxTv1aQm0pE4mi7V43nQEU70LrXgJ+qSQKZ4WLIn30E7PF2f2rt1WD09sqQnJrjA ONSA6WdfGLssDc1zN0xOfSe9gPu+hqTRr5YMh5L3dvmCvh+ff6HCCwmMA= 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 :subject:to:cc:message-id:date:mime-version:content-type; s= default; bh=SHaE7iMx4OtSZjVrhaJeQaUaeWc=; b=AXU0Vs/zxCIsG15r8p0l ruBcgfBs+4Fj32kKRm+3wIhJNxqE9bDeiKDLZPjonqHLQ4/nKKM/lQUiPEPQvyyY GNtV/gxxoYh102tOXh+oe18V+1cA6WSpaPCDT+6wQhstmyrG+A98R4fCgtkHsH65 mQtvTIMyTMUQcxmCoMonyYg= Received: (qmail 47517 invoked by alias); 10 Nov 2016 10:47:21 -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 47420 invoked by uid 89); 10 Nov 2016 10:47:21 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.1 required=5.0 tests=AWL, BAYES_00, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=no version=3.3.2 spammy=gang, nathancodesourcerycom, 2016XXXX, 2016-XX-XX 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, 10 Nov 2016 10:47:18 +0000 Received: from svr-orw-mbx-03.mgc.mentorg.com ([147.34.90.203]) by relay1.mentorg.com with esmtp id 1c4mt2-0002vZ-PV from ChungLin_Tang@mentor.com ; Thu, 10 Nov 2016 02:47:16 -0800 Received: from svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) by svr-orw-mbx-03.mgc.mentorg.com (147.34.90.203) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Thu, 10 Nov 2016 02:47:14 -0800 Received: from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) with Microsoft SMTP Server (TLS) id 15.0.1210.3 via Frontend Transport; Thu, 10 Nov 2016 02:47:10 -0800 From: Chung-Lin Tang Subject: [Patch 5/5] OpenACC tile clause support, libgomp testsuite patches To: gcc-patches , Jakub Jelinek CC: Nathan Sidwell , Cesar Philippidis , Chung-Lin Tang Message-ID: Date: Thu, 10 Nov 2016 18:47:07 +0800 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.11; rv:45.0) Gecko/20100101 Thunderbird/45.4.0 MIME-Version: 1.0 Some additional tests and adjustments to existing ones were made. 2016-XX-XX Nathan Sidwell Chung-Lin Tang libgomp/ * testsuite/libgomp.oacc-c-c++-common/tile-1.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust and add additional case. * testsuite/libgomp.oacc-c-c++-common/vprop.c: XFAIL under "openacc_nvidia_accel_selected". * libgomp.oacc-fortran/nested-function-1.f90 (test2): Add num_workers(8) clause. Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c (revision 241809) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c (working copy) @@ -112,7 +112,7 @@ int vector_1 (int *ary, int size) ary[ix] = place (); } - return check (ary, size, 0, 0, 1); + return check (ary, size, 0, 1, 1); } int vector_2 (int *ary, int size) @@ -196,10 +196,24 @@ int gang_3 (int *ary, int size) ary[ix + jx * 64] = place (); } + return check (ary, size, 1, 1, 1); +} + +int gang_4 (int *ary, int size) +{ + clear (ary, size); + +#pragma acc parallel vector_length(32) copy(ary[0:size]) firstprivate (size) + { +#pragma acc loop auto + for (int jx = 0; jx < size; jx++) + ary[jx] = place (); + } + return check (ary, size, 1, 0, 1); } -#define N (32*32*32) +#define N (32*32*32*2) int main () { int ondev = 0; @@ -227,6 +241,8 @@ int main () return 1; if (gang_3 (ary, N)) return 1; + if (gang_4 (ary, N)) + return 1; return 0; } Index: libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/tile-1.c (revision 0) @@ -0,0 +1,281 @@ +/* This code uses nvptx inline assembly guarded with acc_on_device, which is + not optimized away at -O0, and then confuses the target assembler. + { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */ + +/* { dg-additional-options "-fopenacc-dim=32" } */ + +#include +#include + +static int check (const int *ary, int size, int gp, int wp, int vp) +{ + int exit = 0; + int ix; + int gangs[32], workers[32], vectors[32]; + + for (ix = 0; ix < 32; ix++) + gangs[ix] = workers[ix] = vectors[ix] = 0; + + for (ix = 0; ix < size; ix++) + { + vectors[ary[ix] & 0xff]++; + workers[(ary[ix] >> 8) & 0xff]++; + gangs[(ary[ix] >> 16) & 0xff]++; + } + + for (ix = 0; ix < 32; ix++) + { + if (gp) + { + int expect = gangs[0]; + if (gangs[ix] != expect) + { + exit = 1; + printf ("gang %d not used %d times\n", ix, expect); + } + } + else if (ix && gangs[ix]) + { + exit = 1; + printf ("gang %d unexpectedly used\n", ix); + } + + if (wp) + { + int expect = workers[0]; + if (workers[ix] != expect) + { + exit = 1; + printf ("worker %d not used %d times\n", ix, expect); + } + } + else if (ix && workers[ix]) + { + exit = 1; + printf ("worker %d unexpectedly used\n", ix); + } + + if (vp) + { + int expect = vectors[0]; + if (vectors[ix] != expect) + { + exit = 1; + printf ("vector %d not used %d times\n", ix, expect); + } + } + else if (ix && vectors[ix]) + { + exit = 1; + printf ("vector %d unexpectedly used\n", ix); + } + + } + return exit; +} + +#pragma acc routine seq +static int __attribute__((noinline)) place () +{ + int r = 0; + + if (acc_on_device (acc_device_nvidia)) + { + int g = 0, w = 0, v = 0; + + __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g)); + __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w)); + __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v)); + r = (g << 16) | (w << 8) | v; + } + return r; +} + +static void clear (int *ary, int size) +{ + int ix; + + for (ix = 0; ix < size; ix++) + ary[ix] = -1; +} + +int gang_vector_1 (int *ary, int size) +{ + clear (ary, size); +#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(128) gang vector + for (int jx = 0; jx < size; jx++) + ary[jx] = place (); + } + + return check (ary, size, 1, 0, 1); +} + +int gang_vector_2a (int *ary, int size) +{ + if (size % 256) + return 1; + + clear (ary, size); +#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(64, 64) gang vector + for (int jx = 0; jx < size / 256; jx++) + for (int ix = 0; ix < 256; ix++) + ary[jx * 256 + ix] = place (); + } + + return check (ary, size, 1, 0, 1); +} + +int gang_vector_2b (int *ary, int size) +{ + if (size % 256) + return 1; + + clear (ary, size); +#pragma acc parallel vector_length(32) num_gangs (32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(64, 64) gang vector + for (int jx = 0; jx < size; jx += 256) + for (int ix = 0; ix < 256; ix++) + ary[jx + ix] = place (); + } + + return check (ary, size, 1, 0, 1); +} + +int worker_vector_2a (int *ary, int size) +{ + if (size % 256) + return 1; + + clear (ary, size); +#pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(64, 64) worker vector + for (int jx = 0; jx < size / 256; jx++) + for (int ix = 0; ix < 256; ix++) + ary[jx * 256 + ix] = place (); + } + + return check (ary, size, 0, 1, 1); +} + +int worker_vector_2b (int *ary, int size) +{ + if (size % 256) + return 1; + + clear (ary, size); +#pragma acc parallel vector_length(32) num_workers (32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(64, 64) worker vector + for (int jx = 0; jx < size; jx += 256) + for (int ix = 0; ix < 256; ix++) + ary[jx + ix] = place (); + } + + return check (ary, size, 0, 1, 1); +} + +int gang_worker_vector_2a (int *ary, int size) +{ + if (size % 256) + return 1; + clear (ary, size); +#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(32, 32) + for (int jx = 0; jx < size / 256; jx++) + for (int ix = 0; ix < 256; ix++) + ary[jx * 256 + ix] = place (); + } + + return check (ary, size, 1, 1, 1); +} + +int gang_worker_vector_2b (int *ary, int size) +{ + if (size % 256) + return 1; + clear (ary, size); +#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(32, 32) + for (int jx = 0; jx < size; jx += 256) + for (int ix = 0; ix < 256; ix++) + ary[jx + ix] = place (); + } + + return check (ary, size, 1, 1, 1); +} + +int gang_worker_vector_star_2a (int *ary, int size) +{ + if (size % 256) + return 1; + + clear (ary, size); +#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(*, *) + for (int jx = 0; jx < size / 256; jx++) + for (int ix = 0; ix < 256; ix++) + ary[jx * 256 + ix] = place (); + } + + return check (ary, size, 1, 1, 1); +} + +int gang_worker_vector_star_2b (int *ary, int size) +{ + if (size % 256) + return 1; + + clear (ary, size); +#pragma acc parallel vector_length(32) num_workers (32) num_gangs(32) copy (ary[0:size]) firstprivate (size) + { +#pragma acc loop tile(*, *) + for (int jx = 0; jx < size; jx +=256) + for (int ix = 0; ix < 256; ix++) + ary[jx + ix] = place (); + } + + return check (ary, size, 1, 1, 1); +} + +#define N (32*32*32*8) +int main () +{ + int ondev = 0; + +#pragma acc parallel copy(ondev) + { + ondev = acc_on_device (acc_device_not_host); + } + if (!ondev) + return 0; + + int ary[N]; + if (gang_vector_1 (ary, N)) + return 1; + if (gang_vector_2a (ary, N)) + return 1; + if (worker_vector_2a (ary, N)) + return 1; + if (gang_worker_vector_2a (ary, N)) + return 1; + if (gang_worker_vector_star_2a (ary, N)) + return 1; + if (gang_vector_2b (ary, N)) + return 1; + if (worker_vector_2b (ary, N)) + return 1; + if (gang_worker_vector_2b (ary, N)) + return 1; + if (gang_worker_vector_star_2b (ary, N)) + return 1; + return 0; +} Index: libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c (revision 241809) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c (working copy) @@ -1,3 +1,6 @@ +/* { dg-do run } */ +/* { dg-xfail-run-if "PR78266" { openacc_nvidia_accel_selected } { "*" } { "" } } */ + #include #define test(type) \ Index: libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90 =================================================================== --- libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90 (revision 241809) +++ libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90 (working copy) @@ -33,7 +33,7 @@ firstdo: do i = 1, 3 subroutine test2 integer :: a(3,3,3), k, kk, kkk, l, ll, lll a = 0 - !$acc parallel + !$acc parallel num_workers(8) ! Use "gang(static:1)" here and below to effectively turn gang-redundant ! execution mode into something like gang-single. !$acc loop gang(static:1) collapse(1)