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
[PATCH] [libgomp] Truncate config/nvptx/oacc-parallel.c 2018-XX-YY Cesar Philippidis <ce...@codesourcery.com> Thomas Schwinge <tho...@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
