diff mbox series

[nvptx] Truncate config/nvptx/oacc-parallel.c

Message ID 1e40ad5b-dfc2-fc93-f4df-424fc5e468aa@codesourcery.com
State New
Headers show
Series [nvptx] Truncate config/nvptx/oacc-parallel.c | expand

Commit Message

Cesar Philippidis July 31, 2018, 3:55 p.m. UTC 
Way back in the GCC 5 days when support for OpenACC was in its infancy,
we used to rely on having various GOACC_ thread functions in the runtime
to implement the execution model, or there lack of (that version of GCC
only supported vector level parallelism). However, beginning with GCC 6,
those external functions were replaced with internal functions that get
expanded by the nvptx BE directly.

This patch removes those stale libgomp functions from the nvptx libgomp
target. Is this OK for trunk, or does libgomp still need to maintain
backwards compatibility with GCC 5?

This patch has been bootstrapped and regtested for x86_64 with nvptx
offloading.

Thanks,
Cesar

Comments

Tom de Vries Aug. 1, 2018, 11:33 a.m. UTC | #1 
On 07/31/2018 05:55 PM, Cesar Philippidis wrote:
> Way back in the GCC 5 days when support for OpenACC was in its infancy,
> we used to rely on having various GOACC_ thread functions in the runtime
> to implement the execution model, or there lack of (that version of GCC
> only supported vector level parallelism). However, beginning with GCC 6,
> those external functions were replaced with internal functions that get
> expanded by the nvptx BE directly.
> 
> This patch removes those stale libgomp functions from the nvptx libgomp
> target. Is this OK for trunk, or does libgomp still need to maintain
> backwards compatibility with GCC 5?
> 
> This patch has been bootstrapped and regtested for x86_64 with nvptx
> offloading.

AFAIU, if you use a GCC 5 nvptx offloading compiler that generates calls
to these GOACC_ thread functions, you're also expected to use the GCC 5
nvptx libgomp.a containing these functions, so I don't see any backwards
compatibility issues here.

OK for me.

Jakub, do you have an opinion here?

Thanks,
- Tom
Jakub Jelinek Aug. 1, 2018, 11:55 a.m. UTC | #2 
On Wed, Aug 01, 2018 at 01:33:09PM +0200, Tom de Vries wrote:
> On 07/31/2018 05:55 PM, Cesar Philippidis wrote:
> > Way back in the GCC 5 days when support for OpenACC was in its infancy,
> > we used to rely on having various GOACC_ thread functions in the runtime
> > to implement the execution model, or there lack of (that version of GCC
> > only supported vector level parallelism). However, beginning with GCC 6,
> > those external functions were replaced with internal functions that get
> > expanded by the nvptx BE directly.
> > 
> > This patch removes those stale libgomp functions from the nvptx libgomp
> > target. Is this OK for trunk, or does libgomp still need to maintain
> > backwards compatibility with GCC 5?
> > 
> > This patch has been bootstrapped and regtested for x86_64 with nvptx
> > offloading.
> 
> AFAIU, if you use a GCC 5 nvptx offloading compiler that generates calls
> to these GOACC_ thread functions, you're also expected to use the GCC 5
> nvptx libgomp.a containing these functions, so I don't see any backwards
> compatibility issues here.
> 
> OK for me.
> 
> Jakub, do you have an opinion here?

The ABI compatibility is mainly for libgomp.so which hasn't (ever) bumped
the soname and I don't plan to do that any time soon, but even for the
offloaded libgomp.a I guess one might compile with GCC 5 and link with GCC
9 and expect things not to fail miserably.  This is a *.a library, can't you
e.g. move those functions to a separate *.c file so that they aren't linked
in unless GCC 5 is really used?

	Jakub
