<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body text="#000000" bgcolor="#FFFFFF">
<p>Thanks for all the help!</p>
<p>These benchmarks, especially lulesh, showed me that my approach
of setting the address spaces manually according to their scope
doesn't work(*) and that I have to use the generic address space<br>
like the nvptx backend does.<br>
</p>
<p>Now with this much more robust version, I decided to make my
project public:</p>
<p>My clang fork is available at<br>
<a class="moz-txt-link-freetext" href="https://github.com/daniel-schuermann/clang">https://github.com/daniel-schuermann/clang</a></p>
<p>and the openmp runtime fork at<br>
<a class="moz-txt-link-freetext" href="https://github.com/daniel-schuermann/openmp">https://github.com/daniel-schuermann/openmp</a></p>
<p>and the necessary llvm fork (for generating SPIR-V) is from<br>
<a class="moz-txt-link-freetext" href="https://github.com/thewilsonator/llvm/tree/compute">https://github.com/thewilsonator/llvm/tree/compute</a></p>
<p>All 3 of them might need a pull from upstream as they are not
always synced.</p>
<p>libomptarget-spir needs an OpenCL runtime which supports SPIRV
kernels (with OpenCL 2.1 headers).<br>
Unfortunately, the Intel OpenCL runtime started to segfault with
the change to generic address space.<br>
Therefore, the only working OpenCL runtime I know of is
AMDGPU-Pro.</p>
<p>The follwing pragmas should work for now:<br>
#target (enter/exit data)<br>
#teams<br>
#distribute / parallel for<br>
#master<br>
#barrier</p>
<p>as well as the clauses:<br>
map, shared, private, firstprivate, lastprivate, schedule</p>
<p>A device runtime is not planned for the moment (I also don't
know, how to compile and link OpenCL sources into libomptarget),
but the generated code works independently from runtime functions.<br>
</p>
<p>For those more adventurous, you can (try to) build your sources
with -fopenmp -fopenmp-targets=spir64-unknown-unknown<br>
For C++, I recommend to add -fno-exceptions and -O0 as there seem
to be optimizer passes enabled<br>
which don't work for spir.</p>
<p>I would be very thankful for some feedback (I hope, it won't get
too depressing).<br>
Although I am not able to accept pull requests at the moment for
legal purposes,<br>
I would welcome any hint to make the implementation more robust
and complete<br>
as well as statements about (not) working programs.<br>
(no comments on code style please, this can be fixed later °°)</p>
<p>Kind regards,<br>
Daniel</p>
<p>(*) While this is a valid program snipped,</p>
<p>#pragma omp target map(to:a[0:n])<br>
{ int * b = a; }<br>
</p>
<p>it doesn't work if |a| is a pointer to addrSpace(1) (cl_global)
and |b| gets allocated as pointer to cl_private.<br>
</p>
<br>
<div class="moz-cite-prefix">On 10/11/2017 06:41 PM, Sunita
Chandrasekaran wrote:<br>
</div>
<blockquote type="cite"
cite="mid:CAJOqkcrDOJ8SLX35duGLtgwcGAieyfwm41Ok4UyYi6-+AmA-pA@mail.gmail.com">
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
<div dir="ltr">
<div class="gmail_default" style="font-family:georgia,serif">Yes
these SNU OMP 3.1 C codes are quite useful (and some of the
authors of this software also attend SC and have a booth on
the showfloor. I met one or 2 of the last year at SC). <br>
</div>
<div class="gmail_default" style="font-family:georgia,serif"><br>
</div>
<div class="gmail_default" style="font-family:georgia,serif">On
a related note, some of us (w/ NASA) in my group have just
begun to create OpenMP 4.5 NPB codes (that did not make it to
the SPEC HPG). We will keep you informed. <br>
</div>
<div class="gmail_default" style="font-family:georgia,serif"><br>
</div>
<div class="gmail_default" style="font-family:georgia,serif">Thanks
<br>
</div>
<div class="gmail_default" style="font-family:georgia,serif">Sunita</div>
<div class="gmail_default" style="font-family:georgia,serif"><br>
</div>
<div class="gmail_default" style="font-family:georgia,serif">
<div dir="ltr"
style="color:rgb(34,34,34);font-family:arial,sans-serif;font-size:small;font-style:normal;font-variant-ligatures:normal;font-variant-caps:normal;font-weight:normal;letter-spacing:normal;text-align:start;text-indent:0px;text-transform:none;white-space:normal;word-spacing:0px;background-color:rgb(255,255,255);text-decoration-style:initial;text-decoration-color:initial">
<div style="font-family:georgia,serif"><br
class="m_4929773007851618185gmail-Apple-interchange-newline">
*****************************<wbr>*****</div>
Sunita Chandrasekaran<br>
Asst. Prof. Computer and Information Sciences<br>
Affiliated, Center for Bioinformatics and Computational
Biology<br>
430 Smith Hall, University of Delaware<br>
p: <a href="tel:%28302%29%20831-2714" value="+13028312714"
target="_blank" moz-do-not-send="true">302-831-2714</a>
e: <a href="mailto:schandra@udel.edu"
style="color:rgb(17,85,204)" target="_blank"
moz-do-not-send="true">schandra@udel.edu</a><br>
<div style="font-family:georgia,serif"><span
style="font-family:arial,sans-serif">-----------------------------<wbr>-----------</span></div>
<div style="font-family:georgia,serif"><span
style="font-family:arial,sans-serif">Adjunct Prof. Dept.
of Computer Science</span><br>
</div>
University of Houston, TX<br>
------------------------------<wbr>----------<br>
</div>
<div dir="ltr"
style="color:rgb(34,34,34);font-family:arial,sans-serif;font-size:small;font-style:normal;font-variant-ligatures:normal;font-variant-caps:normal;font-weight:normal;letter-spacing:normal;text-align:start;text-indent:0px;text-transform:none;white-space:normal;word-spacing:0px;background-color:rgb(255,255,255);text-decoration-style:initial;text-decoration-color:initial">t: <a
href="https://twitter.com/chandrasunita"
style="color:rgb(17,85,204)" target="_blank"
moz-do-not-send="true">https://twitter.com/<wbr>chandrasunita</a><br>
w: <a href="https://www.eecis.udel.edu/%7Eschandra/"
style="color:rgb(17,85,204)" target="_blank"
moz-do-not-send="true">https://www.eecis.udel.edu/<wbr>~schandra/</a><br>
CRPL
<div style="font-family:georgia,serif;display:inline">Research
Group: </div>
<a href="http://crpl.cis.udel.edu/"
style="color:rgb(17,85,204)" target="_blank"
moz-do-not-send="true">http://crpl.cis.udel.edu/</a></div>
<br
class="m_4929773007851618185gmail-Apple-interchange-newline">
</div>
<div class="gmail_extra"><br>
<div class="gmail_quote">On Wed, Oct 11, 2017 at 12:24 PM,
Jeff Hammond via Openmp-dev <span dir="ltr"><<a
href="mailto:openmp-dev@lists.llvm.org" target="_blank"
moz-do-not-send="true">openmp-dev@lists.llvm.org</a>></span>
wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0
.8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="ltr">I just found <a
href="http://aces.snu.ac.kr/software/snu-npb/"
target="_blank" moz-do-not-send="true">http://aces.snu.ac.kr/so<wbr>ftware/snu-npb/</a>
today, but have not tried it.
<div><br>
</div>
<div>Jeff</div>
</div>
<div class="gmail_extra">
<div>
<div class="m_4929773007851618185h5"><br>
<div class="gmail_quote">On Tue, Oct 10, 2017 at
7:37 AM, Jeff Hammond <span dir="ltr"><<a
href="mailto:jeff.science@gmail.com"
target="_blank" moz-do-not-send="true">jeff.science@gmail.com</a>></span>
wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0
0 .8ex;border-left:1px #ccc
solid;padding-left:1ex">
<div dir="auto"><a
href="https://github.com/ParRes/Kernels/tree/master/Cxx11"
target="_blank" moz-do-not-send="true">https://github.com/ParRes/Kern<wbr>els/tree/master/Cxx11</a> implemen<wbr>ts
at least two kernels in OpenCL, GPU-oriented
OpenMP4 (#pragma omp <span
style="background-color:rgba(255,255,255,0)">teams
distribute parallel <span
class="m_4929773007851618185m_-5316131450052861093m_-771006556969918392pl-k"
style="box-sizing:border-box">for</span>
simd <span
class="m_4929773007851618185m_-5316131450052861093m_-771006556969918392pl-smi"
style="box-sizing:border-box">collapse</span>(<span
class="m_4929773007851618185m_-5316131450052861093m_-771006556969918392pl-c1"
style="box-sizing:border-box">2</span>) <span
class="m_4929773007851618185m_-5316131450052861093m_-771006556969918392pl-c1"
style="box-sizing:border-box">schedule</span>(<span
class="m_4929773007851618185m_-5316131450052861093m_-771006556969918392pl-k"
style="box-sizing:border-box">static</span>,<span
class="m_4929773007851618185m_-5316131450052861093m_-771006556969918392pl-c1"
style="box-sizing:border-box">1</span>) as
recommended by NVIDIA), RAJA, Kokkos, and
numerous CPU implementations.</span>
<div><span
style="background-color:rgba(255,255,255,0)"><br>
</span></div>
<div><span
style="background-color:rgba(255,255,255,0)">You’ll
need to (1) tune the OpenCL code for your
GPU as it is currently not optimized for
any architecture and (2) tweak the RAJA
and KOKKOS implementations to target GPU
models in addition to CPU ones. Neither
should be very difficult and the PRK team
is happy to provide assistance.</span></div>
<div><span
style="background-color:rgba(255,255,255,0)"><br>
</span></div>
<div>That project contains one C++ CUDA and
two Fortran OpenACC implementations but
these are unfinished and need work. However,
I’m very happy to see others contribute
better versions. It may not be the easiest
route to fame and fortune, but it’s not the
worst either 😂</div>
<div>
<div>
<div>
<div><br>
</div>
<div>Write me privately or create GitHub
issues if you have questions. I am
currently on leave from my day job but
will be responsive via GMail and
GitHub.</div>
<div><br>
</div>
<div>Sorry if this is a repost. I
intended to reply early but didn’t and
can’t find it if I did, but mistakes
have been known to occur.</div>
<div><br>
</div>
<div>Jeff<br>
<br>
<div
id="m_4929773007851618185m_-5316131450052861093m_-771006556969918392AppleMailSignature">Sent
from my iPhone</div>
<div>
<div
class="m_4929773007851618185m_-5316131450052861093h5">
<div><br>
On Oct 10, 2017, at 6:56 AM,
Jonas Hahnfeld via Openmp-dev
<<a
href="mailto:openmp-dev@lists.llvm.org"
target="_blank"
moz-do-not-send="true">openmp-dev@lists.llvm.org</a>>
wrote:<br>
<br>
</div>
<blockquote type="cite">
<div><span>Hi,</span><br>
<span></span><br>
<span>some (public) benchmarks
that come to my mind:</span><br>
<span> * LULESH: <a
href="https://codesign.llnl.gov/lulesh.php"
target="_blank"
moz-do-not-send="true">https://codesign.llnl.gov/lule<wbr>sh.php</a></span><br>
<span> * CloverLeaf: <a
href="http://uk-mac.github.io/CloverLeaf/"
target="_blank"
moz-do-not-send="true">http://uk-mac.github.io/Clover<wbr>Leaf/</a></span><br>
<span> * my own, a Conjugate
Gradient solver: <a
href="https://github.com/hahnjo/CGxx"
target="_blank"
moz-do-not-send="true">https://github.com/hahnjo/CGxx</a></span><br>
<span></span><br>
<span>As James wrote, SPEC
ACCEL also has an OpenMP
suite, but you need a
license.</span><br>
<span></span><br>
<span>Regards</span><br>
<span>Jonas</span><br>
<span></span><br>
<span>Am 2017-10-10 09:43,
schrieb Daniel Schürmann via
Openmp-dev:</span><br>
<blockquote type="cite"><span>Hello
together,</span><br>
</blockquote>
<blockquote type="cite"><span>as
the title states, I am
looking for benchmarks
which are</span><br>
</blockquote>
<blockquote type="cite"><span>particularly
suitable for GPU
accelerators (or at least
make use of</span><br>
</blockquote>
<blockquote type="cite"><span>the
#teams pragma).</span><br>
</blockquote>
<blockquote type="cite"><span>I
already tried rodinia
benchmark suite, but they
seem to be written</span><br>
</blockquote>
<blockquote type="cite"><span>for
CPU acceleration only.</span><br>
</blockquote>
<blockquote type="cite"><span>I
would be very pleased if
someone could provide me
with one or more</span><br>
</blockquote>
<blockquote type="cite"><span>kernels
which can be used with the
NVPTX backend.</span><br>
</blockquote>
<blockquote type="cite"><span>The
purpose is a master thesis
about using OpenCL and
SPIR-V as OpenMP backend.</span><br>
</blockquote>
<blockquote type="cite"><span>Thank
you in advance and</span><br>
</blockquote>
<blockquote type="cite"><span>kind
regards,</span><br>
</blockquote>
<blockquote type="cite"><span>Daniel</span><br>
</blockquote>
<blockquote type="cite"><span>______________________________<wbr>_________________</span><br>
</blockquote>
<blockquote type="cite"><span>Openmp-dev
mailing list</span><br>
</blockquote>
<blockquote type="cite"><span><a
href="mailto:Openmp-dev@lists.llvm.org" target="_blank"
moz-do-not-send="true">Openmp-dev@lists.llvm.org</a></span><br>
</blockquote>
<blockquote type="cite"><span><a
href="http://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev"
target="_blank"
moz-do-not-send="true">http://lists.llvm.org/cgi-bin/<wbr>mailman/listinfo/openmp-dev</a></span><br>
</blockquote>
<span>______________________________<wbr>_________________</span><br>
<span>Openmp-dev mailing list</span><br>
<span><a
href="mailto:Openmp-dev@lists.llvm.org"
target="_blank"
moz-do-not-send="true">Openmp-dev@lists.llvm.org</a></span><br>
<span><a
href="http://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev"
target="_blank"
moz-do-not-send="true">http://lists.llvm.org/cgi-bin/<wbr>mailman/listinfo/openmp-dev</a></span><br>
</div>
</blockquote>
</div>
</div>
</div>
</div>
</div>
</div>
</div>
</blockquote>
</div>
<br>
<br clear="all">
<div><br>
</div>
</div>
</div>
<span class="m_4929773007851618185HOEnZb"><font
color="#888888">-- <br>
<div
class="m_4929773007851618185m_-5316131450052861093gmail_signature"
data-smartmail="gmail_signature">Jeff Hammond<br>
<a href="mailto:jeff.science@gmail.com"
target="_blank" moz-do-not-send="true">jeff.science@gmail.com</a><br>
<a href="http://jeffhammond.github.io/"
target="_blank" moz-do-not-send="true">http://jeffhammond.github.io/</a></div>
</font></span></div>
<br>
______________________________<wbr>_________________<br>
Openmp-dev mailing list<br>
<a href="mailto:Openmp-dev@lists.llvm.org" target="_blank"
moz-do-not-send="true">Openmp-dev@lists.llvm.org</a><br>
<a
href="http://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev"
rel="noreferrer" target="_blank" moz-do-not-send="true">http://lists.llvm.org/cgi-bin/<wbr>mailman/listinfo/openmp-dev</a><br>
<br>
</blockquote>
</div>
<br>
</div>
</div>
</blockquote>
<br>
</body>
</html>