<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 - Possible issue with target enter/exit data calls on objects with pointers inside"
href="https://bugs.llvm.org/show_bug.cgi?id=48218">48218</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>Possible issue with target enter/exit data calls on objects with pointers inside
</td>
</tr>
<tr>
<th>Product</th>
<td>OpenMP
</td>
</tr>
<tr>
<th>Version</th>
<td>unspecified
</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>enhancement
</td>
</tr>
<tr>
<th>Priority</th>
<td>P
</td>
</tr>
<tr>
<th>Component</th>
<td>Clang Compiler Support
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedclangbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>bertoni@anl.gov
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr></table>
<p>
<div>
<pre>With the code below, we see incorrect results (junk data). However, if we
modify it such that the `target exit data` calls are in the inverted order from
the `target enter data` calls, then the code works (see example below).
>From a conversation with Johannes and Kelvin at SC, it seems that this is
expected, but we wanted to file a report to give a concrete example.
Acknowledgements to John Tramm (ANL) for the initial version of this code.
```
<span class="quote">> cat main.cpp </span >
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
typedef struct{
int val;
int * arr;
} MyObject;
int main(void)
{
int N = 5;
// Allocate and initialize an array of MyObjects
MyObject * objects = (MyObject *) malloc(N * sizeof(MyObject));
for( int i = 0; i < N; i++ )
{
objects[i].val = i;
objects[i].arr = (int *) malloc(N * sizeof(int));
for( int j = 0; j < N; j++ )
{
objects[i].arr[j] = i;
}
}
// Send data to device
#pragma omp target enter data map(to: objects[:N])
for( int i = 0; i < N; i++ )
{
#pragma omp target enter data map(to: objects[i].arr[:N])
}
// Execute device kernel
#pragma omp target teams distribute parallel for
for( int i = 0; i < N; i++)
{
objects[i].val *= 2;
for( int j = 0; j < N; j++ )
{
objects[i].arr[j] *= 2;
}
}
// Pull data from device
#ifndef INVERT
#pragma omp target exit data map(from: objects[:N])
for( int i = 0; i < N; i++ )
{
#pragma omp target exit data map(from: objects[i].arr[:N])
}
#else
for( int i = 0; i < N; i++ )
{
#pragma omp target exit data map(from: objects[i].arr[:N])
}
#pragma omp target exit data map(from: objects[:N])
#endif
// Check for any errors
for( int i = 0; i < N; i++)
{
printf( "Expecting: %d, Got: %d\n", i*2, objects[i].val );
assert( objects[i].val == i*2 );
}
for( int i = 0; i < N; i++)
{
for( int j = 0; j < N; j++)
{
printf( "Expecting: %d, Got: %d\n", i*2, objects[i].arr[j] );
assert( objects[i].arr[j] == i*2 );
}
}
return 0;
}
<span class="quote">> clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp</span >
<span class="quote">> ./a.out </span >
Expecting: 0, Got: 0
Expecting: 2, Got: 1
a.out: main.cpp:63: int main(): Assertion `objects[i].val == i*2' failed.
Aborted
<span class="quote">> clang++ -DINVERT -fopenmp -fopenmp-targets=nvptx64 main.cpp</span >
<span class="quote">> ./a.out </span >
Expecting: 0, Got: 0
Expecting: 2, Got: 2
Expecting: 4, Got: 4
Expecting: 6, Got: 6
Expecting: 8, Got: 8
Expecting: 0, Got: 0
Expecting: 0, Got: 0
Expecting: 0, Got: 0
Expecting: 0, Got: 0
Expecting: 0, Got: 0
Expecting: 2, Got: 2
Expecting: 2, Got: 2
Expecting: 2, Got: 2
Expecting: 2, Got: 2
Expecting: 2, Got: 2
Expecting: 4, Got: 4
Expecting: 4, Got: 4
Expecting: 4, Got: 4
Expecting: 4, Got: 4
Expecting: 4, Got: 4
Expecting: 6, Got: 6
Expecting: 6, Got: 6
Expecting: 6, Got: 6
Expecting: 6, Got: 6
Expecting: 6, Got: 6
Expecting: 8, Got: 8
Expecting: 8, Got: 8
Expecting: 8, Got: 8
Expecting: 8, Got: 8
Expecting: 8, Got: 8
<span class="quote">> clang++ --version</span >
clang version 12.0.0
(/gpfs/jlse-fs0/users/yeluo/opt/llvm-clang/llvm-project/clang
15a68fed111f67293f56e47904dc82d661c9e4d8)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /soft/compilers/llvm/master-latest/bin
```</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>