From patchwork Fri Aug 21 20:56:27 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 509644 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 0AB651402A6 for ; Sat, 22 Aug 2015 06:56:42 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=bYuhVwk+; dkim-atps=neutral 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=Ggl+N4+Ymdv+L0Bb+6nd+8nGYS+sJa8H8Ahr3pB03K4Ut/+2C4 Gn4eAV2Af20VgiKGf4LrjKbQgOKp1qgOuJNFOUe84VtkitGcaGUyIOLDaYhxZGjf qEAlQTmZCfmlUpFEal1lFJ2KhuQmNZ+Ct6EGTZYGJZ8F6M7jv4N9Q/HrM= 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=bH+jyZ9qB+G3dFWt9HAwmrCuM+0=; b=bYuhVwk+prIjhsv2XrS3 OCyOAFvsRUsUFZb6riAMfVgi5n1aYlWQV/WFNxqM2GBH7WK7UgyV4kFIAd2H5l4P sHzeajO66kkrQ5CwPEwbGzQ05ItRgDaHlKMybzb+ojvCyA6Q6Zhc956msi4+Zzub J6qEwISBcjSVh83S1rq9ldE= Received: (qmail 62799 invoked by alias); 21 Aug 2015 20:56:32 -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 62786 invoked by uid 89); 21 Aug 2015 20:56:32 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=BAYES_00, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-qk0-f171.google.com Received: from mail-qk0-f171.google.com (HELO mail-qk0-f171.google.com) (209.85.220.171) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Fri, 21 Aug 2015 20:56:31 +0000 Received: by qkbm65 with SMTP id m65so38016964qkb.2 for ; Fri, 21 Aug 2015 13:56:29 -0700 (PDT) X-Received: by 10.55.209.157 with SMTP id o29mr21757434qkl.34.1440190589072; Fri, 21 Aug 2015 13:56:29 -0700 (PDT) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id z202sm5143093qhd.15.2015.08.21.13.56.28 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Fri, 21 Aug 2015 13:56:28 -0700 (PDT) To: GCC Patches Cc: Cesar Philippidis From: Nathan Sidwell Subject: [gomp4] force vector size to 32 Message-ID: <55D7907B.8040002@acm.org> Date: Fri, 21 Aug 2015 16:56:27 -0400 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.1.0 MIME-Version: 1.0 I've committed this patch to force the vector_length to constant 32. This helps in 2 places 1) Cesar's reduction patch set made this assumption, as we can't emit a loop for vector reductions, as he explained. 2) Routines can now optimize on this axis knowing it is 32. The only case lost is where one selects a single element from each gang, which is a bizarre way of operating the GPGPU, so no real loss. Also fixed a testcase that had got a duplicated body. nathan 2015-08-21 Nathan Sidwell * config/nvptx/nvptx.c (nvptx_validate_dims): Set vector length to 32 always. testsuite/ * c-c++-common/goacc/routine-5.c: Remove duplication. Index: gcc/config/nvptx/nvptx.c =================================================================== --- gcc/config/nvptx/nvptx.c (revision 227082) +++ gcc/config/nvptx/nvptx.c (working copy) @@ -3570,29 +3570,18 @@ nvptx_validate_dims (tree decl, int dims { bool changed = false; - if (fn_level >= 0) - /* This is a routine. All dimensions are dynamic and controlled - by the calling function. Because we permit a 1vx1wxNg - geometry, we can't take the opportunity to fix the vector - dimension inside a routine. Perhaps we should? */ - return false; - - /* If the worker size is not 1, the vector size must be 32. If - the vector size is not 1, it must be 32. */ - if ((dims[GOMP_DIM_WORKER] > 1 || dims[GOMP_DIM_WORKER] == 0) - || (dims[GOMP_DIM_VECTOR] > 1 || dims[GOMP_DIM_VECTOR] == 0)) + /* The vector size must be 32, unless this is a SEQ routine. */ + if (fn_level <= GOMP_DIM_VECTOR + && dims[GOMP_DIM_VECTOR] != PTX_VECTOR_LENGTH) { - if (dims[GOMP_DIM_VECTOR] != PTX_VECTOR_LENGTH) - { - if (dims[GOMP_DIM_VECTOR] >= 0) - warning_at (DECL_SOURCE_LOCATION (decl), 0, - dims[GOMP_DIM_VECTOR] - ? "using vector_length (%d), ignoring %d" - : "using vector_length (%d), ignoring runtime setting", - PTX_VECTOR_LENGTH, dims[GOMP_DIM_VECTOR]); - dims[GOMP_DIM_VECTOR] = PTX_VECTOR_LENGTH; - changed = true; - } + if (dims[GOMP_DIM_VECTOR] >= 0 && fn_level < 0) + warning_at (DECL_SOURCE_LOCATION (decl), 0, + dims[GOMP_DIM_VECTOR] + ? "using vector_length (%d), ignoring %d" + : "using vector_length (%d), ignoring runtime setting", + PTX_VECTOR_LENGTH, dims[GOMP_DIM_VECTOR]); + dims[GOMP_DIM_VECTOR] = PTX_VECTOR_LENGTH; + changed = true; } /* Check the num workers is not too large. */ Index: gcc/testsuite/c-c++-common/goacc/routine-5.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/routine-5.c (revision 227082) +++ gcc/testsuite/c-c++-common/goacc/routine-5.c (working copy) @@ -17,24 +17,3 @@ void Foo () #pragma acc routine // { dg-error "not followed by function" } int i; - -/* Test invalid use of clauses with routine. */ -/* { dg-do compile } */ -/* { dg-additional-options "-W */ - -void Bar (); - -void Foo () -{ - Bar (); -} - -#pragma acc routine (Bar) // { dg-error "must be applied before use" } - -#pragma acc routine (Foo) gang // { dg-error "must be applied before definition" } - -#pragma acc routine (Baz) // { dg-error "not been declared" } - -#pragma acc routine // { dg-error "not followed by function" } -int i; -