138 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2 3// REQUIRES: unified_shared_memory4// UNSUPPORTED: clang-6, clang-7, clang-8, clang-95 6// amdgpu runtime crash7// Fails on nvptx with error: an illegal memory access was encountered8// UNSUPPORTED: amdgcn-amd-amdhsa9// UNSUPPORTED: nvptx64-nvidia-cuda10// UNSUPPORTED: nvptx64-nvidia-cuda-LTO11 12#include <omp.h>13#include <stdio.h>14 15#pragma omp requires unified_shared_memory16 17#define N 102418 19int main(int argc, char *argv[]) {20 int fails;21 void *host_alloc, *device_alloc;22 void *host_data, *device_data;23 int *alloc = (int *)malloc(N * sizeof(int));24 int data[N];25 26 for (int i = 0; i < N; ++i) {27 alloc[i] = 10;28 data[i] = 1;29 }30 31 host_data = &data[0];32 host_alloc = &alloc[0];33 34//35// Test that updates on the device are not visible to host36// when only a TO mapping is used.37//38#pragma omp target map(tofrom : device_data, device_alloc) \39 map(close, to : alloc[ : N], data[ : N])40 {41 device_data = &data[0];42 device_alloc = &alloc[0];43 44 for (int i = 0; i < N; i++) {45 alloc[i] += 1;46 data[i] += 1;47 }48 }49 50 // CHECK: Address of alloc on device different from host address.51 if (device_alloc != host_alloc)52 printf("Address of alloc on device different from host address.\n");53 54 // CHECK: Address of data on device different from host address.55 if (device_data != host_data)56 printf("Address of data on device different from host address.\n");57 58 // On the host, check that the arrays have been updated.59 // CHECK: Alloc host values not updated: Succeeded60 fails = 0;61 for (int i = 0; i < N; i++) {62 if (alloc[i] != 10)63 fails++;64 }65 printf("Alloc host values not updated: %s\n",66 (fails == 0) ? "Succeeded" : "Failed");67 68 // CHECK: Data host values not updated: Succeeded69 fails = 0;70 for (int i = 0; i < N; i++) {71 if (data[i] != 1)72 fails++;73 }74 printf("Data host values not updated: %s\n",75 (fails == 0) ? "Succeeded" : "Failed");76 77 //78 // Test that updates on the device are visible on host79 // when a from is used.80 //81 82 for (int i = 0; i < N; i++) {83 alloc[i] += 1;84 data[i] += 1;85 }86 87#pragma omp target map(close, tofrom : alloc[ : N], data[ : N])88 {89 // CHECK: Alloc device values are correct: Succeeded90 fails = 0;91 for (int i = 0; i < N; i++) {92 if (alloc[i] != 11)93 fails++;94 }95 printf("Alloc device values are correct: %s\n",96 (fails == 0) ? "Succeeded" : "Failed");97 // CHECK: Data device values are correct: Succeeded98 fails = 0;99 for (int i = 0; i < N; i++) {100 if (data[i] != 2)101 fails++;102 }103 printf("Data device values are correct: %s\n",104 (fails == 0) ? "Succeeded" : "Failed");105 106 // Update values on the device107 for (int i = 0; i < N; i++) {108 alloc[i] += 1;109 data[i] += 1;110 }111 }112 113 // CHECK: Alloc host values updated: Succeeded114 fails = 0;115 for (int i = 0; i < N; i++) {116 if (alloc[i] != 12)117 fails++;118 }119 printf("Alloc host values updated: %s\n",120 (fails == 0) ? "Succeeded" : "Failed");121 122 // CHECK: Data host values updated: Succeeded123 fails = 0;124 for (int i = 0; i < N; i++) {125 if (data[i] != 3)126 fails++;127 }128 printf("Data host values updated: %s\n",129 (fails == 0) ? "Succeeded" : "Failed");130 131 free(alloc);132 133 // CHECK: Done!134 printf("Done!\n");135 136 return 0;137}138