1 // RUN: %libomptarget-compile-run-and-check-generic 2 3 // REQUIRES: libc 4 5 #include <assert.h> 6 #include <omp.h> 7 #include <stdio.h> 8 9 #pragma omp begin declare variant match(device = {kind(gpu)}) 10 // Extension provided by the 'libc' project. 11 unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size); 12 #pragma omp declare target to(__llvm_omp_host_call) device_type(nohost) 13 #pragma omp end declare variant 14 15 #pragma omp begin declare variant match(device = {kind(cpu)}) 16 // Dummy host implementation to make this work for all targets. 17 unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size) { 18 return ((unsigned long long (*)(void *))fn)(args); 19 } 20 #pragma omp end declare variant 21 22 long long foo(void *data) { return -1; } 23 24 void *fn_ptr = NULL; 25 #pragma omp declare target to(fn_ptr) 26 27 int main() { 28 fn_ptr = (void *)&foo; 29 #pragma omp target update to(fn_ptr) 30 31 for (int i = 0; i < 4; ++i) { 32 #pragma omp target 33 { 34 long long res = __llvm_omp_host_call(fn_ptr, NULL, 0); 35 assert(res == -1 && "RPC call failed\n"); 36 } 37 38 for (int j = 0; j < 128; ++j) { 39 #pragma omp target nowait 40 { 41 long long res = __llvm_omp_host_call(fn_ptr, NULL, 0); 42 assert(res == -1 && "RPC call failed\n"); 43 } 44 } 45 #pragma omp taskwait 46 47 #pragma omp target 48 { 49 long long res = __llvm_omp_host_call(fn_ptr, NULL, 0); 50 assert(res == -1 && "RPC call failed\n"); 51 } 52 } 53 54 // CHECK: PASS 55 puts("PASS"); 56 } 57