[llvm-bugs] [Bug 48218] New: Possible issue with target enter/exit data calls on objects with pointers inside

[via llvm-bugs]

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

            Bug ID: 48218
           Summary: Possible issue with target enter/exit data calls on
                    objects with pointers inside
           Product: OpenMP
           Version: unspecified
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: enhancement
          Priority: P
         Component: Clang Compiler Support
          Assignee: unassignedclangbugs at nondot.org
          Reporter: bertoni at anl.gov
                CC: llvm-bugs at lists.llvm.org

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.

```
> cat main.cpp 
#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;
}

> clang++ -fopenmp -fopenmp-targets=nvptx64 main.cpp

> ./a.out 
Expecting: 0, Got: 0
Expecting: 2, Got: 1
a.out: main.cpp:63: int main(): Assertion `objects[i].val == i*2' failed.
Aborted

> clang++ -DINVERT -fopenmp -fopenmp-targets=nvptx64 main.cpp

> ./a.out 
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

> clang++ --version
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
```

-- 
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/20201118/0835d44a/attachment.html>