================ @@ -0,0 +1,158 @@ +// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 +// %libomptarget-run-generic 2>&1 | %fcheck-generic + +// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9 + +// REQUIRES: amdgcn-amd-amdhsa + +#include <omp.h> +#include <stdio.h> + +#pragma omp requires unified_shared_memory + +#define N 1024 + +int main(int argc, char *argv[]) { + int fails; + void *host_alloc, *device_alloc; + void *host_data, *device_data; + int *alloc = (int *)malloc(N * sizeof(int)); + int data[N]; + + for (int i = 0; i < N; ++i) { + alloc[i] = 10; + data[i] = 1; + } + + host_data = &data[0]; + host_alloc = &alloc[0]; + +// CHECK: Creating new map entry ONLY with +// HstPtrBase=[[DEVICE_DATA_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_DATA_HST_PTR]], +// TgtAllocBegin=[[DEVICE_DATA_HST_PTR]], TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], +// Size=8, DynRefCount=1, HoldRefCount=0 CHECK: Creating new map entry ONLY with +// HstPtrBase=[[DATA_HST_PTR:0x.*]], HstPtrBegin=[[DATA_HST_PTR]], +// TgtAllocBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, +// DynRefCount=1, HoldRefCount=0 CHECK: Creating new map entry ONLY with +// HstPtrBase=[[DEVICE_ALLOC_HST_PTR:0x.*]], +// HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]], +// TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0 + +// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]], +// TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=1 (update +// suppressed), HoldRefCount=0 CHECK: Mapping exists with +// HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, +// DynRefCount=1 (update suppressed), HoldRefCount=0 CHECK: Mapping exists with +// HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], +// Size=8, DynRefCount=1 (update suppressed), HoldRefCount=0 + +// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks +// and 256 threads in Generic mode + +// CHECK: Mapping exists with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], +// TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=0 (decremented, +// delayed deletion), HoldRefCount=0 CHECK: Mapping exists with +// HstPtrBegin=[[DATA_HST_PTR]], TgtPtrBegin=[[DATA_HST_PTR]], Size=4096, +// DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0 CHECK: Mapping +// exists with HstPtrBegin=[[DEVICE_DATA_HST_PTR]], +// TgtPtrBegin=[[DEVICE_DATA_HST_PTR]], Size=8, DynRefCount=0 (decremented, +// delayed deletion), HoldRefCount=0 + +// CHECK: Removing map entry with HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]]{{.*}} +// Size=8 CHECK: Removing map entry with HstPtrBegin=[[DATA_HST_PTR]]{{.*}} +// Size=4096 CHECK: Removing map entry with +// HstPtrBegin=[[DEVICE_DATA_HST_PTR]]{{.*}} Size=8 + +// implicit mapping of data +#pragma omp target map(tofrom : device_data, device_alloc) + { + device_data = &data[0]; + device_alloc = &alloc[0]; + + for (int i = 0; i < N; i++) { + alloc[i] += 1; + data[i] += 1; + } + } + + if (device_alloc == host_alloc) + printf("Address of alloc on device matches host address.\n"); + + if (device_data == host_data) + printf("Address of data on device matches host address.\n"); + + // On the host, check that the arrays have been updated. + fails = 0; + for (int i = 0; i < N; i++) { + if (alloc[i] != 11) + fails++; + } + printf("Alloc device values updated: %s\n", + (fails == 0) ? "Succeeded" : "Failed"); + + fails = 0; + for (int i = 0; i < N; i++) { + if (data[i] != 2) + fails++; + } + printf("Data device values updated: %s\n", + (fails == 0) ? "Succeeded" : "Failed"); + + // + // Test that updates on the host snd on the device are both visible. ---------------- carlobertolli wrote:
snd-->and https://github.com/llvm/llvm-project/pull/69005 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits