48 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2 3#include <assert.h>4#include <stdio.h>5 6#pragma omp begin declare variant match(device = {kind(gpu)})7// Provided by the runtime.8void *__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 variant12 13#pragma omp begin declare variant match(device = {kind(cpu)})14// We assume unified addressing on the CPU target.15void *__llvm_omp_indirect_call_lookup(void *host_ptr) { return host_ptr; }16#pragma omp end declare variant17 18#pragma omp begin declare target indirect19void foo(int *x) { *x = *x + 1; }20void bar(int *x) { *x = *x + 1; }21void baz(int *x) { *x = *x + 1; }22#pragma omp end declare target23 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: PASS46 printf("PASS\n");47}48