From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 70766 invoked by alias); 21 Oct 2015 10:03:45 -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 70706 invoked by uid 89); 21 Oct 2015 10:03:44 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=AWL,BAYES_00,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Wed, 21 Oct 2015 10:03:42 +0000 Received: from int-mx14.intmail.prod.int.phx2.redhat.com (int-mx14.intmail.prod.int.phx2.redhat.com [10.5.11.27]) by mx1.redhat.com (Postfix) with ESMTPS id 758C6354173; Wed, 21 Oct 2015 10:03:41 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-53.ams2.redhat.com [10.36.116.53]) by int-mx14.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t9LA3d74012087 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Wed, 21 Oct 2015 06:03:41 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id t9LA3cj9019223; Wed, 21 Oct 2015 12:03:39 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id t9LA3buV019222; Wed, 21 Oct 2015 12:03:37 +0200 Date: Wed, 21 Oct 2015 10:12:00 -0000 From: Jakub Jelinek To: Alexander Monakov Cc: gcc-patches@gcc.gnu.org, Dmitry Melnik Subject: Re: [gomp4 13/14] libgomp: provide minimal GOMP_teams Message-ID: <20151021100337.GT478@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <1445366076-16082-1-git-send-email-amonakov@ispras.ru> <1445366076-16082-14-git-send-email-amonakov@ispras.ru> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <1445366076-16082-14-git-send-email-amonakov@ispras.ru> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-10/txt/msg02054.txt.bz2 On Tue, Oct 20, 2015 at 09:34:35PM +0300, Alexander Monakov wrote: > On NVPTX, we don't need most of target.c functionality, except for GOMP_teams. > Provide it as a copy of the generic implementation for now (it most likely > will need to change down the line: on NVPTX we do need to spawn several > thread blocks for #pragma omp teams). > > Alternatively, it might make sense to split GOMP_teams out of target.c into > its own file (teams.c?), leaving target.c a 0-size stub in config/nvptx. > > * config/nvptx/target.c: New. This is fine. I bet it will need changes on the compiler side too, GOMP_teams has been written with rough idea of the dynamic parallelism, we'll need to do something with private (per-CTA) vars, shared (global) vars etc., dunno if with tweaks just in the late omp lowering pass after it is lowered/expanded as OpenMP already, or if we'll need Martin's HSA approach. Jakub