Tom de Vries Aug. 1, 2018, 12:31 p.m. UTC | #3 
On 08/01/2018 01:55 PM, Jakub Jelinek wrote:
> On Wed, Aug 01, 2018 at 01:33:09PM +0200, Tom de Vries wrote:
>> On 07/31/2018 05:55 PM, Cesar Philippidis wrote:
>>> Way back in the GCC 5 days when support for OpenACC was in its infancy,
>>> we used to rely on having various GOACC_ thread functions in the runtime
>>> to implement the execution model, or there lack of (that version of GCC
>>> only supported vector level parallelism). However, beginning with GCC 6,
>>> those external functions were replaced with internal functions that get
>>> expanded by the nvptx BE directly.
>>>
>>> This patch removes those stale libgomp functions from the nvptx libgomp
>>> target. Is this OK for trunk, or does libgomp still need to maintain
>>> backwards compatibility with GCC 5?
>>>
>>> This patch has been bootstrapped and regtested for x86_64 with nvptx
>>> offloading.
>>
>> AFAIU, if you use a GCC 5 nvptx offloading compiler that generates calls
>> to these GOACC_ thread functions, you're also expected to use the GCC 5
>> nvptx libgomp.a containing these functions, so I don't see any backwards
>> compatibility issues here.
>>
>> OK for me.
>>
>> Jakub, do you have an opinion here?
> 
> The ABI compatibility is mainly for libgomp.so which hasn't (ever) bumped
> the soname and I don't plan to do that any time soon, but even for the
> offloaded libgomp.a I guess one might compile with GCC 5 and link with GCC
> 9 and expect things not to fail miserably.  This is a *.a library, can't you
> e.g. move those functions to a separate *.c file so that they aren't linked
> in unless GCC 5 is really used?

You're describing the current situation: all those functions sit
together in config/nvptx/oacc-parallel.c. [ Besides, the nvptx .o and .a
files are marked up text files containing ptx functions, so I'm not sure
if the linker is obliged to pull in entire .o files. ]

Thanks,
- Tom
Nathan Sidwell Aug. 1, 2018, 5:13 p.m. UTC | #4 
On 08/01/2018 04:55 AM, Jakub Jelinek wrote:

> The ABI compatibility is mainly for libgomp.so which hasn't (ever) bumped
> the soname and I don't plan to do that any time soon, but even for the
> offloaded libgomp.a I guess one might compile with GCC 5 and link with GCC
> 9 and expect things not to fail miserably.  This is a *.a library, can't you

I think it should fail horribly.  If it succeeded, the performance would 
suck.  Far better to shout at the user sooner.

nathan
Jakub Jelinek Aug. 1, 2018, 5:34 p.m. UTC | #5 
On Wed, Aug 01, 2018 at 10:13:06AM -0700, Nathan Sidwell wrote:
> On 08/01/2018 04:55 AM, Jakub Jelinek wrote:
> 
> > The ABI compatibility is mainly for libgomp.so which hasn't (ever) bumped
> > the soname and I don't plan to do that any time soon, but even for the
> > offloaded libgomp.a I guess one might compile with GCC 5 and link with GCC
> > 9 and expect things not to fail miserably.  This is a *.a library, can't you
> 
> I think it should fail horribly.  If it succeeded, the performance would
> suck.  Far better to shout at the user sooner.

Ok, I can live with that for OpenACC and GCC 5, just would not like to force
people to rebuild everything every time GCC is bumped even for OpenACC (and
once exported symbols from libgomp.so need to stay forever, unless soname is
bumnped).

	Jakub
diff mbox series

Patch

[PATCH] [libgomp] Truncate config/nvptx/oacc-parallel.c

2018-XX-YY  Cesar Philippidis  <cesar@codesourcery.com>
	    Thomas Schwinge <thomas@codesourcery.com>

	libgomp/
	* config/nvptx/oacc-parallel.c: Truncate.

(cherry picked from gomp-4_0-branch r228836)
---
 libgomp/config/nvptx/oacc-parallel.c | 358 -----------------------------------
 1 file changed, 358 deletions(-)

