1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137
|
// RUN: %libomptarget-compile-run-and-check-generic
// REQUIRES: unified_shared_memory
// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
// amdgpu runtime crash
// UNSUPPORTED: 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];
//
// Test that updates on the device are not visible to host
// when only a TO mapping is used.
//
#pragma omp target map(tofrom \
: device_data, device_alloc) map(close, to \
: alloc[:N], data \
[:N])
{
device_data = &data[0];
device_alloc = &alloc[0];
for (int i = 0; i < N; i++) {
alloc[i] += 1;
data[i] += 1;
}
}
// CHECK: Address of alloc on device different from host address.
if (device_alloc != host_alloc)
printf("Address of alloc on device different from host address.\n");
// CHECK: Address of data on device different from host address.
if (device_data != host_data)
printf("Address of data on device different from host address.\n");
// On the host, check that the arrays have been updated.
// CHECK: Alloc host values not updated: Succeeded
fails = 0;
for (int i = 0; i < N; i++) {
if (alloc[i] != 10)
fails++;
}
printf("Alloc host values not updated: %s\n",
(fails == 0) ? "Succeeded" : "Failed");
// CHECK: Data host values not updated: Succeeded
fails = 0;
for (int i = 0; i < N; i++) {
if (data[i] != 1)
fails++;
}
printf("Data host values not updated: %s\n",
(fails == 0) ? "Succeeded" : "Failed");
//
// Test that updates on the device are visible on host
// when a from is used.
//
for (int i = 0; i < N; i++) {
alloc[i] += 1;
data[i] += 1;
}
#pragma omp target map(close, tofrom : alloc[:N], data[:N])
{
// CHECK: Alloc device values are correct: Succeeded
fails = 0;
for (int i = 0; i < N; i++) {
if (alloc[i] != 11)
fails++;
}
printf("Alloc device values are correct: %s\n",
(fails == 0) ? "Succeeded" : "Failed");
// CHECK: Data device values are correct: Succeeded
fails = 0;
for (int i = 0; i < N; i++) {
if (data[i] != 2)
fails++;
}
printf("Data device values are correct: %s\n",
(fails == 0) ? "Succeeded" : "Failed");
// Update values on the device
for (int i = 0; i < N; i++) {
alloc[i] += 1;
data[i] += 1;
}
}
// CHECK: Alloc host values updated: Succeeded
fails = 0;
for (int i = 0; i < N; i++) {
if (alloc[i] != 12)
fails++;
}
printf("Alloc host values updated: %s\n",
(fails == 0) ? "Succeeded" : "Failed");
// CHECK: Data host values updated: Succeeded
fails = 0;
for (int i = 0; i < N; i++) {
if (data[i] != 3)
fails++;
}
printf("Data host values updated: %s\n",
(fails == 0) ? "Succeeded" : "Failed");
free(alloc);
// CHECK: Done!
printf("Done!\n");
return 0;
}
|