66 lines · cpp
1// RUN: %libomptarget-compilexx-generic2// RUN: %libomptarget-run-generic %fcheck-generic3 4// REQUIRES: gpu5// UNSUPPORTED: nvidiagpu6// UNSUPPORTED: amdgpu7 8#include <cstdio>9 10#include <omp.h>11 12extern "C" {13void *llvm_omp_target_lock_mem(void *ptr, size_t size, int device_num);14void llvm_omp_target_unlock_mem(void *ptr, int device_num);15}16 17int 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: PASS61 if (err == 0)62 printf("PASS\n");63 64 return err;65}66