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: fn_ptr, NULL, size: 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: fn_ptr, NULL, size: 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: fn_ptr, NULL, size: 0); |
50 | assert(res == -1 && "RPC call failed\n"); |
51 | } |
52 | } |
53 | |
54 | // CHECK: PASS |
55 | puts(s: "PASS"); |
56 | } |
57 |