1 | // RUN: %libomptarget-compile-run-and-check-generic |
2 | |
3 | #include <assert.h> |
4 | #include <omp.h> |
5 | #include <stdio.h> |
6 | |
7 | #pragma omp begin declare variant match(device = {kind(gpu)}) |
8 | // Extension provided by the 'libc' project. |
9 | unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size); |
10 | #pragma omp declare target to(__llvm_omp_host_call) device_type(nohost) |
11 | #pragma omp end declare variant |
12 | |
13 | #pragma omp begin declare variant match(device = {kind(cpu)}) |
14 | // Dummy host implementation to make this work for all targets. |
15 | unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size) { |
16 | return ((unsigned long long (*)(void *))fn)(args); |
17 | } |
18 | #pragma omp end declare variant |
19 | |
20 | typedef struct args_s { |
21 | int thread_id; |
22 | int block_id; |
23 | } args_t; |
24 | |
25 | // CHECK-DAG: Thread: 0, Block: 0 |
26 | // CHECK-DAG: Thread: 1, Block: 0 |
27 | // CHECK-DAG: Thread: 0, Block: 1 |
28 | // CHECK-DAG: Thread: 1, Block: 1 |
29 | // CHECK-DAG: Thread: 0, Block: 2 |
30 | // CHECK-DAG: Thread: 1, Block: 2 |
31 | // CHECK-DAG: Thread: 0, Block: 3 |
32 | // CHECK-DAG: Thread: 1, Block: 3 |
33 | unsigned long long foo(void *data) { |
34 | assert(omp_is_initial_device() && "Not executing on host?" ); |
35 | args_t *args = (args_t *)data; |
36 | printf(format: "Thread: %d, Block: %d\n" , args->thread_id, args->block_id); |
37 | return 42; |
38 | } |
39 | |
40 | void *fn_ptr = NULL; |
41 | #pragma omp declare target to(fn_ptr) |
42 | |
43 | int main() { |
44 | fn_ptr = (void *)&foo; |
45 | #pragma omp target update to(fn_ptr) |
46 | |
47 | int failed = 0; |
48 | #pragma omp target teams num_teams(4) map(tofrom : failed) |
49 | #pragma omp parallel num_threads(2) |
50 | { |
51 | args_t args = {omp_get_thread_num(), omp_get_team_num()}; |
52 | unsigned long long res = |
53 | __llvm_omp_host_call(fn: fn_ptr, args: &args, size: sizeof(args_t)); |
54 | if (res != 42) |
55 | #pragma omp atomic write |
56 | failed = 1; |
57 | } |
58 | |
59 | // CHECK: PASS |
60 | if (!failed) |
61 | printf(format: "PASS\n" ); |
62 | } |
63 | |