1 // RUN: %libomptarget-compilexx-generic 2 // RUN: %libomptarget-run-generic %fcheck-generic 3 4 // REQUIRES: gpu 5 // UNSUPPORTED: nvidiagpu 6 // UNSUPPORTED: amdgpu 7 8 #include <cstdio> 9 10 #include <omp.h> 11 12 extern "C" { 13 void *llvm_omp_target_lock_mem(void *ptr, size_t size, int device_num); 14 void llvm_omp_target_unlock_mem(void *ptr, int device_num); 15 } 16 17 int main() { 18 int n = 100; 19 int *unlocked = new int[n]; 20 21 for (int i = 0; i < n; i++) 22 unlocked[i] = i; 23 24 int *locked = (int *)llvm_omp_target_lock_mem(unlocked, n * sizeof(int), 25 omp_get_default_device()); 26 if (!locked) 27 return 0; 28 29 #pragma omp target teams distribute parallel for map(tofrom : unlocked[ : n]) 30 for (int i = 0; i < n; i++) 31 unlocked[i] += 1; 32 33 #pragma omp target teams distribute parallel for map(tofrom : unlocked[10 : 10]) 34 for (int i = 10; i < 20; i++) 35 unlocked[i] += 1; 36 37 #pragma omp target teams distribute parallel for map(tofrom : locked[ : n]) 38 for (int i = 0; i < n; i++) 39 locked[i] += 1; 40 41 #pragma omp target teams distribute parallel for map(tofrom : locked[10 : 10]) 42 for (int i = 10; i < 20; i++) 43 locked[i] += 1; 44 45 llvm_omp_target_unlock_mem(unlocked, omp_get_default_device()); 46 47 int err = 0; 48 for (int i = 0; i < n; i++) { 49 if (i < 10 || i > 19) { 50 if (unlocked[i] != i + 2) { 51 printf("Err at %d, got %d, expected %d\n", i, unlocked[i], i + 1); 52 err++; 53 } 54 } else if (unlocked[i] != i + 4) { 55 printf("Err at %d, got %d, expected %d\n", i, unlocked[i], i + 2); 56 err++; 57 } 58 } 59 60 // CHECK: PASS 61 if (err == 0) 62 printf("PASS\n"); 63 64 return err; 65 } 66