llvm/offload/test/mapping/present/target_array_extension.c

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

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

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

// RUN: %libomptarget-compile-generic \
// RUN:   -DEXTENDS=AFTER
// RUN: %libomptarget-run-fail-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_SIZE (SMALL_END - SMALL_BEG)
#define LARGE_SIZE (LARGE_END - LARGE_BEG)

#define SMALL                                                                  \
  SMALL_BEG:                                                                   \
  SMALL_SIZE
#define LARGE                                                                  \
  LARGE_BEG:                                                                   \
  LARGE_SIZE

int main() {
  int arr[SIZE];

  // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]]
  fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG],
          SMALL_SIZE * sizeof arr[0]);

  // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]]
  fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG],
          LARGE_SIZE * sizeof arr[0]);

  // CHECK-NOT: omptarget
#pragma omp target data map(alloc : arr[LARGE])
  {
#pragma omp target map(present, tofrom : arr[SMALL])
    ;
  }

  // CHECK: arr is present
  fprintf(stderr, "arr is present\n");

  // CHECK: omptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
  // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
  // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
  // CHECK: omptarget error: Call to targetDataBegin failed, abort target.
  // CHECK: omptarget error: Failed to process data before launching the kernel.
  // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
#pragma omp target data map(alloc : arr[SMALL])
  {
#pragma omp target map(present, tofrom : arr[LARGE])
    ;
  }

  // CHECK-NOT: arr is present
  fprintf(stderr, "arr is present\n");

  return 0;
}