This is an archive of the discontinued LLVM Phabricator instance.

[OPENMP]Fix PR46012: declare target pointer cannot be accessed in target region.
ClosedPublic

Authored by ABataev on Jul 20 2020, 9:24 AM.

Details

Summary

Need to avoid an optimization for base pointer mapping for target data
directives.

Diff Detail

Event Timeline

ABataev created this revision.Jul 20 2020, 9:24 AM
Herald added a project: Restricted Project. · View Herald TranscriptJul 20 2020, 9:24 AM
ye-luo accepted this revision.Jul 21 2020, 8:29 AM

I verified that 46012 is fixed with this patch

This revision is now accepted and ready to land.Jul 21 2020, 8:29 AM
This revision was automatically updated to reflect the committed changes.
grokos added a subscriber: grokos.Jul 24 2020, 4:28 PM

@ABataev:

After this patch was committed, I tried to run the following example:

#include <stdio.h>

int *yptr;

int main() {
  int y[10];
  y[1] = 1;
  yptr = &y[0];

  printf("&yptr = %p\n", &yptr);
  printf("&y[0] = %p\n", &y[0]);

  #pragma omp target data map(to: yptr[0:5])
  #pragma omp target
  {
    printf("y = %d\n", yptr[1]);
    yptr[1] = 10;
    printf("y = %d\n", yptr[1]);
  }

  printf("y = %d\n", yptr[1]);
  return 0;
}

The arguments clang generates are:

1) base = &y[0], begin = &yptr, size = 8, type = TARGET_PARAM | TO
2) base = &yptr, begin = &y[0], size = 8, type = PTR_AND_OBJ | TO

The second argument is correct, the first argument doesn't make much sense. I believe it should have its base set to &yptr, not &y[0].
y[0] is not the base for anything, it's only the pointee object.

@ABataev:

After this patch was committed, I tried to run the following example:

#include <stdio.h>

int *yptr;

int main() {
  int y[10];
  y[1] = 1;
  yptr = &y[0];

  printf("&yptr = %p\n", &yptr);
  printf("&y[0] = %p\n", &y[0]);

  #pragma omp target data map(to: yptr[0:5])
  #pragma omp target
  {
    printf("y = %d\n", yptr[1]);
    yptr[1] = 10;
    printf("y = %d\n", yptr[1]);
  }

  printf("y = %d\n", yptr[1]);
  return 0;
}

The arguments clang generates are:

1) base = &y[0], begin = &yptr, size = 8, type = TARGET_PARAM | TO
2) base = &yptr, begin = &y[0], size = 8, type = PTR_AND_OBJ | TO

The second argument is correct, the first argument doesn't make much sense. I believe it should have its base set to &yptr, not &y[0].
y[0] is not the base for anything, it's only the pointee object.

Yes, I know, investigating it.

@ABataev:

After this patch was committed, I tried to run the following example:

#include <stdio.h>

int *yptr;

int main() {
  int y[10];
  y[1] = 1;
  yptr = &y[0];

  printf("&yptr = %p\n", &yptr);
  printf("&y[0] = %p\n", &y[0]);

  #pragma omp target data map(to: yptr[0:5])
  #pragma omp target
  {
    printf("y = %d\n", yptr[1]);
    yptr[1] = 10;
    printf("y = %d\n", yptr[1]);
  }

  printf("y = %d\n", yptr[1]);
  return 0;
}

The arguments clang generates are:

1) base = &y[0], begin = &yptr, size = 8, type = TARGET_PARAM | TO
2) base = &yptr, begin = &y[0], size = 8, type = PTR_AND_OBJ | TO

The second argument is correct, the first argument doesn't make much sense. I believe it should have its base set to &yptr, not &y[0].
y[0] is not the base for anything, it's only the pointee object.

I hit the same issue and it results incorrect mapping.
https://bugs.llvm.org/show_bug.cgi?id=46868
In addition, only yptr is touched inside the target region. After this patch, two mappings instead of one occurs.

@ABataev:

After this patch was committed, I tried to run the following example:

#include <stdio.h>

int *yptr;

int main() {
  int y[10];
  y[1] = 1;
  yptr = &y[0];

  printf("&yptr = %p\n", &yptr);
  printf("&y[0] = %p\n", &y[0]);

  #pragma omp target data map(to: yptr[0:5])
  #pragma omp target
  {
    printf("y = %d\n", yptr[1]);
    yptr[1] = 10;
    printf("y = %d\n", yptr[1]);
  }

  printf("y = %d\n", yptr[1]);
  return 0;
}

The arguments clang generates are:

1) base = &y[0], begin = &yptr, size = 8, type = TARGET_PARAM | TO
2) base = &yptr, begin = &y[0], size = 8, type = PTR_AND_OBJ | TO

The second argument is correct, the first argument doesn't make much sense. I believe it should have its base set to &yptr, not &y[0].
y[0] is not the base for anything, it's only the pointee object.

I hit the same issue and it results incorrect mapping.
https://bugs.llvm.org/show_bug.cgi?id=46868
In addition, only yptr is touched inside the target region. After this patch, two mappings instead of one occurs.

The fix is almost ready, working on tests update