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.
9unsigned 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.
15unsigned 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
20typedef 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
33unsigned 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
40void *fn_ptr = NULL;
41#pragma omp declare target to(fn_ptr)
42
43int 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

source code of offload/test/api/omp_host_call.c