offload/test/mapping/ompx_hold/struct.c

// RUN: %libomptarget-compile-generic -fopenmp-extensions
// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
// XFAIL: intelgpu

#include <omp.h>
#include <stdio.h>

#pragma omp begin declare target
#define MAX_NAME_SIZE 100
char N1[MAX_NAME_SIZE], N2[MAX_NAME_SIZE];
int V1, V2;
void copy_name(char *dst, char *src) {
  int i;
  for (i = 0; i < MAX_NAME_SIZE - 1 && src[i]; ++i)
    dst[i] = src[i];
  dst[i] = 0;
}
#pragma omp end declare target

#define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
  printf("    presence of %s, %s, %s: %d, %d, %d\n", #Var1, #Var2, #Var3,      \
         omp_target_is_present(&(Var1), omp_get_default_device()),             \
         omp_target_is_present(&(Var2), omp_get_default_device()),             \
         omp_target_is_present(&(Var3), omp_get_default_device()))

#define CHECK_VALUES_HELPER(N1, N2, Var1, Var2)                                \
  printf("    values of %s, %s: %d, %d\n", N1, N2, (Var1), (Var2))

#define CHECK_VALUES_DELAYED(Var1, Var2)                                       \
  copy_name(N1, #Var1);                                                        \
  copy_name(N2, #Var2);                                                        \
  V1 = (Var1);                                                                 \
  V2 = (Var2);

#define CHECK_DELAYED_VALUES()                                                 \
  _Pragma("omp target update from(N1, N2, V1, V2)")                            \
      CHECK_VALUES_HELPER(N1, N2, V1, V2)

#define CHECK_VALUES(Var1, Var2)                                               \
  CHECK_VALUES_HELPER(#Var1, #Var2, (Var1), (Var2))

int main() {
  struct S {
    int i;
    int j;
  } s;
  // CHECK: presence of s, s.i, s.j: 0, 0, 0
  CHECK_PRESENCE(s, s.i, s.j);

  // =======================================================================
  // Check that ompx_hold keeps entire struct present.

  // -----------------------------------------------------------------------
  // CHECK-LABEL: check:{{.*}}
  printf("check: ompx_hold only on first member\n");
  s.i = 20;
  s.j = 30;
#pragma omp target data map(tofrom : s) map(ompx_hold, tofrom : s.i)           \
    map(tofrom : s.j)
  {
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(tofrom : s)
    {
      s.i = 21;
      s.j = 31;
    }
#pragma omp target exit data map(delete : s, s.i)
    // ompx_hold on s.i applies to all of s.
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    // CHECK-NEXT: values of s.i, s.j: 20, 30
    CHECK_PRESENCE(s, s.i, s.j);
    CHECK_VALUES(s.i, s.j);
  }
  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
  // CHECK-NEXT: values of s.i, s.j: 21, 31
  CHECK_PRESENCE(s, s.i, s.j);
  CHECK_VALUES(s.i, s.j);

  // -----------------------------------------------------------------------
  // CHECK-LABEL: check:{{.*}}
  printf("check: ompx_hold only on last member\n");
  s.i = 20;
  s.j = 30;
#pragma omp target data map(tofrom : s) map(tofrom : s.i)                      \
    map(ompx_hold, tofrom : s.j)
  {
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(tofrom : s)
    {
      s.i = 21;
      s.j = 31;
    }
#pragma omp target exit data map(delete : s, s.i)
    // ompx_hold on s.j applies to all of s.
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    // CHECK-NEXT: values of s.i, s.j: 20, 30
    CHECK_PRESENCE(s, s.i, s.j);
    CHECK_VALUES(s.i, s.j);
  }
  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
  // CHECK-NEXT: values of s.i, s.j: 21, 31
  CHECK_PRESENCE(s, s.i, s.j);
  CHECK_VALUES(s.i, s.j);

  // -----------------------------------------------------------------------
  // CHECK-LABEL: check:{{.*}}
  printf("check: ompx_hold only on struct\n");
  s.i = 20;
  s.j = 30;
#pragma omp target data map(ompx_hold, tofrom : s) map(tofrom : s.i)           \
    map(tofrom : s.j)
  {
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(tofrom : s)
    {
      s.i = 21;
      s.j = 31;
    }
#pragma omp target exit data map(delete : s, s.i)
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    // CHECK-NEXT: values of s.i, s.j: 20, 30
    CHECK_PRESENCE(s, s.i, s.j);
    CHECK_VALUES(s.i, s.j);
  }
  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
  // CHECK-NEXT: values of s.i, s.j: 21, 31
  CHECK_PRESENCE(s, s.i, s.j);
  CHECK_VALUES(s.i, s.j);

  // =======================================================================
  // Check that transfer to/from host checks reference count correctly.

  // -----------------------------------------------------------------------
  // CHECK-LABEL: check:{{.*}}
  printf("check: parent DynRefCount=1 is not sufficient for transfer\n");
  s.i = 20;
  s.j = 30;
#pragma omp target data map(ompx_hold, tofrom : s)
#pragma omp target data map(ompx_hold, tofrom : s)
  {
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(from : s.i, s.j)
    {
      s.i = 21;
      s.j = 31;
    } // No transfer here even though parent's DynRefCount=1.
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    // CHECK-NEXT: values of s.i, s.j: 20, 30
    CHECK_PRESENCE(s, s.i, s.j);
    CHECK_VALUES(s.i, s.j);
#pragma omp target map(to : s.i, s.j)
    { // No transfer here even though parent's DynRefCount=1.
      // CHECK-NEXT: values of s.i, s.j: 21, 31
      CHECK_VALUES_DELAYED(s.i, s.j);
    }
    CHECK_DELAYED_VALUES();
  }
  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
  // CHECK-NEXT: values of s.i, s.j: 21, 31
  CHECK_PRESENCE(s, s.i, s.j);
  CHECK_VALUES(s.i, s.j);

  // -----------------------------------------------------------------------
  // CHECK-LABEL: check:{{.*}}
  printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");
  s.i = 20;
  s.j = 30;
#pragma omp target data map(tofrom : s)
#pragma omp target data map(tofrom : s)
  {
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(ompx_hold, from : s.i, s.j)
    {
      s.i = 21;
      s.j = 31;
    } // No transfer here even though parent's HoldRefCount=1.
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    // CHECK-NEXT: values of s.i, s.j: 20, 30
    CHECK_PRESENCE(s, s.i, s.j);
    CHECK_VALUES(s.i, s.j);
#pragma omp target map(ompx_hold, to : s.i, s.j)
    { // No transfer here even though parent's HoldRefCount=1.
      // CHECK-NEXT: values of s.i, s.j: 21, 31
      CHECK_VALUES_DELAYED(s.i, s.j);
    }
    CHECK_DELAYED_VALUES();
  }
  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
  // CHECK-NEXT: values of s.i, s.j: 21, 31
  CHECK_PRESENCE(s, s.i, s.j);
  CHECK_VALUES(s.i, s.j);

  // -----------------------------------------------------------------------
  // CHECK-LABEL: check:{{.*}}
  //
  // At the beginning of a region, if the parent's TotalRefCount=1, then the
  // transfer should happen.
  //
  // At the end of a region, it also must be true that the reference count being
  // decremented is the reference count that is 1.
  printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");
  s.i = 20;
  s.j = 30;
#pragma omp target data map(ompx_hold, tofrom : s)
  {
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(ompx_hold, tofrom : s.i, s.j)
    {
      s.i = 21;
      s.j = 31;
    }
#pragma omp target exit data map(from : s.i, s.j)
    // No transfer here even though parent's TotalRefCount=1.
    // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
    // CHECK-NEXT: values of s.i, s.j: 20, 30
    CHECK_PRESENCE(s, s.i, s.j);
    CHECK_VALUES(s.i, s.j);
  }
  // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
  // CHECK-NEXT: values of s.i, s.j: 21, 31
  CHECK_PRESENCE(s, s.i, s.j);
  CHECK_VALUES(s.i, s.j);

  return 0;
}