[llvm-bugs] [Bug 51367] New: Target region mapping error when calling a custom target offload function that takes a lambda

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=51367

            Bug ID: 51367
           Summary: Target region mapping error when calling a custom
                    target offload function that takes a lambda
           Product: clang
           Version: 12.0
          Hardware: Other
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: OpenMP
          Assignee: unassignedclangbugs at nondot.org
          Reporter: khaled.3ttia at gmail.com
                CC: llvm-bugs at lists.llvm.org

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)

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20210805/a4c4438b/attachment.html>