62 lines · cpp
1// clang-format off2// RUN: %libomptarget-compilexx-generic3// RUN: env LIBOMPTARGET_INFO=32 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=NO-USM4//5// RUN: %libomptarget-compilexxx-generic-force-usm6// RUN: env HSA_XNACK=1 LIBOMPTARGET_INFO=32 \7// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=FORCE-USM8//9// REQUIRES: unified_shared_memory10//11// UNSUPPORTED: nvptx64-nvidia-cuda12// UNSUPPORTED: nvptx64-nvidia-cuda-LTO13// clang-format on14 15#include <cassert>16#include <cstdio>17#include <cstdlib>18 19int GI;20#pragma omp declare target21int *pGI;22#pragma omp end declare target23 24int main(void) {25 26 GI = 0;27 // Implicit mappings28 int alpha = 1;29 int beta[3] = {2, 5, 8};30 31 // Require map clauses for non-USM execution32 pGI = (int *)malloc(sizeof(int));33 *pGI = 42;34 35#pragma omp target map(pGI[ : 1], GI)36 {37 GI = 1 * alpha;38 *pGI = 2 * beta[1];39 }40 41 assert(GI == 1);42 assert(*pGI == 10);43 44 printf("SUCCESS\n");45 46 return 0;47}48 49// clang-format off50// NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=451// NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=1252// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=453// NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=8, Name=pGI54// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=455// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=1256// NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=457// NO-USM-NEXT: SUCCESS58 59// FORCE-USM: SUCCESS60//61// clang-format on62