* 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).