From patchwork Tue Feb 17 18:06:35 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 440677 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 1713A1401DD for ; Wed, 18 Feb 2015 05:07:01 +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:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type; q=dns; s=default; b=BHsisYQmRMBG7g2UJKajRUcaVtzYL oibr1cU5s3di/tJGHCKo/dKI1hOomu+yWeAHIy6OWmumV6rrk8CHGYG6lPYoGN83 6wGn54UlprzMwHAkx2/596Wv5W+bNsjbG2Cv195Uu6QUb1QhO3rJmyr3STwU3cls 9RzQGyMu5vUl+k= 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 :to:subject:in-reply-to:references:date:message-id:mime-version :content-type; s=default; bh=JDobxjrWw1Wyg8XO6rPiUJdtdKw=; b=teV TlKBOPDR1lYeiTadsVhp/vGoRX6NjNn5+d2mgrBoAG/USfMpkV9ES4TXYEyaJuqI 1RJrT+3OlbP03godEGUkqExzH8j4DZI8AePHHXqwRDUKHdhZV2L37PYenO0GOxu1 3Kz8xWpBUMHzDBGgcFktgufhykDAMevyoS3B0waM= Received: (qmail 8573 invoked by alias); 17 Feb 2015 18:06:53 -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 8561 invoked by uid 89); 17 Feb 2015 18:06:53 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.3 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.2 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; Tue, 17 Feb 2015 18:06:51 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1YNmXn-0001sY-Mw from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Tue, 17 Feb 2015 10:06:48 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.224.2; Tue, 17 Feb 2015 18:06:46 +0000 From: Thomas Schwinge To: Subject: Re: Merge current set of OpenACC changes from gomp-4_0-branch In-Reply-To: <87egqvu77s.fsf@schwinge.name> References: <87egqvu77s.fsf@schwinge.name> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (i586-pc-linux-gnu) Date: Tue, 17 Feb 2015 19:06:35 +0100 Message-ID: <878ufwh0pg.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Thu, 15 Jan 2015 21:20:07 +0100, I wrote: > In r219682, I have committed to trunk our current set of OpenACC changes, > which we had prepared on gomp-4_0-branch. This whole file is scheduled to go away: the routines are to be replaced by builtins which are expanded in the nvptx backend, but until we're there, here's a patch to make it at least work; committed to trunk in r220768: commit f816d7a63c8bc11c81080a0b34bf587d46b6f4c6 Author: tschwinge Date: Tue Feb 17 18:05:24 2015 +0000 libgomp: Make nvptx helper routines self-contained. libgomp/ * oacc-ptx.h (GOACC_INTERNAL_PTX): Add GOACC_tid, GOACC_ntid, GOACC_ctaid, and GOACC_nctaid routines. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@220768 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 6 ++ libgomp/oacc-ptx.h | 224 ++++++++++++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 230 insertions(+) Grüße, Thomas diff --git libgomp/ChangeLog libgomp/ChangeLog index 6c24531..2c32d9e 100644 --- libgomp/ChangeLog +++ libgomp/ChangeLog @@ -1,3 +1,9 @@ +2015-02-17 Thomas Schwinge + Cesar Philippidis + + * oacc-ptx.h (GOACC_INTERNAL_PTX): Add GOACC_tid, GOACC_ntid, + GOACC_ctaid, and GOACC_nctaid routines. + 2015-02-11 Jakub Jelinek PR c/64824 diff --git libgomp/oacc-ptx.h libgomp/oacc-ptx.h index 13ff86f..2419a46 100644 --- libgomp/oacc-ptx.h +++ libgomp/oacc-ptx.h @@ -101,9 +101,233 @@ ".version 3.1\n" \ ".target sm_30\n" \ ".address_size 64\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1);\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1);\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1);\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1);\n" \ ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads;\n" \ ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num;\n" \ ".extern .func abort;\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1)\n" \ + "{\n" \ + ".reg .u32 %ar1;\n" \ + ".reg .u32 %retval;\n" \ + ".reg .u64 %hr10;\n" \ + ".reg .u32 %r22;\n" \ + ".reg .u32 %r23;\n" \ + ".reg .u32 %r24;\n" \ + ".reg .u32 %r25;\n" \ + ".reg .u32 %r26;\n" \ + ".reg .u32 %r27;\n" \ + ".reg .u32 %r28;\n" \ + ".reg .u32 %r29;\n" \ + ".reg .pred %r30;\n" \ + ".reg .u32 %r31;\n" \ + ".reg .pred %r32;\n" \ + ".reg .u32 %r33;\n" \ + ".reg .pred %r34;\n" \ + ".local .align 8 .b8 %frame[4];\n" \ + "ld.param.u32 %ar1,[%in_ar1];\n" \ + "mov.u32 %r27,%ar1;\n" \ + "st.local.u32 [%frame],%r27;\n" \ + "ld.local.u32 %r28,[%frame];\n" \ + "mov.u32 %r29,1;\n" \ + "setp.eq.u32 %r30,%r28,%r29;\n" \ + "@%r30 bra $L4;\n" \ + "mov.u32 %r31,2;\n" \ + "setp.eq.u32 %r32,%r28,%r31;\n" \ + "@%r32 bra $L5;\n" \ + "mov.u32 %r33,0;\n" \ + "setp.eq.u32 %r34,%r28,%r33;\n" \ + "@!%r34 bra $L8;\n" \ + "mov.u32 %r23,%tid.x;\n" \ + "mov.u32 %r22,%r23;\n" \ + "bra $L7;\n" \ + "$L4:\n" \ + "mov.u32 %r24,%tid.y;\n" \ + "mov.u32 %r22,%r24;\n" \ + "bra $L7;\n" \ + "$L5:\n" \ + "mov.u32 %r25,%tid.z;\n" \ + "mov.u32 %r22,%r25;\n" \ + "bra $L7;\n" \ + "$L8:\n" \ + "{\n" \ + "{\n" \ + "call abort;\n" \ + "}\n" \ + "}\n" \ + "$L7:\n" \ + "mov.u32 %r26,%r22;\n" \ + "mov.u32 %retval,%r26;\n" \ + "st.param.u32 [%out_retval],%retval;\n" \ + "ret;\n" \ + "}\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1)\n" \ + "{\n" \ + ".reg .u32 %ar1;\n" \ + ".reg .u32 %retval;\n" \ + ".reg .u64 %hr10;\n" \ + ".reg .u32 %r22;\n" \ + ".reg .u32 %r23;\n" \ + ".reg .u32 %r24;\n" \ + ".reg .u32 %r25;\n" \ + ".reg .u32 %r26;\n" \ + ".reg .u32 %r27;\n" \ + ".reg .u32 %r28;\n" \ + ".reg .u32 %r29;\n" \ + ".reg .pred %r30;\n" \ + ".reg .u32 %r31;\n" \ + ".reg .pred %r32;\n" \ + ".reg .u32 %r33;\n" \ + ".reg .pred %r34;\n" \ + ".local .align 8 .b8 %frame[4];\n" \ + "ld.param.u32 %ar1,[%in_ar1];\n" \ + "mov.u32 %r27,%ar1;\n" \ + "st.local.u32 [%frame],%r27;\n" \ + "ld.local.u32 %r28,[%frame];\n" \ + "mov.u32 %r29,1;\n" \ + "setp.eq.u32 %r30,%r28,%r29;\n" \ + "@%r30 bra $L11;\n" \ + "mov.u32 %r31,2;\n" \ + "setp.eq.u32 %r32,%r28,%r31;\n" \ + "@%r32 bra $L12;\n" \ + "mov.u32 %r33,0;\n" \ + "setp.eq.u32 %r34,%r28,%r33;\n" \ + "@!%r34 bra $L15;\n" \ + "mov.u32 %r23,%ntid.x;\n" \ + "mov.u32 %r22,%r23;\n" \ + "bra $L14;\n" \ + "$L11:\n" \ + "mov.u32 %r24,%ntid.y;\n" \ + "mov.u32 %r22,%r24;\n" \ + "bra $L14;\n" \ + "$L12:\n" \ + "mov.u32 %r25,%ntid.z;\n" \ + "mov.u32 %r22,%r25;\n" \ + "bra $L14;\n" \ + "$L15:\n" \ + "{\n" \ + "{\n" \ + "call abort;\n" \ + "}\n" \ + "}\n" \ + "$L14:\n" \ + "mov.u32 %r26,%r22;\n" \ + "mov.u32 %retval,%r26;\n" \ + "st.param.u32 [%out_retval],%retval;\n" \ + "ret;\n" \ + "}\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1)\n" \ + "{\n" \ + ".reg .u32 %ar1;\n" \ + ".reg .u32 %retval;\n" \ + ".reg .u64 %hr10;\n" \ + ".reg .u32 %r22;\n" \ + ".reg .u32 %r23;\n" \ + ".reg .u32 %r24;\n" \ + ".reg .u32 %r25;\n" \ + ".reg .u32 %r26;\n" \ + ".reg .u32 %r27;\n" \ + ".reg .u32 %r28;\n" \ + ".reg .u32 %r29;\n" \ + ".reg .pred %r30;\n" \ + ".reg .u32 %r31;\n" \ + ".reg .pred %r32;\n" \ + ".reg .u32 %r33;\n" \ + ".reg .pred %r34;\n" \ + ".local .align 8 .b8 %frame[4];\n" \ + "ld.param.u32 %ar1,[%in_ar1];\n" \ + "mov.u32 %r27,%ar1;\n" \ + "st.local.u32 [%frame],%r27;\n" \ + "ld.local.u32 %r28,[%frame];\n" \ + "mov.u32 %r29,1;\n" \ + "setp.eq.u32 %r30,%r28,%r29;\n" \ + "@%r30 bra $L18;\n" \ + "mov.u32 %r31,2;\n" \ + "setp.eq.u32 %r32,%r28,%r31;\n" \ + "@%r32 bra $L19;\n" \ + "mov.u32 %r33,0;\n" \ + "setp.eq.u32 %r34,%r28,%r33;\n" \ + "@!%r34 bra $L22;\n" \ + "mov.u32 %r23,%ctaid.x;\n" \ + "mov.u32 %r22,%r23;\n" \ + "bra $L21;\n" \ + "$L18:\n" \ + "mov.u32 %r24,%ctaid.y;\n" \ + "mov.u32 %r22,%r24;\n" \ + "bra $L21;\n" \ + "$L19:\n" \ + "mov.u32 %r25,%ctaid.z;\n" \ + "mov.u32 %r22,%r25;\n" \ + "bra $L21;\n" \ + "$L22:\n" \ + "{\n" \ + "{\n" \ + "call abort;\n" \ + "}\n" \ + "}\n" \ + "$L21:\n" \ + "mov.u32 %r26,%r22;\n" \ + "mov.u32 %retval,%r26;\n" \ + "st.param.u32 [%out_retval],%retval;\n" \ + "ret;\n" \ + "}\n" \ + ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1)\n" \ + "{\n" \ + ".reg .u32 %ar1;\n" \ + ".reg .u32 %retval;\n" \ + ".reg .u64 %hr10;\n" \ + ".reg .u32 %r22;\n" \ + ".reg .u32 %r23;\n" \ + ".reg .u32 %r24;\n" \ + ".reg .u32 %r25;\n" \ + ".reg .u32 %r26;\n" \ + ".reg .u32 %r27;\n" \ + ".reg .u32 %r28;\n" \ + ".reg .u32 %r29;\n" \ + ".reg .pred %r30;\n" \ + ".reg .u32 %r31;\n" \ + ".reg .pred %r32;\n" \ + ".reg .u32 %r33;\n" \ + ".reg .pred %r34;\n" \ + ".local .align 8 .b8 %frame[4];\n" \ + "ld.param.u32 %ar1,[%in_ar1];\n" \ + "mov.u32 %r27,%ar1;\n" \ + "st.local.u32 [%frame],%r27;\n" \ + "ld.local.u32 %r28,[%frame];\n" \ + "mov.u32 %r29,1;\n" \ + "setp.eq.u32 %r30,%r28,%r29;\n" \ + "@%r30 bra $L25;\n" \ + "mov.u32 %r31,2;\n" \ + "setp.eq.u32 %r32,%r28,%r31;\n" \ + "@%r32 bra $L26;\n" \ + "mov.u32 %r33,0;\n" \ + "setp.eq.u32 %r34,%r28,%r33;\n" \ + "@!%r34 bra $L29;\n" \ + "mov.u32 %r23,%nctaid.x;\n" \ + "mov.u32 %r22,%r23;\n" \ + "bra $L28;\n" \ + "$L25:\n" \ + "mov.u32 %r24,%nctaid.y;\n" \ + "mov.u32 %r22,%r24;\n" \ + "bra $L28;\n" \ + "$L26:\n" \ + "mov.u32 %r25,%nctaid.z;\n" \ + "mov.u32 %r22,%r25;\n" \ + "bra $L28;\n" \ + "$L29:\n" \ + "{\n" \ + "{\n" \ + "call abort;\n" \ + "}\n" \ + "}\n" \ + "$L28:\n" \ + "mov.u32 %r26,%r22;\n" \ + "mov.u32 %retval,%r26;\n" \ + "st.param.u32 [%out_retval],%retval;\n" \ + "ret;\n" \ + "}\n" \ ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads\n" \ "{\n" \ ".reg .u32 %retval;\n" \