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 | void rpc_host_call(void *fn, void *args, size_t size); |
12 | #pragma omp declare target to(rpc_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 | void rpc_host_call(void *fn, void *args, size_t size) { |
18 | ((void (*)(void *))fn)(args); |
19 | } |
20 | #pragma omp end declare variant |
21 | |
22 | typedef struct args_s { |
23 | int thread_id; |
24 | int block_id; |
25 | } args_t; |
26 | |
27 | // CHECK-DAG: Thread: 0, Block: 0 |
28 | // CHECK-DAG: Thread: 1, Block: 0 |
29 | // CHECK-DAG: Thread: 0, Block: 1 |
30 | // CHECK-DAG: Thread: 1, Block: 1 |
31 | // CHECK-DAG: Thread: 0, Block: 2 |
32 | // CHECK-DAG: Thread: 1, Block: 2 |
33 | // CHECK-DAG: Thread: 0, Block: 3 |
34 | // CHECK-DAG: Thread: 1, Block: 3 |
35 | void foo(void *data) { |
36 | assert(omp_is_initial_device() && "Not executing on host?" ); |
37 | args_t *args = (args_t *)data; |
38 | printf(format: "Thread: %d, Block: %d\n" , args->thread_id, args->block_id); |
39 | } |
40 | |
41 | void *fn_ptr = NULL; |
42 | #pragma omp declare target to(fn_ptr) |
43 | |
44 | int main() { |
45 | fn_ptr = (void *)&foo; |
46 | #pragma omp target update to(fn_ptr) |
47 | |
48 | #pragma omp target teams num_teams(4) |
49 | #pragma omp parallel num_threads(2) |
50 | { |
51 | args_t args = {omp_get_thread_num(), omp_get_team_num()}; |
52 | rpc_host_call(fn: fn_ptr, args: &args, size: sizeof(args_t)); |
53 | } |
54 | } |
55 | |