diff --git a/libgomp/config/nvptx/oacc-parallel.c b/libgomp/config/nvptx/oacc-parallel.c
index 5dc53da..e69de29 100644
--- a/libgomp/config/nvptx/oacc-parallel.c
+++ b/libgomp/config/nvptx/oacc-parallel.c
@@ -1,358 +0,0 @@ 
-/* OpenACC constructs
-
-   Copyright (C) 2014-2018 Free Software Foundation, Inc.
-
-   Contributed by Mentor Embedded.
-
-   This file is part of the GNU Offloading and Multi Processing Library
-   (libgomp).
-
-   Libgomp is free software; you can redistribute it and/or modify it
-   under the terms of the GNU General Public License as published by
-   the Free Software Foundation; either version 3, or (at your option)
-   any later version.
-
-   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
-   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
-   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
-   more details.
-
-   Under Section 7 of GPL version 3, you are granted additional
-   permissions described in the GCC Runtime Library Exception, version
-   3.1, as published by the Free Software Foundation.
-
-   You should have received a copy of the GNU General Public License and
-   a copy of the GCC Runtime Library Exception along with this program;
-   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
-   <http://www.gnu.org/licenses/>.  */
-
-#include "libgomp_g.h"
-
-__asm__ (".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"
-	 "// BEGIN GLOBAL FUNCTION DECL: GOACC_get_num_threads\n"
-	 ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads;\n"
-	 "// BEGIN GLOBAL FUNCTION DECL: GOACC_get_thread_num\n"
-	 ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num;\n"
-	 "// BEGIN GLOBAL FUNCTION DECL: abort\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"
-	 "// BEGIN GLOBAL FUNCTION DEF: GOACC_get_num_threads\n"
-	 ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads\n"
-	 "{\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"
-	 "mov.u32 %r26,0;\n"
-	 "{\n"
-	 ".param .u32 %retval_in;\n"
-	 "{\n"
-	 ".param .u32 %out_arg0;\n"
-	 "st.param.u32 [%out_arg0],%r26;\n"
-	 "call (%retval_in),GOACC_ntid,(%out_arg0);\n"
-	 "}\n"
-	 "ld.param.u32 %r27,[%retval_in];\n"
-	 "}\n"
-	 "mov.u32 %r22,%r27;\n"
-	 "mov.u32 %r28,0;\n"
-	 "{\n"
-	 ".param .u32 %retval_in;\n"
-	 "{\n"
-	 ".param .u32 %out_arg0;\n"
-	 "st.param.u32 [%out_arg0],%r28;\n"
-	 "call (%retval_in),GOACC_nctaid,(%out_arg0);\n"
-	 "}\n"
-	 "ld.param.u32 %r29,[%retval_in];\n"
-	 "}\n"
-	 "mov.u32 %r23,%r29;\n"
-	 "mul.lo.u32 %r24,%r22,%r23;\n"
-	 "mov.u32 %r25,%r24;\n"
-	 "mov.u32 %retval,%r25;\n"
-	 "st.param.u32 [%out_retval],%retval;\n"
-	 "ret;\n"
-	 "}\n"
-	 "// BEGIN GLOBAL FUNCTION DEF: GOACC_get_thread_num\n"
-	 ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num\n"
-	 "{\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 .u32 %r30;\n"
-	 ".reg .u32 %r31;\n"
-	 ".reg .u32 %r32;\n"
-	 ".reg .u32 %r33;\n"
-	 "mov.u32 %r28,0;\n"
-	 "{\n"
-	 ".param .u32 %retval_in;\n"
-	 "{\n"
-	 ".param .u32 %out_arg0;\n"
-	 "st.param.u32 [%out_arg0],%r28;\n"
-	 "call (%retval_in),GOACC_ntid,(%out_arg0);\n"
-	 "}\n"
-	 "ld.param.u32 %r29,[%retval_in];\n"
-	 "}\n"
-	 "mov.u32 %r22,%r29;\n"
-	 "mov.u32 %r30,0;\n"
-	 "{\n"
-	 ".param .u32 %retval_in;\n"
-	 "{\n"
-	 ".param .u32 %out_arg0;\n"
-	 "st.param.u32 [%out_arg0],%r30;\n"
-	 "call (%retval_in),GOACC_ctaid,(%out_arg0);\n"
-	 "}\n"
-	 "ld.param.u32 %r31,[%retval_in];\n"
-	 "}\n"
-	 "mov.u32 %r23,%r31;\n"
-	 "mul.lo.u32 %r24,%r22,%r23;\n"
-	 "mov.u32 %r32,0;\n"
-	 "{\n"
-	 ".param .u32 %retval_in;\n"
-	 "{\n"
-	 ".param .u32 %out_arg0;\n"
-	 "st.param.u32 [%out_arg0],%r32;\n"
-	 "call (%retval_in),GOACC_tid,(%out_arg0);\n"
-	 "}\n"
-	 "ld.param.u32 %r33,[%retval_in];\n"
-	 "}\n"
-	 "mov.u32 %r25,%r33;\n"
-	 "add.u32 %r26,%r24,%r25;\n"
-	 "mov.u32 %r27,%r26;\n"
-	 "mov.u32 %retval,%r27;\n"
-	 "st.param.u32 [%out_retval],%retval;\n"
-	 "ret;\n"
-	 "}\n");
-- 
2.7.4