<html>
<head>
<meta content="text/html; charset=utf-8" http-equiv="Content-Type">
</head>
<body bgcolor="#FFFFFF" text="#000000">
<p>Hi, Jeff,</p>
<p>First, I don't think we should spam all of cfe-dev the source
listings of every project using OpenACC :-) -- We actually do need
to understand what's out there (both open source and otherwise),
for many different reasons, but this is not the place to collect
that information. Of course, one of the reasons that we'd like
OpenACC support in Clang is to make it easier to do this kind of
analysis (current internal work is using GCC).<br>
</p>
Second, we have a significant body of code using OpenACC, and
moreover, the rate at which OpenACC code is being written is
increasing. As I mentioned, this is because the compiler OpenACC
support has now reached a point where using it makes sense. On many
machines with GPUs, including large DOE machines, OpenACC is the
recommended way to exploit them at this point. As OpenMP accelerator
support matures, that may change. We're (DOE is) investing heavily
in OpenMP for the future. Nevertheless, OpenACC is a cross-vendor
standard (<a class="moz-txt-link-freetext" href="https://www.openacc.org">https://www.openacc.org</a>), which is actively developed,
supported by the Cray and PGI compilers, and also by GCC (as of GCC
v7). The GCC support is new, and less mature than what is provided
by the vendor compilers, but I expect that it will get better over
time as well.<br>
<br>
I think that it makes sense for Clang to support OpenACC. We'd like
to use Clang's tooling capabilities on this code. Moreover, I'd like
to do this in a way that enables maximum interoperability with
OpenMP. Over time, this will make it easier to transition code from
OpenACC to OpenMP when that's desired. There's a significant overlap
in functionality between the two sets of directives. That should
lower the implementation cost significantly. OpenACC is not a pure
subset of OpenMP, however, and that's also important.<br>
<br>
Thanks again,<br>
Hal<br>
<br>
<div class="moz-cite-prefix">On 12/08/2017 12:02 PM, Jeff Hammond
wrote:<br>
</div>
<blockquote
cite="mid:CAGKz=u+KdpSW=T9u0288hQKpNNcEzyw-UFr107TKsy9Av_r7Jw@mail.gmail.com"
type="cite">
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
<div dir="ltr"><br>
<br>
On Fri, Dec 8, 2017 at 9:00 AM, Joel E. Denny <<a
moz-do-not-send="true" href="mailto:jdenny.ornl@gmail.com">jdenny.ornl@gmail.com</a>>
wrote:<br>
><br>
> On Fri, Dec 8, 2017 at 11:32 AM, Jeff Hammond <<a
moz-do-not-send="true" href="mailto:jeff.science@gmail.com">jeff.science@gmail.com</a>>
wrote:<br>
>><br>
>><br>
>><br>
>> On Fri, Dec 8, 2017 at 7:51 AM, Joel E. Denny <<a
moz-do-not-send="true" href="mailto:jdenny.ornl@gmail.com">jdenny.ornl@gmail.com</a>>
wrote:<br>
>> ><br>
>> > Hi Jeff, Hal,<br>
>> ><br>
>> > Thanks for your feedback. My comments are inline
below.<br>
>> ><br>
>> > On Tue, Dec 5, 2017 at 6:43 PM, Hal Finkel <<a
moz-do-not-send="true" href="mailto:hfinkel@anl.gov">hfinkel@anl.gov</a>>
wrote:<br>
>> >><br>
>> >> On 12/05/2017 05:11 PM, Jeff Hammond via
cfe-dev wrote:<br>
>> >><br>
>> >> All of the usage of OpenACC outside of
benchmarks/research that I know about is done in Fortran.<br>
>> ><br>
>> > I agree that it's easier to find real apps that
use OpenACC in Fortran than those that use OpenACC in C/C++.
However, the latter certainly exist. For example:<br>
>><br>
>> Two of the three examples you cite are primarily
Fortran and using OpenACC exclusively in Fortran subroutines.<br>
><br>
><br>
> Are you saying that the occurrences of "pragma acc" in
Nek5000 and NekCEM are unused?<br>
><br>
<br>
The instances of "pragma acc" in those - it's the same code in
both projects - are either (1) only causing host-device data
synchronization or (2) commented-out.
<div><br>
</div>
<div>It's unclear to me what actually happens in the code as
currently written. The OpenACC C/C++ code does not more than
copy data to/from the device. I didn't trace the entire code
execution but I can't tell if any code touches the device data
that OpenACC is updating. If it is updated, it is updated by
Fortran OpenACC code somewhere else in the source tree.<br>
<div><br>
</div>
<div>What does the OpenACC standard say about interoperability
of compilers+runtimes, as would be required if one used
Clang OpenACC for C/C++ and Fortran OpenACC implemented by
PGI, Cray, or GCC. OpenMP definitely does not support this,
even if a subset of usage may work when one uses the same
runtime library with different compilers.</div>
<div><br>
<div>/tmp/Nek5000$ git grep "pragma acc"</div>
<div>jl/gs.c:#pragma acc update
host(sendbuf[0:unit_size*bufSize/2]) if(acc)</div>
<div>jl/gs.c:#pragma acc update
device(buf[0:unit_size*bufSize/2]) if(acc)</div>
<div>jl/gs.c:#pragma acc update
host(sendbuf[0:unit_size*bufSize/2]) if(acc)</div>
<div>jl/gs.c:#pragma acc update
device(buf[0:unit_size*bufSize/2]) if(acc)</div>
<div>jl/gs.c:#pragma acc exit data delete(map0,map1)</div>
<div>jl/gs.c:#pragma acc update
host(buf[0:unit_size*bufSize]) if(acc)</div>
<div>jl/gs.c:#pragma acc update
device(buf[0:unit_size*bufSize]) if(acc)</div>
<div>jl/gs.c:#pragma acc update
host(buf[0:unit_size*bufSize]) if(acc)</div>
<div>jl/gs.c:#pragma acc update
device(buf[0:unit_size*bufSize]) if(acc)</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage[0].scatter_map[0:stage[0].s_size],stage[0].scatter_mapf[0:stage[0].s_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage[i].scatter_map[i:stage[i].s_size],stage[i].scatter_mapf[i:stage[i].s_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage[i].gather_map[i:stage[i].g_size],stage[i].gather_mapf[i:stage[i].g_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage[i].scatter_map[i:stage[i].s_size],stage[i].scatter_mapf[i:stage[i].s_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage[i].gather_map[i:stage[i].g_size],stage[i].gather_mapf[i:stage[i].g_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage2[0].scatter_map[0:stage2[0].s_size],stage2[0].scatter_mapf[0:stage2[0].s_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage2[i].scatter_map[i:stage2[i].s_size],stage2[i].scatter_mapf[i:stage2[i].s_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage2[i].gather_map[i:stage2[i].g_size],stage2[i].gather_mapf[i:stage2[i].g_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage2[i].scatter_map[i:stage2[i].s_size],stage2[i].scatter_mapf[i:stage2[i].s_nt])</div>
<div>jl/gs.c://#pragma acc enter data
copyin(stage2[i].gather_map[i:stage2[i].g_size],stage2[i].gather_mapf[i:stage2[i].g_nt])</div>
<div>jl/gs.c:#pragma acc exit data delete(map,mapf)</div>
<div>jl/gs.c:#pragma acc exit data delete(map,mapf)</div>
<div>jl/gs.c:#pragma acc exit data delete(map,mapf)</div>
<div>jl/gs.c:#pragma acc exit data delete(map,mapf)</div>
<div>jl/gs.c:#pragma acc update
host(buf[0:vn*unit_size*bufSize]) if(acc)</div>
<div>jl/gs.c:#pragma acc update
device(buf[0:vn*unit_size*bufSize]) if(acc)</div>
<div>jl/gs.c: //#pragma acc exit data
delete(ard->map_to_buf[0],ard->map_to_buf[1],ard->map_from_buf[0],ard->map_from_buf[1])</div>
<div>jl/gs.c: //#pragma acc enter data
copyin(ard->map_to_buf[0][0:ard->mt_size[0]],ard->map_from_buf[0][0:ard->mf_size[0]],ard->map_to_buf_f[0][0:ard->mt_nt[0]],ard->map_from_buf_f[0][0:ard->mf_nt[0]],ard->map_to_buf[1][0:ard->mt_size[1]],ard->map_from_buf[1][0:ard->mf_size[1]],ard->map_to_buf_f[1][0:ard->mt_nt[1]],ard->map_from_buf_f[1][0:ard->mf_nt[1]])</div>
<div>jl/gs.c:#pragma acc update host(a[0:n])</div>
<div>jl/gs.c:#pragma acc update host(a[0:n])</div>
<div>jl/gs.c:#pragma acc exit data delete(bufPtr)</div>
<div>jl/gs.c:#pragma acc enter data
create(bufPtr[0:vn*gs_dom_size[dom]*gsh->r.buffer_size])</div>
<div>jl/gs.c:#pragma acc exit data delete(bufPtr)</div>
<div>jl/gs.c:#pragma acc enter data
create(bufPtr[0:vn*gs_dom_size[dom]*gsh->r.buffer_size])</div>
<div>jl/gs.c:#pragma acc exit data
delete(map_local0,map_local1,flagged_primaries)</div>
<div>jl/gs.c:#pragma acc enter data
pcopyin(map[0:*m_size],mapf2[0:2*mf_temp])</div>
<div>jl/gs_acc.c://#pragma acc data present(buf[0:l])</div>
<div>jl/gs_acc.c://#pragma acc host_data use_device(buf)</div>
<div>jl/gs_acc.c://#pragma acc data present(buf[0:l])</div>
<div>jl/gs_acc.c://#pragma acc host_data use_device(buf)</div>
<div>jl/gs_acc.c: //#pragma acc enter data
copyin(t_mapf[0:t_m_nt*2],mapf[0:m_nt*2],snd_mapf[0:snd_m_nt*2],rcv_mapf[0:rcv_m_nt*2],fp_mapf[0:fp_m_nt*2],
t_map[0:t_m_size],map[0:m_size],fp_map[0:fp_m_size],snd_map[0:snd_m_size],rcv_map[0:rcv_m_size])</div>
<div>jl/gs_acc.c: //#pragma acc enter data
copyin(t_mapf[0:t_m_nt*2],mapf[0:m_nt*2],snd_mapf[0:snd_m_nt*2],rcv_mapf[0:rcv_m_nt*2],fp_mapf[0:fp_m_nt*2],
t_map[0:t_m_size],map[0:m_size],fp_map[0:fp_m_size],snd_map[0:snd_m_size],rcv_map[0:rcv_m_size])</div>
<div>jl/gs_acc.c://#pragma acc enter data
pcopyin(t_mapf[0:t_m_nt*2],mapf[0:m_nt*2],snd_mapf[0:snd_m_nt*2],rcv_mapf[0:rcv_m_nt*2],fp_mapf[0:fp_m_nt*2],
t_map[0:t_m_size],map[0:m_size],fp_map[0:fp_m_size],snd_map[0:snd_m_size],rcv_map[0:rcv_m_size])</div>
<div>jl/gs_acc.c://#pragma acc data
present(u[0:uds],mapf[0:m_nt*2],snd_mapf[0:snd_m_nt*2],rcv_mapf[0:rcv_m_nt*2],fp_mapf[0:fp_m_nt*2],t_map[0:t_m_size],map[0:m_size],fp_map[0:fp_m_size],snd_map[0:snd_m_size],rcv_map[0:rcv_m_size]) </div>
<div>jl/gs_acc.c://#pragma acc data
create(sbuf[0:bl],rbuf[0:bl]) if(bl!=0)</div>
<div>jl/gs_acc.c://#pragma acc parallel loop gang vector
present(u[0:uds],map[0:m_size],mapf[0:m_nt*2])
private(i,j,t) async(k+1)</div>
<div>jl/gs_acc.c://#pragma acc loop seq</div>
<div>jl/gs_acc.c://#pragma acc wait </div>
<div>jl/gs_acc.c://#pragma acc parallel loop gang vector
present(u[0:uds],fp_map[0:fp_m_size],fp_mapf[0:fp_m_nt*2])
private(i,j) async(k+1)</div>
<div>jl/gs_acc.c://#pragma acc loop seq</div>
<div>jl/gs_acc.c://#pragma acc wait</div>
<div>jl/gs_acc.c://#pragma acc parallel loop gang vector
present(u[0:uds],fp_map[0:fp_m_size]) private(i,k)</div>
<div>jl/gs_acc.c://#pragma acc parallel loop gang vector
present(u[0:uds],snd_map[0:snd_m_size],snd_mapf[0:snd_m_nt*2],sbuf[0:bl])
private(i,j,t) async(k+1)</div>
<div>jl/gs_acc.c://#pragma acc loop seq</div>
<div>jl/gs_acc.c://#pragma acc wait </div>
<div>jl/gs_acc.c://#pragma acc parallel loop gang vector
present(u[0:uds],snd_map[0:snd_m_size],sbuf[0:bl])
private(i,j,k)</div>
<div>jl/gs_acc.c://#pragma acc update host(sbuf[0:bl])
async(vn+2)</div>
<div>jl/gs_acc.c://#pragma acc wait</div>
<div>jl/gs_acc.c://#pragma acc update device(rbuf[0:bl])
async(vn+2)</div>
<div>jl/gs_acc.c://#pragma acc wait</div>
<div>jl/gs_acc.c://#pragma acc parallel loop gang vector
present(u[0:uds],rcv_map[0:rcv_m_size],rcv_mapf[0:rcv_m_nt*2],rbuf[0:bl])
private(i,j,t) async(k+1)</div>
<div>jl/gs_acc.c://#pragma acc loop seq</div>
<div>jl/gs_acc.c://#pragma acc wait </div>
<div>jl/gs_acc.c: //#pragma acc parallel loop gang vector
present(u[0:uds],rcv_map[0:rcv_m_size],rbuf[0:bl])
private(i,j,k)</div>
<div>jl/gs_acc.c://#pragma acc parallel loop gang vector
present(u[0:uds],t_map[0:t_m_size],t_mapf[0:t_m_nt*2])
private(i,j,t) async(k+1)</div>
<div>jl/gs_acc.c://#pragma acc loop seq</div>
<div>jl/gs_acc.c://#pragma acc wait</div>
<br>
>><br>
>><br>
>> > <a moz-do-not-send="true"
href="http://mrfil.github.io/PowerGrid/">http://mrfil.github.io/PowerGrid/</a><br>
>><br>
>> /tmp/PowerGrid$ git grep -il "pragma acc"<br>
>> PowerGrid/Gfft.hpp<br>
>> PowerGrid/Gnufft.hpp<br>
>> PowerGrid/ftCpu.hpp<br>
>> PowerGrid/gridding.hpp<br>
>> PowerGrid/griddingSupport.hpp<br>
>><br>
>> From <a moz-do-not-send="true"
href="http://mrfil.github.io/PowerGrid/docs/Installation">http://mrfil.github.io/PowerGrid/docs/Installation</a>:<br>
>><br>
>> We have experience with PGC++ 15.7 from NVIDIA/The
Portland Group as the version we have used most extensively.
There is a free license available as part of the OpenACC
Toolkit for academic users.<br>
>><br>
>> GCC 6.1 has OpenACC support but has not yet been
tested by the developers, we welcome reports of anyone
trying to compile with it. We hope to support it alongside
PGI compilers in the near future.<br>
>><br>
>> For those lucky enough to have access to Cray
supercomputers, the Cray compiler does support OpenACC, but
we have not tried to build with it. Because the Cray
compilers are not available on desktops, workstations, or
non-Cray branded clusters, we have not dedicated resources
to testing PowerGrid on it.<br>
>><br>
>> So these folks support OpenACC, but haven't
bothered to try the GCC implementation in the 1+ year that
it's been available. How likely are they to use Clang's?<br>
><br>
><br>
> I cannot answer that. Perhaps they were waiting for GCC
support to mature?<br>
<br>
Or maybe they aren't interested using in OpenACC compiler
support outside of PGI.</div>
<div><br>
</div>
<div>What I'm really getting at here is who is going to use
OpenACC support in Clang, particularly if there is no
compatible Fortran OpenACC compiler? In addition to
justifying the code maintenance effort, users who are not
developers are essential for implementation hardening.</div>
<div><br>
</div>
<div>Best,</div>
<div><br>
</div>
<div>Jeff<br>
<br>
> Thanks.<br>
><br>
> Joel<br>
><br>
>><br>
>> > <a moz-do-not-send="true"
href="https://nek5000.mcs.anl.gov/">https://nek5000.mcs.anl.gov/</a>
(look at the openacc branch in github)<br>
>><br>
>> (on the openacc branch)<br>
>><br>
>> /tmp/Nek5000$ git grep -il "\$acc "<br>
>> core/acc.f<br>
>> core/comm_mpi.f<br>
>> core/gmres.f<br>
>> core/hmholtz.f<br>
>> core/hsmg.f<br>
>> core/math.f<br>
>> core/navier1.f<br>
>> core/navier4.f<br>
>> core/plan4.f<br>
>> core/prepost.f<br>
>> core/subs2.f<br>
>><br>
>> ><br>
>> > <a moz-do-not-send="true"
href="https://nekcem.mcs.anl.gov/">https://nekcem.mcs.anl.gov/</a><br>
>><br>
>> (on master)<br>
>> /tmp/svn$ git grep -il "\$acc"<br>
>> branches/maxwell-experimental/src/cem_dg.F<br>
>> branches/maxwell-experimental/src/dssum2.F<br>
>> branches/maxwell-experimental/src/io.F<br>
>> branches/maxwell-experimental/src/mat1.F<br>
>> branches/maxwell-experimental/src/maxwell.F<br>
>> branches/maxwell-experimental/src/maxwell_acc.F<br>
>> branches/maxwell-experimental/src/mxm_acc.F<br>
>> branches/trunkQu/src/quantum_csr.F<br>
>> branches/trunkQu/src/quantum_setup.f<br>
>> branches/trunkQu/src/quantum_time.F<br>
>> trunk/examples/openacc_gpu=1/box.usr<br>
>> trunk/examples/openacc_gpu=8/box.usr<br>
>> trunk/src/acoustic.F<br>
>> trunk/src/cem_dg2.F<br>
>> trunk/src/complex.F<br>
>> trunk/src/drift1.F<br>
>> trunk/src/drift1_maud.F<br>
>> trunk/src/drive.F<br>
>> trunk/src/drive_maud.F<br>
>> trunk/src/dssum2.F<br>
>> trunk/src/hmholtz.F<br>
>> trunk/src/io.F<br>
>> trunk/src/mat1.F<br>
>> trunk/src/maxwell.F<br>
>> trunk/src/maxwell_acc.F<br>
>> trunk/src/mg_r2204.F<br>
>> trunk/src/mxm_acc.F<br>
>> trunk/src/poisson.F<br>
>> trunk/src/quantum2.F<br>
>> www/examples/libs/phpThumb/phpthumb.functions.php<br>
>> www/examples/phpthumb.functions.php<br>
>><br>
>> >> Can you provide a list of C/C++
applications using OpenACC today and estimate the number of
users that will benefit from this feature?<br>
>> >><br>
>> >><br>
>> >> Such lists exist, although I don't know
what can be shared (and Oak Ridge likely has better lists in
this regard than I do).<br>
>> ><br>
>> > I'll look for a better list that I can share.<br>
>><br>
>> That would be helpful.<br>
>><br>
>> Best,<br>
>><br>
>> Jeff<br>
>><br>
>><br>
>><br>
>> >> I can tell you, from my own experience,
that we're seeing an increase in development using OpenACC,
in both C/C++ and Fortran, over the last couple of years
(essentially because the compiler technology has improved to
the point where that is now a potentially-productive
choice).<br>
>> ><br>
>> ><br>
>> > Providing support in a production-quality,
open-source compiler tool chain like LLVM will hopefully
accelerate this trend.<br>
>> ><br>
>> > Joel<br>
>> ><br>
>> >><br>
>> >> Also, we have a strong desire to enable
tooling over code bases using OpenACC. Among many other
things, at some point we'll likely want the option to
automatically migrate much of this code to using OpenMP.
Having an OpenACC-enabled Clang, with an implementation that
maps to OpenMP, is an important step in that process.<br>
>> >><br>
>> >> -Hal<br>
>> >><br>
>> >><br>
>> >><br>
>> >> Thanks,<br>
>> >><br>
>> >> Jeff<br>
>> >><br>
>> >> On Tue, Dec 5, 2017 at 11:06 AM, Joel E.
Denny via cfe-dev <<a moz-do-not-send="true"
href="mailto:cfe-dev@lists.llvm.org">cfe-dev@lists.llvm.org</a>>
wrote:<br>
>> >>><br>
>> >>> Hi,<br>
>> >>><br>
>> >>> We are working on a new project,
clacc, that extends clang with OpenACC support. Clacc's
approach is to translate OpenACC (a descriptive language) to
OpenMP (a prescriptive language) and thus to build on
clang's existing OpenMP support. While we plan to develop
clacc to support our own research, an important goal is to
contribute clacc as a production-quality component of
upstream clang.<br>
>> >>><br>
>> >>> We have begun implementing an early
prototype of clacc. Before we get too far into the
implementation, we would like to get feedback from the LLVM
community to help ensure our design would ultimately be
acceptable for contribution. For that purpose, below is an
analysis of several high-level design alternatives we have
considered and their various features. We welcome any
feedback.<br>
>> >>><br>
>> >>> Thanks.<br>
>> >>><br>
>> >>> Joel E. Denny<br>
>> >>> Future Technologies Group<br>
>> >>> Oak Ridge National Laboratory<br>
>> >>><br>
>> >>><br>
>> >>> Design Alternatives<br>
>> >>> -------------------<br>
>> >>><br>
>> >>> We have considered three design
alternatives for the clacc compiler:<br>
>> >>><br>
>> >>> 1. acc src --parser-->
omp AST --codegen--> LLVM IR + omp rt calls<br>
>> >>> 2. acc src --parser--> acc AST
--codegen--> LLVM IR + omp rt calls<br>
>> >>> 3. acc src --parser--> acc AST
--ttx--> omp AST --codegen--> LLVM IR + omp rt
calls<br>
>> >>><br>
>> >>> In the above diagram:<br>
>> >>><br>
>> >>> * acc src = C source code containing
acc constructs.<br>
>> >>> * acc AST = a clang AST in which acc
constructs are represented by<br>
>> >>> nodes with acc node types. Of
course, such node types do not<br>
>> >>> already exist in clang's
implementation.<br>
>> >>> * omp AST = a clang AST in which acc
constructs have been lowered<br>
>> >>> to omp constructs represented by
nodes with omp node types. Of<br>
>> >>> course, such node types do already
exist in clang's<br>
>> >>> implementation.<br>
>> >>> * parser = the existing clang parser
and semantic analyzer,<br>
>> >>> extended to handle acc constructs.<br>
>> >>> * codegen = the existing clang backend
that translates a clang AST<br>
>> >>> to LLVM IR, extended if necessary
(depending on which design is<br>
>> >>> chosen) to perform codegen from acc
nodes.<br>
>> >>> * ttx (tree transformer) = a new clang
component that transforms<br>
>> >>> acc to omp in clang ASTs.<br>
>> >>><br>
>> >>> Design Features<br>
>> >>> ---------------<br>
>> >>><br>
>> >>> There are several features to consider
when choosing among the designs<br>
>> >>> in the previous section:<br>
>> >>><br>
>> >>> 1. acc AST as an artifact -- Because
they create acc AST nodes,<br>
>> >>> designs 2 and 3 best facilitate the
creation of additional acc<br>
>> >>> source-level tools (such as pretty
printers, analyzers, lint-like<br>
>> >>> tools, and editor extensions).
Some of these tools, such as pretty<br>
>> >>> printing, would be available
immediately or as minor extensions of<br>
>> >>> tools that already exist in clang's
ecosystem.<br>
>> >>><br>
>> >>> 2. omp AST/source as an artifact --
Because they create omp AST<br>
>> >>> nodes, designs 1 and 3 best
facilitate the use of source-level<br>
>> >>> tools to help an application
developer discover how clacc has<br>
>> >>> mapped his acc to omp, possibly in
order to debug a mapping<br>
>> >>> specification he has supplied.
With design 2 instead, an<br>
>> >>> application developer has to
examine low-level LLVM IR + omp rt<br>
>> >>> calls. Moreover, with designs 1
and 3, permanently migrating an<br>
>> >>> application's acc source to omp
source can be automated.<br>
>> >>><br>
>> >>> 3. omp AST for mapping implementation
-- Designs 1 and 3 might<br>
>> >>> also make it easier for the
compiler developer to reason about and<br>
>> >>> implement mappings from acc to
omp. That is, because acc and omp<br>
>> >>> syntax is so similar, implementing
the translation at the level of<br>
>> >>> a syntactic representation is
probably easier than translating to<br>
>> >>> LLVM IR.<br>
>> >>><br>
>> >>> 4. omp AST for codegen -- Designs 1
and 3 simplify the<br>
>> >>> compiler implementation by enabling
reuse of clang's existing omp<br>
>> >>> support for codegen. In contrast,
design 2 requires at least some<br>
>> >>> extensions to clang codegen to
support acc nodes.<br>
>> >>><br>
>> >>> 5. Full acc AST for mapping -- Designs
2 and 3 potentially<br>
>> >>> enable the compiler to analyze the
entire source (as opposed to<br>
>> >>> just the acc construct currently
being parsed) while choosing the<br>
>> >>> mapping to omp. It is not clear if
this feature will prove useful,<br>
>> >>> but it might enable more
optimizations and compiler research<br>
>> >>> opportunities.<br>
>> >>><br>
>> >>> 6. No acc node classes -- Design 1
simplifies the compiler<br>
>> >>> implementation by eliminating the
need to implement many acc node<br>
>> >>> classes. While we have so far
found that implementing these<br>
>> >>> classes is mostly mechanical, it
does take a non-trivial amount of<br>
>> >>> time.<br>
>> >>><br>
>> >>> 7. No omp mapping -- Design 2 does not
require acc to be mapped to<br>
>> >>> omp. That is, it is conceivable
that, for some acc constructs,<br>
>> >>> there will prove to be no omp
syntax to capture the semantics we<br>
>> >>> wish to implement. It is also
conceivable that we might one day<br>
>> >>> want to represent some acc
constructs directly as extensions to<br>
>> >>> LLVM IR, where some acc analyses or
optimizations might be more<br>
>> >>> feasible to implement. This
possibility dovetails with recent<br>
>> >>> discussions in the LLVM community
about developing LLVM IR<br>
>> >>> extensions for various parallel
programming models.<br>
>> >>><br>
>> >>> Because of features 4 and 6, design 1
is likely the fastest design to<br>
>> >>> implement, at least at first while we
focus on simple acc features and<br>
>> >>> simple mappings to omp. However, we
have so far found no advantage<br>
>> >>> that design 1 has but that design 3
does not have except for feature<br>
>> >>> 6, which we see as the least important
of the above features in the<br>
>> >>> long term.<br>
>> >>><br>
>> >>> The only advantage we have found that
design 2 has but that design 3<br>
>> >>> does not have is feature 7. It should
be possible to choose design 3<br>
>> >>> as the default but, for certain acc
constructs or scenarios where<br>
>> >>> feature 7 proves important (if any),
incorporate design 2. In other<br>
>> >>> words, if we decide not to map a
particular acc construct to any omp<br>
>> >>> construct, ttx would leave it alone,
and we would extend codegen to<br>
>> >>> handle it directly.<br>
>> >>><br>
>> >>> Conclusions<br>
>> >>> -----------<br>
>> >>><br>
>> >>> For the above reasons, and because
design 3 offers the cleanest<br>
>> >>> separation of concerns, we have chosen
design 3 with the possibility<br>
>> >>> of incorporating design 2 where it
proves useful.<br>
>> >>><br>
>> >>> Because of the immutability of clang's
AST, the design of our proposed<br>
>> >>> ttx component requires careful
consideration. To shorten this initial<br>
>> >>> email, we have omitted those details
for now, but we will be happy to<br>
>> >>> include them as the discussion
progresses.<br>
>> >>><br>
>> >>>
_______________________________________________<br>
>> >>> cfe-dev mailing list<br>
>> >>> <a moz-do-not-send="true"
href="mailto:cfe-dev@lists.llvm.org">cfe-dev@lists.llvm.org</a><br>
>> >>> <a moz-do-not-send="true"
href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev">http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev</a><br>
>> >>><br>
>> >><br>
>> >><br>
>> >><br>
>> >> --<br>
>> >> Jeff Hammond<br>
>> >> <a moz-do-not-send="true"
href="mailto:jeff.science@gmail.com">jeff.science@gmail.com</a><br>
>> >> <a moz-do-not-send="true"
href="http://jeffhammond.github.io/">http://jeffhammond.github.io/</a><br>
>> >><br>
>> >><br>
>> >>
_______________________________________________<br>
>> >> cfe-dev mailing list<br>
>> >> <a moz-do-not-send="true"
href="mailto:cfe-dev@lists.llvm.org">cfe-dev@lists.llvm.org</a><br>
>> >> <a moz-do-not-send="true"
href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev">http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev</a><br>
>> >><br>
>> >><br>
>> >> --<br>
>> >> Hal Finkel<br>
>> >> Lead, Compiler Technology and Programming
Languages<br>
>> >> Leadership Computing Facility<br>
>> >> Argonne National Laboratory<br>
>> ><br>
>> ><br>
>><br>
>><br>
>><br>
>> --<br>
>> Jeff Hammond<br>
>> <a moz-do-not-send="true"
href="mailto:jeff.science@gmail.com">jeff.science@gmail.com</a><br>
>> <a moz-do-not-send="true"
href="http://jeffhammond.github.io/">http://jeffhammond.github.io/</a><br>
><br>
><br>
<br>
<br>
<br>
--<br>
Jeff Hammond<br>
<a moz-do-not-send="true"
href="mailto:jeff.science@gmail.com">jeff.science@gmail.com</a><br>
<a moz-do-not-send="true"
href="http://jeffhammond.github.io/">http://jeffhammond.github.io/</a></div>
</div>
</div>
</blockquote>
<br>
<pre class="moz-signature" cols="72">--
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory</pre>
</body>
</html>