<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 - Target region mapping error when calling a custom target offload function that takes a lambda"
href="https://bugs.llvm.org/show_bug.cgi?id=51367">51367</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>Target region mapping error when calling a custom target offload function that takes a lambda
</td>
</tr>
<tr>
<th>Product</th>
<td>clang
</td>
</tr>
<tr>
<th>Version</th>
<td>12.0
</td>
</tr>
<tr>
<th>Hardware</th>
<td>Other
</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>khaled.3ttia@gmail.com
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr></table>
<p>
<div>
<pre>Hello,
I have been trying to build a custom function that takes a lambda and offloads
it using omp target directive. When doing it this way, mapping of objects used
inside the lambda is done incorrectly and ignores the specified declare mapper
rules. Full code example below
#include <omp.h>
// A wrapper class around raw pointers
class MyClass {
public:
int *_data;
size_t _len;
// ctor
MyClass(int *src, size_t len) : _data(src), _len(len) {}
// overload [] operator
int &operator[](int idx) { return _data[idx]; }
};
// declare mapper for MyClass objects
#pragma omp declare mapper(MyClass a) map(from : a._data [0:a._len])
// Function to offload a lambda expression to device
template <typename Func> inline void offloadme(Func logic) {
#pragma omp target map(tofrom : logic)
logic();
}
int main() {
int *a = new int[10];
MyClass acc(a, 10);
// call the offload function, passing a lambda
offloadme([&]() { acc[0] = 5; });
delete[] a;
return 0;
}
I compile using
clang++ -fopenmp -fopenmp-targets=nvptx64 -gline-tables-only -O3 main.cpp
when running LIBOMPTARGET_INFO=-1 ./a.out I get the following output
CUDA device 0 info: Device supports up to 65536 CUDA blocks and 1024 threads
with a warp size of 32
Libomptarget device 0 info: Entering OpenMP kernel at main.cpp:22:1 with 2
arguments:
Libomptarget device 0 info: tofrom(logic)[8]
Libomptarget device 0 info: firstprivate(acc)[16] (implicit)
Libomptarget device 0 info: Creating new map entry with
HstPtrBegin=0x00007ffe155c9488, TgtPtrBegin=0x00007f67c6400000, Size=8,
RefCount=1, Name=logic
Libomptarget device 0 info: Copying data from host to device,
HstPtr=0x00007ffe155c9488, TgtPtr=0x00007f67c6400000, Size=8, Name=logic
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x00007ffe155c9488,
TgtPtrBegin=0x00007f67c6400000, Size=8, RefCount=1 (update suppressed)
CUDA device 0 info: Launching kernel
__omp_offloading_10302_15410ef__Z9offloadmeIZ4mainE3$_0EvT__l22 with 1 blocks
and 33 threads in Generic mode
Libomptarget device 0 info: Mapping exists with HstPtrBegin=0x00007ffe155c9488,
TgtPtrBegin=0x00007f67c6400000, Size=8, RefCount=1 (deferred final decrement)
Libomptarget device 0 info: Copying data from device to host,
TgtPtr=0x00007f67c6400000, HstPtr=0x00007ffe155c9488, Size=8, Name=logic
CUDA error: an illegal memory access was encountered
Libomptarget error: Copying data from device failed.
Libomptarget error: Call to targetDataEnd failed, abort target.
Libomptarget error: Failed to process data after launching the kernel.
Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at
main.cpp:22:1:
Libomptarget device 0 info: Host Ptr Target Ptr Size (B)
RefCount Declaration
Libomptarget device 0 info: 0x00007ffe155c9488 0x00007f67c6400000 8 1
logic at main.cpp:20:53
main.cpp:22:1: Libomptarget fatal error 1: failure of target construct while
offloading is mandatory
Aborted (core dumped)</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>