offload/test/api/omp_indirect_call.c

   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.
  15 void *__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
  19 void foo(int *x) { *x = *x + 1; }
  20 void bar(int *x) { *x = *x + 1; }
  21 void baz(int *x) { *x = *x + 1; }
  22 #pragma omp end declare target
  23
  24 int 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 }