From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 127013 invoked by alias); 7 Jul 2015 14:13:03 -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 127003 invoked by uid 89); 7 Jul 2015 14:13:03 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.6 required=5.0 tests=BAYES_00,FREEMAIL_FROM,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=ham version=3.3.2 X-HELO: mail-ie0-f177.google.com Received: from mail-ie0-f177.google.com (HELO mail-ie0-f177.google.com) (209.85.223.177) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Tue, 07 Jul 2015 14:13:01 +0000 Received: by iecuq6 with SMTP id uq6so135036617iec.2 for ; Tue, 07 Jul 2015 07:12:59 -0700 (PDT) X-Received: by 10.107.5.1 with SMTP id 1mr6262420iof.88.1436278379362; Tue, 07 Jul 2015 07:12:59 -0700 (PDT) Received: from ?IPv6:2600:1000:b012:8aaa:a2a8:cdff:fe3e:b48? ([2600:1000:b012:8aaa:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id ot6sm11924803igb.11.2015.07.07.07.12.57 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Tue, 07 Jul 2015 07:12:58 -0700 (PDT) Message-ID: <559BDE68.9010302@acm.org> Date: Tue, 07 Jul 2015 14:13:00 -0000 From: Nathan Sidwell User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.7.0 MIME-Version: 1.0 To: Jakub Jelinek CC: GCC Patches Subject: Re: [gomp] Move openacc vector& worker single handling to RTL References: <5597120D.2080308@acm.org> <20150703231159.GP10247@tucnak.redhat.com> <559844EF.6010208@acm.org> <559AD85B.2050102@acm.org> <20150707095408.GD10247@tucnak.redhat.com> In-Reply-To: <20150707095408.GD10247@tucnak.redhat.com> Content-Type: text/plain; charset=windows-1252; format=flowed Content-Transfer-Encoding: 7bit X-SW-Source: 2015-07/txt/msg00518.txt.bz2 On 07/07/15 05:54, Jakub Jelinek wrote: > On Mon, Jul 06, 2015 at 03:34:51PM -0400, Nathan Sidwell wrote: > How does this interact with > #pragma acc routine {gang,worker,vector,seq} ? > Or is that something to be added later on? That is to be added later on. I suspect such routines will trivially work, as they'll be marked up with the loop head/tail functions and levels builtin (the latter might need a bit of reworking). What will need additional work at that point is the callers of routines -- they're typically called from a foo-single mode, but need to get all threads into the called function. I'm thinking each call site will look like a mini-loop[*] surrounded by a hesd/tail marker. (all that can be done in the device-side compiler once real call sites are known.) nathan [*] of course it won't be a loop. Perhaps fork/join are less confusing names after all. WDYT?