public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* Initial submission of OpenACC support integrated into OpenMP's lowering and expansion passes
@ 2013-11-06 19:44 Thomas Schwinge
  2013-11-06 19:44 ` [gomp4 1/9] Add missing include thomas
  2013-11-07 18:14 ` Initial submission of OpenACC support integrated into OpenMP's lowering and expansion passes Evgeny Gavrin
  0 siblings, 2 replies; 44+ messages in thread
From: Thomas Schwinge @ 2013-11-06 19:44 UTC (permalink / raw)
  To: gcc-patches; +Cc: jakub, e.gavrin

[-- Attachment #1: Type: text/plain, Size: 2906 bytes --]

Hi!

Here is our take of adding OpenACC support to GCC's front end and middle
end.  What I post here (that is, as replies to this email, to gcc-patches
only) is still very incomplete, but hopefully enough to outline the
general approach, and I'd like to commit these patches to gomp-4_0-branch
-- OK?  Also, may I additionally commit those patches to trunk that don't
have anything specific to OpenACC in them (such as the first two
preparation patches), which I assume will save some work when merging
between trunk and gomp-4_0-branch later on?

Due to the conceptual similarity compared to OpenMP, and that (later) it
is reasonable to expect to embed OpenACC directives into OpenMP ones, the
approach I've chosen is to directly embed OpenACC handling into the
existing OpenMP infrastructure/passes.  I had first begun implementing
separate OpenACC lowering and expansion passes, but then found myself
re-implementing/copying more and more of the existing OpenMP passes'
code, realized that makes no sense, and we shall instead handle these
directly together in one set of passes -- which also happens to
facilitate the interoperation requirement.

For the same reasons, OpenACC's runtime library shall be implemented in
libgomp next to OpenMP's.  Jakub's already approved this approach
generally.  This patch series doesn't contain any substantial rumtime
library work yet; the last patch adds GOACC_parallel, which so far simply
branches to GOMP_target.  There's more to come.

This is in contrast to Samsung's work, who are implementing OpenACC
separately from the existing OpenMP support.  Yet, I hope we'll be able
to share/re-use at least front end and some middle end code.

Our work builds on what has (briefly) been discussed before, for example
at the GNU Tools Cauldron's Accelerator BOF in this summer.  For the
moment, while still getting the concepts clear, I'm adding new
GENERIC/GIMPLE codes and their handling for OpenACC constructs (for the
implementation of OpenACC parallel "borrowing" from OpenMP target) -- the
long-term goal still is to evolve all these into general "acceleration"
abstractions, evolve the GCC OpenMP infrastructure and passes into
general "acceleration" passes.  As discussed, it makes sense to do this
generalization work later on, once we have a better understanding of what
we actually need.

We directly strive for OpenACC 2.0 support, skipping OpenACC 1.  We're
focussing on the C front end implementation first, following on with C++
and Fortran later on.

The patches I send here only contain a bare minimum of documentation
changes.  I already have further documentation changes prepared (source
code comments as well as Texinfo manual updates); will polish and send
these only once the approach we've chosen has generally been approved.

Comments on the approach and/or patches?


Grüße,
 Thomas

[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]

^ permalink raw reply	[flat|nested] 44+ messages in thread

end of thread, other threads:[~2017-09-28 12:12 UTC | newest]

Thread overview: 44+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2013-11-06 19:44 Initial submission of OpenACC support integrated into OpenMP's lowering and expansion passes Thomas Schwinge
2013-11-06 19:44 ` [gomp4 1/9] Add missing include thomas
2013-11-06 19:44   ` [gomp4 2/9] libgomp: Prepare for testcases without -fopenmp thomas
2013-11-06 19:47     ` [gomp4 3/9] OpenACC: Recognize -fopenacc thomas
2013-11-06 20:39       ` [gomp4 4/9] OpenACC: The runtime library will be implemented in libgomp, too thomas
2013-11-06 19:45         ` [gomp4 5/9] OpenACC: preprocessor definition, Fortran integer parameter thomas
2013-11-06 20:37           ` [gomp4 6/9] OpenACC: Infrastructure for builtins thomas
2013-11-06 19:56             ` [gomp4 7/9] OpenACC: Use OpenMP's lowering and expansion passes thomas
2013-11-06 20:29               ` [gomp4 8/9] OpenACC: Basic support for #pragma acc in the C front end thomas
2013-11-06 19:53                 ` [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel thomas
2013-11-06 20:01                   ` Thomas Schwinge
2013-11-07  8:48                     ` Jakub Jelinek
2014-02-27 19:58                     ` [gomp4] Gimplification: Merge gimplify_oacc_parallel into gimplify_omp_workshare. (was: [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel.) Thomas Schwinge
2014-03-20 14:28                     ` [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel Thomas Schwinge
2014-11-12 13:35                     ` OpenACC GIMPLE_OACC_* -- or not? (was: [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel.) Thomas Schwinge
2014-11-12 13:46                       ` Jakub Jelinek
2014-12-10  9:58                         ` OpenACC GIMPLE_OACC_* -- or not? Thomas Schwinge
2014-12-10 10:07                           ` Thomas Schwinge
2014-12-10 10:10                             ` Jakub Jelinek
2014-12-12 20:02                               ` Thomas Schwinge
2014-12-10 10:02                         ` Nested OpenACC/OpenMP constructs (was: OpenACC GIMPLE_OACC_* -- or not?) Thomas Schwinge
2014-12-10 10:10                           ` Thomas Schwinge
2014-12-10 10:16                             ` Jakub Jelinek
2014-12-12 20:04                               ` Nested OpenACC/OpenMP constructs Thomas Schwinge
2014-12-17 22:21                                 ` OpenACC/OpenMP testing: Revise variable usage in constructs (was: Nested OpenACC/OpenMP constructs) Thomas Schwinge
2013-11-06 20:03                   ` [gomp4 9/9] OpenACC: Basic support for #pragma acc parallel Jakub Jelinek
2013-11-07 19:00                     ` Thomas Schwinge
2013-11-07  8:41                 ` [gomp4 8/9] OpenACC: Basic support for #pragma acc in the C front end Jakub Jelinek
2013-11-07  8:39               ` [gomp4 7/9] OpenACC: Use OpenMP's lowering and expansion passes Jakub Jelinek
2013-11-07  8:23             ` [gomp4 6/9] OpenACC: Infrastructure for builtins Jakub Jelinek
2013-11-07  8:19           ` [gomp4 5/9] OpenACC: preprocessor definition, Fortran integer parameter Jakub Jelinek
2013-11-07  8:18         ` [gomp4 4/9] OpenACC: The runtime library will be implemented in libgomp, too Jakub Jelinek
2013-11-07 14:05           ` Generally link to libgomp for -ftree-parallelize-loops=* (was: [gomp4 4/9] OpenACC: The runtime library will be implemented in libgomp, too.) Thomas Schwinge
2013-11-07 14:48             ` Jakub Jelinek
2013-11-07  8:17       ` [gomp4 3/9] OpenACC: Recognize -fopenacc Jakub Jelinek
2013-11-07  8:16     ` [gomp4 2/9] libgomp: Prepare for testcases without -fopenmp Jakub Jelinek
2017-09-28  7:41       ` [libgomp, testsuite] Remove superfluous -fopenmp from libgomp testcases Tom de Vries
2017-09-28  8:37         ` Thomas Schwinge
2017-09-28 12:12           ` Tom de Vries
2017-09-28  8:46         ` Jakub Jelinek
2013-11-07  8:15   ` [gomp4 1/9] Add missing include Jakub Jelinek
2013-11-07 18:14 ` Initial submission of OpenACC support integrated into OpenMP's lowering and expansion passes Evgeny Gavrin
2013-11-27 23:23   ` Thomas Schwinge
     [not found]   ` <004b01cedc68$a5f3c950$f1db5bf0$%b@samsung.com>
2013-11-27 23:25     ` 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).