<html>
<head>
<base href="https://bugs.llvm.org/">
</head>
<body><table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Bug ID</th>
<td><a class="bz_bug_link
bz_status_NEW "
title="NEW - Clang Incorrectly Determines SPMD vs Generic OpenMP Kernel"
href="https://bugs.llvm.org/show_bug.cgi?id=48851">48851</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>Clang Incorrectly Determines SPMD vs Generic OpenMP Kernel
</td>
</tr>
<tr>
<th>Product</th>
<td>clang
</td>
</tr>
<tr>
<th>Version</th>
<td>trunk
</td>
</tr>
<tr>
<th>Hardware</th>
<td>PC
</td>
</tr>
<tr>
<th>OS</th>
<td>Linux
</td>
</tr>
<tr>
<th>Status</th>
<td>NEW
</td>
</tr>
<tr>
<th>Severity</th>
<td>normal
</td>
</tr>
<tr>
<th>Priority</th>
<td>P
</td>
</tr>
<tr>
<th>Component</th>
<td>OpenMP
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedclangbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>estewart08@gmail.com
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr></table>
<p>
<div>
<pre>Created <span class=""><a href="attachment.cgi?id=24410" name="attach_24410" title="SPMD vs Generic Bug Reproducer">attachment 24410</a> <a href="attachment.cgi?id=24410&action=edit" title="SPMD vs Generic Bug Reproducer">[details]</a></span>
SPMD vs Generic Bug Reproducer
According to the docs, <a href="https://clang.llvm.org/docs/OpenMPSupport.html">https://clang.llvm.org/docs/OpenMPSupport.html</a>:
The compiler will always attempt to use the SPMD mode wherever possible. SPMD
mode will not be used if:
The target region contains user code (other than OpenMP-specific
directives) in between the target and the parallel directives.
#pragma omp target teams thread_limit(THREADS) num_teams(TEAMS)
map(from:gpu_results)
{
int dist[THREADS];
// Uncomment line below to trigger generic kernel
//dist[0] = 0;
#pragma omp parallel
{
Clang determines this is a SPMD kernel, even though there is user code between
the target and parallel pragmas, and the dist array becomes thread private. In
the attached code, there is a reduction later on that assumes dist is team
private.
When looking for the child of the target teams pragma, the compiler skips the
declaration of the dist array and sees the parallel pragma for SPMD mode. This
occurs due to the logic in clang/lib/CodeGen/CGOpenMPRuntime.cpp on line 6525,
VD->getType().isTrivialType(Ctx). Since this declaration is of trivial type it
is skipped, which has adverse effects.
To summarize:
SPMD (dist array is thread private)
int dist[THREADS]; //skipped
#pragma omp parallel //child that is returned
Generic (dist array is team private)
int dist[THREADS]; //skipped
dist[0] = 0 // child 1
#pragma omp parallel //child 2, multiple children, return nullptr
Attachment is derived from:
<a href="https://github.com/zjin-lcf/oneAPI-DirectProgramming/blob/master/all-pairs-distance-omp/main.cpp">https://github.com/zjin-lcf/oneAPI-DirectProgramming/blob/master/all-pairs-distance-omp/main.cpp</a></pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are on the CC list for the bug.</li>
</ul>
</body>
</html>