1 // RUN: %libomptarget-compile-run-and-check-generic 2 3 #include <assert.h> 4 #include <stdio.h> 5 6 #pragma omp begin declare variant match(device = {kind(gpu)}) 7 // Provided by the runtime. 8 void *__llvm_omp_indirect_call_lookup(void *host_ptr); 9 #pragma omp declare target to(__llvm_omp_indirect_call_lookup) \ 10 device_type(nohost) 11 #pragma omp end declare variant 12 13 #pragma omp begin declare variant match(device = {kind(cpu)}) 14 // We assume unified addressing on the CPU target. __llvm_omp_indirect_call_lookup(void * host_ptr)15void *__llvm_omp_indirect_call_lookup(void *host_ptr) { return host_ptr; } 16 #pragma omp end declare variant 17 18 #pragma omp begin declare target indirect foo(int * x)19void foo(int *x) { *x = *x + 1; } bar(int * x)20void bar(int *x) { *x = *x + 1; } baz(int * x)21void baz(int *x) { *x = *x + 1; } 22 #pragma omp end declare target 23 main()24int main() { 25 void *foo_ptr = foo; 26 void *bar_ptr = bar; 27 void *baz_ptr = baz; 28 29 int count = 0; 30 void *foo_res; 31 void *bar_res; 32 void *baz_res; 33 #pragma omp target map(to : foo_ptr, bar_ptr, baz_ptr) map(tofrom : count) 34 { 35 foo_res = __llvm_omp_indirect_call_lookup(foo_ptr); 36 ((void (*)(int *))foo_res)(&count); 37 bar_res = __llvm_omp_indirect_call_lookup(bar_ptr); 38 ((void (*)(int *))bar_res)(&count); 39 baz_res = __llvm_omp_indirect_call_lookup(baz_ptr); 40 ((void (*)(int *))baz_res)(&count); 41 } 42 43 assert(count == 3 && "Calling failed"); 44 45 // CHECK: PASS 46 printf("PASS\n"); 47 } 48