* Remove unused openacc call @ 2015-11-09 21:58 Nathan Sidwell 2016-04-20 8:47 ` Thomas Schwinge 0 siblings, 1 reply; 10+ messages in thread From: Nathan Sidwell @ 2015-11-09 21:58 UTC (permalink / raw) To: GCC Patches; +Cc: Thomas Schwinge [-- Attachment #1: Type: text/plain, Size: 165 bytes --] I've committed this to trunk. It nuke the now unused GOACC_GET_NUM_THREADS and GOACC_GET_THREAD_NUM calls. Also fixed up some comment typos I noticed nathan [-- Attachment #2: trunk-remove.patch --] [-- Type: text/x-patch, Size: 3240 bytes --] 2015-11-09 Nathan Sidwell <nathan@codesourcery.com> * omp-low.c: Fix some OpenACC comment typos. (lower_reduction_clauses): Remove BUILT_IN_GOACC_GET_THREAD_NUM call. * omp-builtins.def (BUILT_IN_GOACC_GET_THREAD_NUM, BUILT_IN_GOACC_GET_NUM_THREADS): Delete. Index: omp-low.c =================================================================== --- omp-low.c (revision 230038) +++ omp-low.c (working copy) @@ -5559,7 +5559,7 @@ lower_reduction_clauses (tree clauses, g { gimple_seq sub_seq = NULL; gimple *stmt; - tree x, c, tid = NULL_TREE; + tree x, c; int count = 0; /* OpenACC loop reductions are handled elsewhere. */ @@ -5589,17 +5589,6 @@ lower_reduction_clauses (tree clauses, g if (count == 0) return; - /* Initialize thread info for OpenACC. */ - if (is_gimple_omp_oacc (ctx->stmt)) - { - /* Get the current thread id. */ - tree call = builtin_decl_explicit (BUILT_IN_GOACC_GET_THREAD_NUM); - tid = create_tmp_var (TREE_TYPE (TREE_TYPE (call))); - gimple *stmt = gimple_build_call (call, 0); - gimple_call_set_lhs (stmt, tid); - gimple_seq_add_stmt (stmt_seqp, stmt); - } - for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) { tree var, ref, new_var, orig_var; @@ -12266,7 +12255,7 @@ expand_omp_atomic (struct omp_region *re } -/* Encode an oacc launc argument. This matches the GOMP_LAUNCH_PACK +/* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK macro on gomp-constants.h. We do not check for overflow. */ static tree @@ -12292,7 +12281,7 @@ oacc_launch_pack (unsigned code, tree de The attribute value is a TREE_LIST. A set of dimensions is represented as a list of INTEGER_CST. Those that are runtime - expres are represented as an INTEGER_CST of zero. + exprs are represented as an INTEGER_CST of zero. TOOO. Normally the attribute will just contain a single such list. If however it contains a list of lists, this will represent the use of @@ -14311,7 +14300,7 @@ lower_omp_for (gimple_stmt_iterator *gsi gimple_omp_for_clauses (stmt), &oacc_head, &oacc_tail, ctx); - /* Add OpenACC partitioning markers just before the loop */ + /* Add OpenACC partitioning and reduction markers just before the loop */ if (oacc_head) gimple_seq_add_seq (&body, oacc_head); @@ -19524,7 +19513,7 @@ public: return execute_oacc_device_lower (); } -}; // class pass_oacc_transform +}; // class pass_oacc_device_lower } // anon namespace Index: omp-builtins.def =================================================================== --- omp-builtins.def (revision 230038) +++ omp-builtins.def (working copy) @@ -47,10 +47,6 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait", BT_FN_VOID_INT_INT_VAR, ATTR_NOTHROW_LIST) -DEF_GOACC_BUILTIN (BUILT_IN_GOACC_GET_THREAD_NUM, "GOACC_get_thread_num", - BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) -DEF_GOACC_BUILTIN (BUILT_IN_GOACC_GET_NUM_THREADS, "GOACC_get_num_threads", - BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device", BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST) ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: Remove unused openacc call 2015-11-09 21:58 Remove unused openacc call Nathan Sidwell @ 2016-04-20 8:47 ` Thomas Schwinge 2016-04-20 10:27 ` Jakub Jelinek 0 siblings, 1 reply; 10+ messages in thread From: Thomas Schwinge @ 2016-04-20 8:47 UTC (permalink / raw) To: gcc-patches, Jakub Jelinek; +Cc: Nathan Sidwell [-- Attachment #1: Type: text/plain, Size: 12222 bytes --] Hi! On Mon, 9 Nov 2015 16:58:48 -0500, Nathan Sidwell <nathan@acm.org> wrote: > I've committed this to trunk. It nuke the now unused GOACC_GET_NUM_THREADS and > GOACC_GET_THREAD_NUM calls. > * omp-low.c: [...] > (lower_reduction_clauses): Remove BUILT_IN_GOACC_GET_THREAD_NUM call. > * omp-builtins.def (BUILT_IN_GOACC_GET_THREAD_NUM, > BUILT_IN_GOACC_GET_NUM_THREADS): Delete. Given that in GCC 6 we only provide host fallback execution for executables compiled with GCC 5 (and thus using the legacy entry points), we can further clean this up as follows. OK for gcc-6-branch and trunk? commit a3993b5fae8b430e1dca23179aaa23a6c53ea2fb Author: Thomas Schwinge <thomas@codesourcery.com> Date: Tue Nov 10 16:59:46 2015 +0100 Clean up libgomp GCC 5 legacy support libgomp/ * config/nvptx/oacc-parallel.c: Empty file. * oacc-parallel.c: Add comments to legacy entry points (GCC 5). --- libgomp/config/nvptx/oacc-parallel.c | 358 ----------------------------------- libgomp/oacc-parallel.c | 6 +- 2 files changed, 5 insertions(+), 359 deletions(-) diff --git libgomp/config/nvptx/oacc-parallel.c libgomp/config/nvptx/oacc-parallel.c index 243f3e5..e69de29 100644 --- libgomp/config/nvptx/oacc-parallel.c +++ libgomp/config/nvptx/oacc-parallel.c @@ -1,358 +0,0 @@ -/* OpenACC constructs - - Copyright (C) 2014-2016 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"); diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index 1fdb01d..9fe5020 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -194,7 +194,7 @@ GOACC_parallel_keyed (int device, void (*fn) (void *), acc_dev->openacc.async_set_async_func (acc_async_sync); } -/* Legacy entry point, only provide host execution. */ +/* Legacy entry point (GCC 5). Only provide host fallback execution. */ void GOACC_parallel (int device, void (*fn) (void *), @@ -494,12 +494,16 @@ GOACC_wait (int async, int num_waits, ...) goacc_thread ()->dev->openacc.async_wait_all_async_func (acc_async_noval); } +/* Legacy entry point (GCC 5). */ + int GOACC_get_num_threads (void) { return 1; } +/* Legacy entry point (GCC 5). */ + int GOACC_get_thread_num (void) { Grüße Thomas [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: Remove unused openacc call 2016-04-20 8:47 ` Thomas Schwinge @ 2016-04-20 10:27 ` Jakub Jelinek 2018-07-31 15:55 ` [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c Cesar Philippidis 0 siblings, 1 reply; 10+ messages in thread From: Jakub Jelinek @ 2016-04-20 10:27 UTC (permalink / raw) To: Thomas Schwinge; +Cc: gcc-patches, Nathan Sidwell On Wed, Apr 20, 2016 at 10:46:38AM +0200, Thomas Schwinge wrote: > On Mon, 9 Nov 2015 16:58:48 -0500, Nathan Sidwell <nathan@acm.org> wrote: > > I've committed this to trunk. It nuke the now unused GOACC_GET_NUM_THREADS and > > GOACC_GET_THREAD_NUM calls. > > > * omp-low.c: [...] > > (lower_reduction_clauses): Remove BUILT_IN_GOACC_GET_THREAD_NUM call. > > * omp-builtins.def (BUILT_IN_GOACC_GET_THREAD_NUM, > > BUILT_IN_GOACC_GET_NUM_THREADS): Delete. > > Given that in GCC 6 we only provide host fallback execution for > executables compiled with GCC 5 (and thus using the legacy entry points), > we can further clean this up as follows. OK for gcc-6-branch and trunk? > > commit a3993b5fae8b430e1dca23179aaa23a6c53ea2fb > Author: Thomas Schwinge <thomas@codesourcery.com> > Date: Tue Nov 10 16:59:46 2015 +0100 > > Clean up libgomp GCC 5 legacy support > > libgomp/ > * config/nvptx/oacc-parallel.c: Empty file. Ok for trunk and 6.2, we don't need this for 6.1. Jakub ^ permalink raw reply [flat|nested] 10+ messages in thread
* [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c @ 2018-07-31 15:55 ` Cesar Philippidis 2018-08-01 11:32 ` Tom de Vries 2019-05-06 8:52 ` [committed] Clean up libgomp GCC 5 legacy support (was: Remove unused openacc call) Thomas Schwinge 0 siblings, 2 replies; 10+ messages in thread From: Cesar Philippidis @ 2018-07-31 15:55 UTC (permalink / raw) To: gcc-patches, Tom de Vries, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 672 bytes --] 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 [-- Attachment #2: 0001-libgomp-Truncate-config-nvptx-oacc-parallel.c.patch --] [-- Type: text/x-patch, Size: 10281 bytes --] [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 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c 2018-07-31 15:55 ` [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c Cesar Philippidis @ 2018-08-01 11:32 ` Tom de Vries 2018-08-01 11:55 ` Jakub Jelinek 2019-05-06 8:52 ` [committed] Clean up libgomp GCC 5 legacy support (was: Remove unused openacc call) Thomas Schwinge 1 sibling, 1 reply; 10+ messages in thread From: Tom de Vries @ 2018-08-01 11:32 UTC (permalink / raw) To: Cesar Philippidis, gcc-patches, Jakub Jelinek 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 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c 2018-08-01 11:32 ` Tom de Vries @ 2018-08-01 11:55 ` Jakub Jelinek 2018-08-01 12:31 ` Tom de Vries 2018-08-01 17:13 ` Nathan Sidwell 0 siblings, 2 replies; 10+ messages in thread From: Jakub Jelinek @ 2018-08-01 11:55 UTC (permalink / raw) To: Tom de Vries; +Cc: Cesar Philippidis, gcc-patches 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 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c 2018-08-01 11:55 ` Jakub Jelinek @ 2018-08-01 12:31 ` Tom de Vries 2018-08-01 17:13 ` Nathan Sidwell 1 sibling, 0 replies; 10+ messages in thread From: Tom de Vries @ 2018-08-01 12:31 UTC (permalink / raw) To: Jakub Jelinek; +Cc: Cesar Philippidis, gcc-patches 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 ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c 2018-08-01 11:55 ` Jakub Jelinek 2018-08-01 12:31 ` Tom de Vries @ 2018-08-01 17:13 ` Nathan Sidwell 2018-08-01 17:34 ` Jakub Jelinek 1 sibling, 1 reply; 10+ messages in thread From: Nathan Sidwell @ 2018-08-01 17:13 UTC (permalink / raw) To: Jakub Jelinek, Tom de Vries; +Cc: Cesar Philippidis, gcc-patches 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 -- Nathan Sidwell ^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c 2018-08-01 17:13 ` Nathan Sidwell @ 2018-08-01 17:34 ` Jakub Jelinek 0 siblings, 0 replies; 10+ messages in thread From: Jakub Jelinek @ 2018-08-01 17:34 UTC (permalink / raw) To: Nathan Sidwell; +Cc: Tom de Vries, Cesar Philippidis, gcc-patches 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 ^ permalink raw reply [flat|nested] 10+ messages in thread
* [committed] Clean up libgomp GCC 5 legacy support (was: Remove unused openacc call) 2018-07-31 15:55 ` [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c Cesar Philippidis 2018-08-01 11:32 ` Tom de Vries @ 2019-05-06 8:52 ` Thomas Schwinge 1 sibling, 0 replies; 10+ messages in thread From: Thomas Schwinge @ 2019-05-06 8:52 UTC (permalink / raw) To: gcc-patches; +Cc: Jakub Jelinek, Tom de Vries [-- Attachment #1.1: Type: text/plain, Size: 698 bytes --] Hi! On Wed, 20 Apr 2016 12:27:27 +0200, Jakub Jelinek <jakub@redhat.com> wrote: > On Wed, Apr 20, 2016 at 10:46:38AM +0200, Thomas Schwinge wrote: > > Clean up libgomp GCC 5 legacy support > > > > libgomp/ > > * config/nvptx/oacc-parallel.c: Empty file. | * oacc-parallel.c: Add comments to legacy entry points (GCC 5). > > Ok for trunk and 6.2, we don't need this for 6.1. I don't know anymore why I didn't commit this back then. Cesar has since committed to trunk in r263236 the patch that truncates 'libgomp/config/nvptx/oacc-parallel.c', and I've now committed the documentation update to trunk in r270901, see attached. Grüße Thomas [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #1.2: 0001-Clean-up-libgomp-GCC-5-legacy-support.trunk.patch --] [-- Type: text/x-diff, Size: 1745 bytes --] From 75a3472e37d26152afb7b08a6fda2691c4b0ee74 Mon Sep 17 00:00:00 2001 From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Mon, 6 May 2019 08:49:55 +0000 Subject: [PATCH] Clean up libgomp GCC 5 legacy support libgomp/ * oacc-parallel.c: Add comments to legacy entry points (GCC 5). git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@270901 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog | 4 ++++ libgomp/oacc-parallel.c | 6 +++++- 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 30bb2d4beba..64e0a8ad8df 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2019-05-06 Thomas Schwinge <thomas@codesourcery.com> + + * oacc-parallel.c: Add comments to legacy entry points (GCC 5). + 2019-03-27 Kevin Buettner <kevinb@redhat.com> * team.c (gomp_team_start): Initialize pool->threads[0]. diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index b77c5e8b9c5..f5fb63c5b5a 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -275,7 +275,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), acc_dev->openacc.async_set_async_func (acc_async_sync); } -/* Legacy entry point, only provide host execution. */ +/* Legacy entry point (GCC 5). Only provide host fallback execution. */ void GOACC_parallel (int flags_m, void (*fn) (void *), @@ -649,12 +649,16 @@ GOACC_wait (int async, int num_waits, ...) acc_wait_all_async (async); } +/* Legacy entry point (GCC 5). */ + int GOACC_get_num_threads (void) { return 1; } +/* Legacy entry point (GCC 5). */ + int GOACC_get_thread_num (void) { -- 2.17.1 [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 658 bytes --] ^ permalink raw reply [flat|nested] 10+ messages in thread
end of thread, other threads:[~2019-05-06 8:52 UTC | newest] Thread overview: 10+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2015-11-09 21:58 Remove unused openacc call Nathan Sidwell 2016-04-20 8:47 ` Thomas Schwinge 2016-04-20 10:27 ` Jakub Jelinek 2018-07-31 15:55 ` [PATCH,nvptx] Truncate config/nvptx/oacc-parallel.c Cesar Philippidis 2018-08-01 11:32 ` Tom de Vries 2018-08-01 11:55 ` Jakub Jelinek 2018-08-01 12:31 ` Tom de Vries 2018-08-01 17:13 ` Nathan Sidwell 2018-08-01 17:34 ` Jakub Jelinek 2019-05-06 8:52 ` [committed] Clean up libgomp GCC 5 legacy support (was: Remove unused openacc call) Thomas Schwinge
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).