Hi, all!
My colleagues and I are going to share our initial support of OpenACC soon
(in a week, I hope). We're on stage of internal approval before making first
public commit, FE+BE patches are prepared already.
So, I think we need to discuss our future collaboration in case we all want
to implement ACC.
Our current status is the following:
- Fortran finished and applies OpenACC 1.0 specs - we're testing it
right now.
- C is mostly finished - in development
- C++ is on initial stage - in development
- Run-time library is finished
- Also we have prototype of OpenCL code generator as back-end
solution
- Tests for FE generated by script, and reduced to acceptable count.
For front-ends it was decided not to touch exciting OpenMP implementation
for now for the reason:
- perform refactoring with achieved understanding of what should be
done will be much easier, than merge it after global redesign and some bunch
of further development
I agree that there are some front-end specific things that can be reused by
OpenACC, but most of them are related to clauses processing.
For back-end, as was discussed earlier, we developed prototype of code
generator that emits OpenCL code for kernel/parallel directives and for
copy* clauses.
For today we can compile and execute simple test like these ones on GPU
using Intel's and Nvidia's CL drivers:
1 #include <stdio.h>
2
3 #define SIZE 64
4
5 int main() {
6 int i;
7 float a[SIZE], b[SIZE], c[SIZE];
8
9 for (i = 0; i < SIZE; i++) {
10 a[i] = (i + 1) / 10.0;
11 b[i] = (SIZE - i) / 10.0;
12 }
13
14 for (i = 0; i < SIZE; i++) {
15 printf ("%f ", a[i]);
16 }
17 printf ("\n");
18
19 for (i = 0; i < SIZE; i++) {
20 printf ("%f ", b[i]);
21 }
22 printf ("\n");
23
24
25 #pragma acc kernels
26 for (i = 0; i < SIZE; i++) {
27 c[i] = a[i] / b[i];
28 }
29
30 for (i = 0; i < SIZE; i++) {
31 printf ("%f ", c[i]);
32 }
33 printf ("\n");
34
35 printf ("test finished\n");
36 return 0;
37 }
1 PROGRAM test
2 IMPLICIT NONE
3 INTEGER :: i
4 REAL :: a(64), b(64), c(64)
5
6 DO i = 1, 64
7 a(i) = (i)/10.0
8 b(i) = (64 - i + 1)/10.0
9 ENDDO
10
11 PRINT *, a
12 PRINT *, b
13
14 !$ACC KERNELS
15 DO i = 1, 64
16 c(i) = a(i)/b(i)
17 ENDDO
18 !$ACC END KERNELS
19
20 PRINT *, c
21 END PROGRAM test
-
Thanks, Evgeny.
-----Original Message-----
From: gcc-ow...@gcc.gnu.org [mailto:gcc-ow...@gcc.gnu.org] On Behalf Of
Jakub Jelinek
Sent: Thursday, September 05, 2013 7:02 PM
To: Thomas Schwinge
Cc: gcc@gcc.gnu.org
Subject: Re: Reusing OpenMP front end infrastructure for OpenACC -- how?
On Thu, Sep 05, 2013 at 04:51:14PM +0200, Thomas Schwinge wrote:
> Implementing OpenACC support in GCC's frontends, there are things that
> I trivially have to reimplement, that are already present for OpenMP.
> For example, the infrastructure for parsing clauses (as attached to
> OpenACC and OpenMP directives), and their representation in the GCC
> internal data structures. Given the syntactical similarity of OpenACC
> and OpenMP, this infrastructure is similar, too, if not even
> identical. Thus, my inner engineer tells me we ought to share this
> infrastructure instead of having a copy of it only different for
s%omp%oacc.
>
> For a specific example, I'd like to reuse everything known by
> omp_clause* and similar, and extend it with the things OpenACC does
> additionally/differently (adding safe-guards, if not already present,
> to the OpenMP code so that no clauses specific to OpenACC end up there).
Most of the omp-low.c code has asserts to verify no unknown clauses are
passed to it, and during parsing it has bitmasks of allowable clauses.
> Is this generally an accepted approach? If yes, should I rename omp*
> to, say, oacc_omp* (putting the OpenACC tag first for the simple
> reason that it sorts first alphabetically), or can we think of a
> better tag than oacc_omp, or should I leave the names as omp* and add
> comments that these things are used for OpenACC, too?
I think there is no point in renaming the existing stuff, we use it for some
Cilk+ stuff too these days, renaming could only complicate maintainance,
making it harder to backport OpenMP bugfixes to older release branches etc.
IMHO just use from the OpenMP parsing, trees and gimple stuff whatever is
usable for OpenACC too, and just for stuff that doesn't have counterparts
add oacc/OACC stuff.
Say, for clauses IMHO it is just fine if you use OMP_CLAUSE tree code, and
just add additional OACC_CLAUSE_SOMETHING, OACC_CLAUSE_SOMETHINGELSE etc.
to the list of omp clauses and descriptors, just put it likely at the end of
the list with a comment that it is OpenACC specific stuff. If some clause
is used with the same meaning in both standards, you can either just always
use OMP_CLAUSE_WHATEVER, or #define OACC_CLAUSE_WHATEVER
OMP_CLAUSE_WHATEVER.
Jakub
