<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=us-ascii">
<meta name="Generator" content="Microsoft Word 15 (filtered medium)">
<style><!--
/* Font Definitions */
@font-face
{font-family:"Cambria Math";
panose-1:2 4 5 3 5 4 6 3 2 4;}
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0in;
margin-bottom:.0001pt;
font-size:11.0pt;
font-family:"Calibri",sans-serif;}
span.EmailStyle20
{mso-style-type:personal-compose;
font-family:"Calibri",sans-serif;
color:windowtext;}
.MsoChpDefault
{mso-style-type:export-only;
font-size:10.0pt;
font-family:"Calibri",sans-serif;}
@page WordSection1
{size:8.5in 11.0in;
margin:1.0in 1.0in 1.0in 1.0in;}
div.WordSection1
{page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
</head>
<body lang="EN-US" link="#0563C1" vlink="#954F72">
<div class="WordSection1">
<p class="MsoNormal">Hi,<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">I have a general question on the behavior of `omp_get_num_devices()`, which seems to be interpreted differently across different compiler vendors (I tried Clang/LLVM and IBM XL).<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Consider this simple example:<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><span style="font-size:9.0pt;font-family:"Courier New"">#include <omp.h><o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:9.0pt;font-family:"Courier New""><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:9.0pt;font-family:"Courier New"">int main() {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:9.0pt;font-family:"Courier New""> printf("num devices: %ld\n", omp_get_num_devices());<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:9.0pt;font-family:"Courier New"">}<o:p></o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">With clang, when this program is compiled for the host (-fopenmp) “omp_get_num_devices() == 0”, even if there are GPUs in the system. If it is compiled with ` -fopenmp-targets=nvptx64` it returns a value >0 (if there are GPUs in the system
of course).<o:p></o:p></p>
<p class="MsoNormal">With IBM XL, no matter if you compile with offload support or not, it seems to always return a value >0 if there are GPUs in the system.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">It looks like in the OpenMP spec is not clear what’s the right behavior in this case. I can see that both clang and IBM XL approach are ok (especially if this is implementation dependent). However, this could cause ambiguity in some tests,
that expect either behaviors.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Consider this sollve test: `https://raw.githubusercontent.com/SOLLVE/sollve_vv/master/tests/4.5/target_teams_distribute_parallel_for/test_target_teams_distribute_parallel_for_devices.c`:<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> int num_dev = omp_get_num_devices();<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> ...<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> ...<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black">
<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> for (dev = 0; dev < num_dev; ++dev) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black">#pragma omp target enter data map(to: a[0:SIZE_N]) device(dev)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> }<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> for (dev = 0; dev < num_dev; ++dev) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> #pragma omp target teams distribute parallel for device(dev) map(tofrom: isHost)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> for (i = 0; i < SIZE_N; i++) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> if (omp_get_team_num() == 0 && omp_get_thread_num() == 0) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> isHost[dev] = omp_is_initial_device();// Checking if running on a device<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> }<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> a[i] += dev;<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> }<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> }<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"><o:p> </o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> for (dev = 0; dev < num_dev; ++dev) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black">#pragma omp target exit data map(from: a[0:SIZE_N]) device(dev)<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> OMPVV_INFOMSG("Device %d ran on the %s", dev, isHost[dev] ? "host" : "device");<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> OMPVV_TEST_AND_SET(errors, isHost[dev] && dev != omp_get_initial_device());<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> for (i = 0; i < SIZE_N; i++) {<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> OMPVV_TEST_AND_SET(errors, a[i] != 1 + dev);<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> }<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:10.0pt;font-family:"Courier New";color:black"> }<o:p></o:p></span></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">This test fails with IBM XL if compiled for host only because it returns a number of devices > 0, but passes with clang because it returns a number of devices == 0 and it basically skips both for loops.<o:p></o:p></p>
<p class="MsoNormal">Of course the test can be changed, but I have seen this pattern quite often in different benchmarks and this difference will make it harder to write a test that is compatible on host and device (I know that OpenMP spec doesn’t guarantee
this) and with different compilers.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">I think it is always nice to have consistent behaviors across compilers as much as possible, but maybe I am missing something, so I wanted to hear your thoughts.<o:p></o:p></p>
<p class="MsoNormal"><o:p> </o:p></p>
<p class="MsoNormal">Thanks.<o:p></o:p></p>
<p class="MsoNormal">Simone<o:p></o:p></p>
</div>
</body>
</html>