Libomptarget message: explicit extension not allowed

[RaviNarayanaswamy]

When the following code is compiled and run generate the following message Libomptarget message: explicit extension not allowed: host address specified is 0x00007fff66cc45a4 (12 bytes), but device allocation maps to host at 0x00007fff66cc4530 (120 bytes) The ref count may not be correct when mappers are used the memory mapped in foo is not freed and when bar is called and uses the same stack memory runtime sees it as trying to extend an already mapped memory.

compiled using -fopenmp -fopenmp-targets=x86_64 Source 👍

#include <stdio.h>

#define TRUE 1
#define FALSE 0

typedef struct {
  int i1, i2, i3;
} TY1;

#pragma omp declare mapper(TY1 t) map(to:t.i1) map(from: t.i3)

unsigned foo() {
  TY1 t1[10];
  int i;
  for(int i = 0; i < 10; i++)
    t1[i].i1 = 1;

#pragma omp target map(tofrom:t1)
  for (i = 0; i < 10; i++)
    t1[i].i3 = t1[i].i1;

  for (int i = 0; i < 10; i++) {
    if (t1[i].i3 != t1[i].i1) {
      printf("failed. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", i, t1[i].i3, i, t1[i].i1);
      return FALSE;
    }
  }

  return TRUE;
}

unsigned bar() {
  TY1 t2[5];

  for(int i = 0; i < 5; i++)
    t2[i].i1 = 2;

#pragma omp target map(tofrom:t2)
  for (int i = 0; i < 5; i++)
    t2[i].i3 = t2[i].i1;

  for (int i = 0; i < 5; i++) {
    if (t2[i].i3 != t2[i].i1) {
      printf("failed. t2[%d].i3 (%d) != t2[%d].i1 (%d)\n", i, t2[i].i3, i, t2[i].i1);
      return FALSE;
    }
  }

  return TRUE;
}

int main() {
  if (foo() != TRUE)
    return 1;
  if (bar() != TRUE)
    return 2;
  printf("passed\n");
  return 0;
}

[llvmbot]

@llvm/issue-subscribers-openmp

[jdoerfert]

We fail to remove the t1 mapping and that clashes with t2 later. Looking at it further, there is dead code and way to little testing.

I'll propose another workaround but we need more tests because this simply breaks something else otherwise...

[jdoerfert]

https://github.com/llvm/llvm-project/commit/dbd6759dac52fa91d3321aaf0d9d78fbe3772dd4

has this test but also one that shows the problem w/o relying on the stack to be allocated at the same place.

[ceseo]

@jdoerfert do you have any plans to submit a fix for this? I'lm also seeing this on aarch64.

[ceseo]

I'm seeing something that I think it's same issue as here in target_dereference_array_pointers.cpp on AArch64:

void foo(int **t1d) {
  int ***t2d = &t1d;
  int ****t3d = &t2d;
  *t1d = (int *)malloc(3 * sizeof(int));
  int j, a = 0, b = 0;

  for (j = 0; j < 3; j++)
    (*t1d)[j] = 0;
#pragma omp target map(tofrom : (*t1d)[0 : 3])
  { (*t1d)[1] = 1; }
  // CHECK: 1
  printf("%d\n", (*t1d)[1]);
#pragma omp target map(tofrom : (**t2d)[0 : 3])
  { (**t2d)[1] = 2; }
  // CHECK: 2
  printf("%d\n", (**t2d)[1]);
#pragma omp target map(tofrom : (***t3d)[0 : 3])
  { (***t3d)[1] = 3; }
  // CHECK: 3
  printf("%d\n", (***t3d)[1]);
#pragma omp target map(tofrom : (**t1d))
  { (*t1d)[0] = 4; }
  // CHECK: 4
  printf("%d\n", (*t1d)[0]);
#pragma omp target map(tofrom : (*(*(t1d + a) + b)))
  { *(*(t1d + a) + b) = 5; }
  // CHECK: 5
  printf("%d\n", *(*(t1d + a) + b));
}

Looks like ****t3d is clashing with ***t2d.

Is this testcase correct? I'm not 100% sure...

[jdoerfert]

I'm seeing something that I think it's same issue as here in target_dereference_array_pointers.cpp on AArch64:

void foo(int **t1d) {
  int ***t2d = &t1d;
  int ****t3d = &t2d;
  *t1d = (int *)malloc(3 * sizeof(int));
  int j, a = 0, b = 0;

  for (j = 0; j < 3; j++)
    (*t1d)[j] = 0;
#pragma omp target map(tofrom : (*t1d)[0 : 3])
  { (*t1d)[1] = 1; }
  // CHECK: 1
  printf("%d\n", (*t1d)[1]);
#pragma omp target map(tofrom : (**t2d)[0 : 3])
  { (**t2d)[1] = 2; }
  // CHECK: 2
  printf("%d\n", (**t2d)[1]);
#pragma omp target map(tofrom : (***t3d)[0 : 3])
  { (***t3d)[1] = 3; }
  // CHECK: 3
  printf("%d\n", (***t3d)[1]);
#pragma omp target map(tofrom : (**t1d))
  { (*t1d)[0] = 4; }
  // CHECK: 4
  printf("%d\n", (*t1d)[0]);
#pragma omp target map(tofrom : (*(*(t1d + a) + b)))
  { *(*(t1d + a) + b) = 5; }
  // CHECK: 5
  printf("%d\n", *(*(t1d + a) + b));
}

Looks like ****t3d is clashing with ***t2d.

Is this testcase correct? I'm not 100% sure...

can you open a new issue. I have a fix for the original problem and it’s related to user defined mapper handling, your issue is not.

[ceseo]

@jdoerfert done!