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