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("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_ptr, &args, 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("PASS\n"); 62 } 63