From patchwork Mon Aug 17 07:31:40 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 507873 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 1C65E14029C for ; Mon, 17 Aug 2015 17:32:04 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=cFgqiMd/; 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:from :to:cc:subject:date:message-id:mime-version:content-type; q=dns; s=default; b=t3GGwmyYs6uCzbXJC2y0HhtgbKm3ZRaNW8wFI6JILZwAVdYz9M vSaxJDfyjUOUixtAq1NCDAu/yzem5YCwcoi7ivT4HuLpD5TMgK4sQXZQ1yZb9tXB Izz6O/ZIJNKfasKCXZ3IGFZKmtBuek0UtmNNEWOuc6YHanvifX49KwOt4= 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:cc:subject:date:message-id:mime-version:content-type; s= default; bh=gWXGLbszM5au/gg4sotu3SAMh5g=; b=cFgqiMd/qu15MSSeyBQW lQVBvUcPZN7iodcDPV0FYVjX7c3g+KsX2xd8yDua2YeEhTYXRPlfAZpbKyRLXDhQ TP3zSIJDuUOSF37fQNbA+SQi9TFzNsW10MrLnjbaGx8fwNLXVPw5rt4l0Kd5XLRZ boxD1Cc+wOQpSMKd7SAspK4= Received: (qmail 33503 invoked by alias); 17 Aug 2015 07:31:56 -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 33490 invoked by uid 89); 17 Aug 2015 07:31:56 -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; Mon, 17 Aug 2015 07:31:54 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZRF2H-000512-PE from Thomas_Schwinge@mentor.com ; Mon, 17 Aug 2015 00:40:50 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Mon, 17 Aug 2015 08:31:49 +0100 From: Thomas Schwinge To: Bernd Schmidt , Nathan Sidwell CC: Subject: [nvptx] Incorrect %retval usage in C++ code User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (x86_64-pc-linux-gnu) Date: Mon, 17 Aug 2015 09:31:40 +0200 Message-ID: <87fv3is74z.fsf@schwinge.name> MIME-Version: 1.0 Hi! I observed that for a (slowly increasing?) number of C++ testcases (gcc/testsuite/g++.*/), their nvptx compilation fails as follows: spawn [...]/build-gcc/gcc/testsuite/g++/../../xg++ -B[...]/build-gcc/gcc/testsuite/g++/../../ [...]/source-gcc/gcc/testsuite/g++.dg/cpp0x/nsdmi4.C -fno-diagnostics-show-caret -fdiagnostics-color=never --sysroot=[...]/install/nvptx-none -fmessage-length=0 -std=c++11 -pedantic-errors -Wno-long-long -DNO_LABEL_VALUES -DNO_TRAMPOLINES -isystem [...]/build-gcc/nvptx-none/./newlib/targ-include -isystem [...]/source-gcc/newlib/libc/include -B[...]/build-gcc/nvptx-none/./newlib/ -L[...]/build-gcc/nvptx-none/./newlib -mmainkernel -lm -o ./nsdmi4.exe ptxas /tmp/cclwz6Zm.o, line 52; error : Arguments mismatch for instruction 'mov' ptxas /tmp/cclwz6Zm.o, line 52; error : Unknown symbol '%retval' ptxas /tmp/cclwz6Zm.o, line 52; error : Label expected for forward reference of '%retval' ptxas fatal : Ptx assembly aborted due to errors nvptx-as: ptxas returned 255 exit status compiler exited with status 1 Reduced from g++.dg/cpp0x/nsdmi4.C: $ < source-gcc/gcc/testsuite/g++.dg/cpp0x/nsdmi4_.C cat struct A { A() { } A(const A&) { } }; A f() { return A(); } $ build-gcc/gcc/xg++ -Bbuild-gcc/gcc/ source-gcc/gcc/testsuite/g++.dg/cpp0x/nsdmi4_.C -std=c++11 -S $ < nsdmi4_.s c++filt // BEGIN PREAMBLE .version 3.1 .target sm_30 .address_size 64 // END PREAMBLE // BEGIN FUNCTION DECL: A::A() .func A::A()(.param.u64 %in_ar1); // BEGIN FUNCTION DEF: A::A() .func A::A()(.param.u64 %in_ar1) { .reg.u64 %ar1; .reg.u64 %hr10; .reg.u64 %r22; .reg.u64 %frame; .local.align 8 .b8 %farray[8]; cvta.local.u64 %frame, %farray; ld.param.u64 %ar1, [%in_ar1]; mov.u64 %r22, %ar1; st.u64 [%frame], %r22; ret; } // BEGIN GLOBAL FUNCTION DECL: f() .visible .func f()(.param.u64 %in_ar1); // BEGIN GLOBAL FUNCTION DEF: f() .visible .func f()(.param.u64 %in_ar1) { .reg.u64 %ar1; .reg.u64 %hr10; .reg.u64 %r22; .reg.u64 %r23; ld.param.u64 %ar1, [%in_ar1]; mov.u64 %r22, %ar1; mov.u64 %r23, %r22; { .param.u64 %out_arg0; st.param.u64 [%out_arg0], %r23; call A::A(), (%out_arg0); } mov.u64 %retval, %r22; ret; } Notice the stray %retval usage very near the end. Note that before r226901, »[PR64164] Drop copyrename, use coalescible partition as base when optimizing.«, , this test case did set up a %frame, and did not assign to %retval: @@ -31,21 +31,18 @@ .reg.u32 %r23; .reg.u64 %r24; .reg.u64 %r25; - .reg.u64 %frame; - .local.align 8 .b8 %farray[8]; - cvta.local.u64 %frame, %farray; ld.param.u64 %ar1, [%in_ar1]; mov.u64 %r24, %ar1; - st.u64 [%frame], %r24; ld.global.u32 %r22, [c]; add.u32 %r23, %r22, 1; st.global.u32 [c], %r23; - ld.u64 %r25, [%frame]; + mov.u64 %r25, %r24; { .param.u64 %out_arg0; st.param.u64 [%out_arg0], %r25; call _ZN1AC1Ev, (%out_arg0); } + mov.u64 %retval, %r24; ret; } But I don't think that this recent commit is directly related to the problem at hand, but it just exposes it some more: before that recent commit, there have already been test cases failing with the same stray %retval usage. I suspect that this mov is incorrectly generated by gcc/function.c:expand_function_end, but can't tell if something's wrong in there, or rather in the nvptx backend's NVPTX_RETURN_REGNUM handling (which I can't claim to really understand), and as I'm unlikely to spend more time on this before leaving for vacations soon, I wanted to dump my state now. Maybe one of you has an idea about this. Also, I guess the following cleanup (untested) is in order: Grüße, Thomas diff --git gcc/config/nvptx/nvptx.h gcc/config/nvptx/nvptx.h index afe4fcd..d846ec3 100644 --- gcc/config/nvptx/nvptx.h +++ gcc/config/nvptx/nvptx.h @@ -103,8 +103,8 @@ enum reg_class #define N_REG_CLASSES (int) LIM_REG_CLASSES #define REG_CLASS_NAMES { \ - "RETURN_REG", \ "NO_REGS", \ + "RETURN_REG", \ "ALL_REGS" } #define REG_CLASS_CONTENTS \ @@ -119,7 +119,7 @@ enum reg_class #define GENERAL_REGS ALL_REGS -#define REGNO_REG_CLASS(R) ((R) == 4 ? RETURN_REG : ALL_REGS) +#define REGNO_REG_CLASS(R) ((R) == NVPTX_RETURN_REGNUM ? RETURN_REG : ALL_REGS) #define BASE_REG_CLASS ALL_REGS #define INDEX_REG_CLASS NO_REGS