libomptarget/test/offloading/force-usm.cpp

// clang-format off
// RUN: %libomptarget-compilexx-generic
// RUN: env LIBOMPTARGET_INFO=32 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=NO-USM
//
// RUN: %libomptarget-compilexxx-generic-force-usm
// RUN: env HSA_XNACK=1 LIBOMPTARGET_INFO=32 \
// RUN:       %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=FORCE-USM
//
// REQUIRES: unified_shared_memory
//
// UNSUPPORTED: nvptx64-nvidia-cuda
// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
// clang-format on

#include <cassert>
#include <cstdio>
#include <cstdlib>

int GI;
#pragma omp declare target
int *pGI;
#pragma omp end declare target

int main(void) {

  GI = 0;
  // Implicit mappings
  int alpha = 1;
  int beta[3] = {2, 5, 8};

  // Require map clauses for non-USM execution
  pGI = (int *)malloc(sizeof(int));
  *pGI = 42;

#pragma omp target map(pGI[ : 1], GI)
  {
    GI = 1 * alpha;
    *pGI = 2 * beta[1];
  }

  assert(GI == 1);
  assert(*pGI == 10);

  printf("SUCCESS\n");

  return 0;
}

// clang-format off
// NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=12
// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=8, Name=pGI
// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=12
// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
// NO-USM-NEXT: SUCCESS

// FORCE-USM: SUCCESS
//
// clang-format on