offload/test/mapping/target_data_array_extension_at_exit.c

// --------------------------------------------------
// Check extends before
// --------------------------------------------------

// RUN: %libomptarget-compile-generic \
// RUN:   -DEXTENDS=BEFORE
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic

// --------------------------------------------------
// Check extends after
// --------------------------------------------------

// RUN: %libomptarget-compile-generic \
// RUN:   -DEXTENDS=AFTER
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic

// END.

#include <stdio.h>

#define BEFORE 0
#define AFTER 1

#define SIZE 100

#if EXTENDS == BEFORE
#define SMALL_BEG (SIZE - 2)
#define SMALL_END SIZE
#define LARGE_BEG 0
#define LARGE_END SIZE
#elif EXTENDS == AFTER
#define SMALL_BEG 0
#define SMALL_END 2
#define LARGE_BEG 0
#define LARGE_END SIZE
#else
#error EXTENDS undefined
#endif

#define SMALL                                                                  \
  SMALL_BEG:                                                                   \
  (SMALL_END - SMALL_BEG)
#define LARGE                                                                  \
  LARGE_BEG:                                                                   \
  (LARGE_END - LARGE_BEG)

void check_not_present() {
  int arr[SIZE];

  for (int i = 0; i < SIZE; ++i)
    arr[i] = 99;

  // CHECK-LABEL: checking not present
  fprintf(stderr, "checking not present\n");

  // arr[LARGE] isn't (fully) present at the end of the target data region, so
  // the device-to-host transfer should not be performed, or it might fail.
#pragma omp target data map(tofrom : arr[LARGE])
  {
#pragma omp target exit data map(delete : arr[LARGE])
#pragma omp target enter data map(alloc : arr[SMALL])
#pragma omp target map(alloc : arr[SMALL])
    for (int i = SMALL_BEG; i < SMALL_END; ++i)
      arr[i] = 88;
  }

  // CHECK-NOT: omptarget
  // CHECK-NOT: error
  for (int i = 0; i < SIZE; ++i) {
    if (arr[i] != 99)
      fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
  }
}

void check_is_present() {
  int arr[SIZE];

  for (int i = 0; i < SIZE; ++i)
    arr[i] = 99;

  // CHECK-LABEL: checking is present
  fprintf(stderr, "checking is present\n");

  // arr[SMALL] is (fully) present at the end of the target data region, and the
  // device-to-host transfer should be performed only for it even though more
  // of the array is then present.
#pragma omp target data map(tofrom : arr[SMALL])
  {
#pragma omp target exit data map(delete : arr[SMALL])
#pragma omp target enter data map(alloc : arr[LARGE])
#pragma omp target map(alloc : arr[LARGE])
    for (int i = LARGE_BEG; i < LARGE_END; ++i)
      arr[i] = 88;
  }

  // CHECK-NOT: omptarget
  // CHECK-NOT: error
  for (int i = 0; i < SIZE; ++i) {
    if (SMALL_BEG <= i && i < SMALL_END) {
      if (arr[i] != 88)
        fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
    } else if (arr[i] != 99) {
      fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
    }
  }
}

int main() {
  check_not_present();
  check_is_present();
  return 0;
}