Hi! I'm sorry for the very late answer!
On Thu, 7 Nov 2013 22:10:15 +0400, Evgeny Gavrin <e.gav...@samsung.com> wrote: > I've looked at your changes and in most of front-end parts it looks > reasonable to me. As you know, we're like-minded with you about how OpenACC's > front-ends should look like. So, I think it's good that you have working flow > for your implementation. Now we can start to merge front-ends from > openacc-1_0-branch to GOMP, while keeping all understandings of ACC in sync > and working. > > Jakub, don't you mind if I prepare a few front-end patches likewise for > review to extend Thomas' vision with the things that was already done by my > colleagues on the ACC branch? I hope it'll much speed up the development. I'd be happy if you did that, sure! (And, as Jakub didn't argue against this, but has allowed for OpenACC development to happen on the gomp-4_0-branch, I suppose he's fine with that, too.) > > the last patch adds GOACC_parallel, which so far simply branches to > > GOMP_target. There's more to come. > I've reviewed ACC-related code from [openacc-1_0-branch] and from the > patches, I noticed there are some mess around marking OpenACC specific code. > There are ACC, GACC, OACC, GOACC abbreviations used across code, which one is > the best? I prefer ACC, this is not critical, I'm just curious how to resolve > this. In the "GNU or GCC world", generally the G prefix is used for any additions/extensions/modifications that are, for example, part of the external GNU ABI but not in the standard. So, for example, in libgomp there is a omp_get_team_size symbol exported as that function name is mandated by the OpenMP standard. On the other hand, for the internal »#pragma omp parallel« implementation, libgomp exports a GNU-only GOMP_parallel symbol, which GCC knows and uses. The same scheme I started using for OpenACC. As for ACC vs. OACC, my reasoning is two-fold: OACC makes it look nicely next to OMP ;-), and there are already quite some ACC* things in the compiler (meaning: accept, access, account, accumulate, ...), making it a bit hard to grep for ACC* things -- a problem we don't have with OACC. I don't insist on this however: can for example also use ACC or maybe even OPENACC if there's consensus that'd be better. > > This patch series doesn't contain any substantial rumtime library work > > yet; > Can you share some technical details/ideas about further implementation, it's > not going to be trivial. On the [openacc-1_0-branch] we resolved it with > OpenCL-driven runtime. Also, OpenACC 2.0 standard declares some > vendor-specific routines, this should be noted too. Yes, we're working on some kind of design sketch similar to what you sent for the OpenACC to OpenCL translation. I think a page on the GCC wiki would be a good place to keep this information for the moment, and then later transfer it into the (internals) documentation of libgomp (or any new library). I have just created a stub page, <http://gcc.gnu.org/wiki/OpenACC>, and will add to that later. > > This is in contrast to Samsung's work, who are implementing OpenACC > > separately from the existing OpenMP support. > I'm not fully agree that this in contrast to things are taking place in > [openacc-1_0-branch]. The main reason for keeping middle-end and back-end > solutions far from GOMP is the understanding of OpenACC semantics and unclear > understanding of what should be done to generalize support of accelerators in > GCC. OpenACC and OpenOMP has some semantically differences, but such things > matters only for middle-end and further code generation. The main one, is the > memory concept. The My approach was (and still is), to first use what we already have (for example, GOMP_target/»#pragma omp target« for implementing #pragma acc parallel«), so that we get something that does something, and then extend/rework that to implement what OpenACC actually requires, for example. > > Yet, I hope we'll be able to share/re-use at least front end and some > middle end code. > I absolutely agree, it's needed to merge our efforts regarding front-ends. > Many things from ACC in [openacc-1_0-branch] is already done, so it doesn't > necessary to re-implement same things. Maybe we can polish front-ends > together and then commit changes to GOMP? Yep, that'd be good to do. > > 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. > How do you think, does it make sense to adopt current Fortran implementation > from the OpenACC branch? This can be done quite soon. Also, I may help to > extend C and C++ front-end the same way it's done now, by adding other > implemented directives, too. Sure, that sounds good. > Regarding proposed patches: > [8/9] c_parser_omp_all_clauses and cp_parser_omp_all_clauses > - c_parser_error (parser, "expected %<#pragma omp%> clause"); > + c_parser_error (parser, "expected clause"); > > - cp_parser_error (parser, "expected %<#pragma omp%> clause"); > + cp_parser_error (parser, "expected clause"); > Does it really needed to remove %<#pragma omp%> in these cases? We handle > both, I think. Hmm, how do you mean »both«? I originally thought this code would be shared with OpenACC -- but maybe that's actually not the best approach here, so I might revert these two changes later. > [9/9] > * gimple.h (gimple_build_oacc_*): New declaration. > gimple_oacc_parallel_* and other functions like that. Let's move these > functions out of gimple.* to gimple-oacc.*! There are going to be much more > of the stuff like this, maybe it'll make sense to keep it externally? May be > not only gimple manipulation functions but also pretty-print Also, have a > look on openacc_1-0_branch please, for gimple-oacc.*, is OK if I'll add some > of this? I'll have a look. Again, here I just followed the pattern how things were done for OpenMP, and added to that. But keeping all these little helper functions in a separate file may be a sensible thing to do. > Handling clauses. > + GIMPLE_CHECK (gs, GIMPLE_OACC_PARALLEL); > + gs->gimple_omp_parallel.clauses = clauses; > I think, we can freely add some kind of flag to gimple_omp_parallel to > indicate that we're are working with ACC. > Or even add new gimple_acc_* statements. We can, but so far I had no need to, so I just used what was already there instead of adding new infrastructure that does the same thing as the existing one. > On gimplification, lowering stuff and further, I think it's okay if targeting > GOMP_target, but on my view, this is going to prevent us from emiting OpenCL > or do some extended analysis of the code for error analysis to be offloaded. > May be introduce some switch that turns on OpenCL generation or > GOMP_targeting dependently on the state? Am I wrong assuming that targeting OpenCL would be another instance of the "offloading" functionality we're currently discussing elsewhere? The idea is to stream out (LTO streaming) the intermediate representation of the structured block attached to a OpenACC/OpenMP/Cilk Plus/... directive, and then have the GCC driver invoke the lto1 targeting the acceleration device's architecture, which will then read it in again, and process it further. For example, in our case, this would be a nvptx-none lto1 that then produces PTX code, which is then fed back into the "main" GCC driver, and embedded (as a text string, basically) into the x86-64 object file that this one is creating. Would it make sense to see OpenCL as another target "architecture"; does it fit into this design, or should this be handled differently? > I hope my colleagues will correct me, if I forgot something. Should I, by the way, always add all of you to the CC field when I send emails that I think your group might be interested in? (No problem doing that, if you'd like me to.) Grüße, Thomas
pgpSrDINn94YV.pgp
Description: PGP signature
