111 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2 3// REQUIRES: unified_shared_memory4 5// amdgpu runtime crash6// Fails on nvptx with error: an illegal memory access was encountered7// UNSUPPORTED: amdgcn-amd-amdhsa8// UNSUPPORTED: nvptx64-nvidia-cuda9// UNSUPPORTED: nvptx64-nvidia-cuda-LTO10 11#include <omp.h>12#include <stdio.h>13 14// End of definitions copied from OpenMP RTL.15// ---------------------------------------------------------------------------16 17#pragma omp requires unified_shared_memory18 19#define N 102420 21int main(int argc, char *argv[]) {22 int fails;23 void *host_alloc, *device_alloc;24 void *host_data, *device_data;25 int *alloc = (int *)malloc(N * sizeof(int));26 int data[N];27 28 for (int i = 0; i < N; ++i) {29 alloc[i] = 10;30 data[i] = 1;31 }32 33 host_data = &data[0];34 host_alloc = &alloc[0];35 36// implicit mapping of data37#pragma omp target map(tofrom : device_data, device_alloc)38 {39 device_data = &data[0];40 device_alloc = &alloc[0];41 42 for (int i = 0; i < N; i++) {43 alloc[i] += 1;44 data[i] += 1;45 }46 }47 48 // CHECK: Address of alloc on device matches host address.49 if (device_alloc == host_alloc)50 printf("Address of alloc on device matches host address.\n");51 52 // CHECK: Address of data on device matches host address.53 if (device_data == host_data)54 printf("Address of data on device matches host address.\n");55 56 // On the host, check that the arrays have been updated.57 // CHECK: Alloc device values updated: Succeeded58 fails = 0;59 for (int i = 0; i < N; i++) {60 if (alloc[i] != 11)61 fails++;62 }63 printf("Alloc device values updated: %s\n",64 (fails == 0) ? "Succeeded" : "Failed");65 66 // CHECK: Data device values updated: Succeeded67 fails = 0;68 for (int i = 0; i < N; i++) {69 if (data[i] != 2)70 fails++;71 }72 printf("Data device values updated: %s\n",73 (fails == 0) ? "Succeeded" : "Failed");74 75 //76 // Test that updates on the host snd on the device are both visible.77 //78 79 // Update on the host.80 for (int i = 0; i < N; ++i) {81 alloc[i] += 1;82 data[i] += 1;83 }84 85#pragma omp target86 {87 // CHECK: Alloc host values updated: Succeeded88 fails = 0;89 for (int i = 0; i < N; i++) {90 if (alloc[i] != 12)91 fails++;92 }93 printf("Alloc host values updated: %s\n",94 (fails == 0) ? "Succeeded" : "Failed");95 // CHECK: Data host values updated: Succeeded96 fails = 0;97 for (int i = 0; i < N; i++) {98 if (data[i] != 3)99 fails++;100 }101 printf("Data host values updated: %s\n",102 (fails == 0) ? "Succeeded" : "Failed");103 }104 105 free(alloc);106 107 printf("Done!\n");108 109 return 0;110}111