57 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2 3// REQUIRES: libc4 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.11unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size);12#pragma omp declare target to(__llvm_omp_host_call) device_type(nohost)13#pragma omp end declare variant14 15#pragma omp begin declare variant match(device = {kind(cpu)})16// Dummy host implementation to make this work for all targets.17unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size) {18 return ((unsigned long long (*)(void *))fn)(args);19}20#pragma omp end declare variant21 22long long foo(void *data) { return -1; }23 24void *fn_ptr = NULL;25#pragma omp declare target to(fn_ptr)26 27int main() {28 fn_ptr = (void *)&foo;29#pragma omp target update to(fn_ptr)30 31 for (int i = 0; i < 4; ++i) {32#pragma omp target33 {34 long long res = __llvm_omp_host_call(fn_ptr, NULL, 0);35 assert(res == -1 && "RPC call failed\n");36 }37 38 for (int j = 0; j < 128; ++j) {39#pragma omp target nowait40 {41 long long res = __llvm_omp_host_call(fn_ptr, NULL, 0);42 assert(res == -1 && "RPC call failed\n");43 }44 }45#pragma omp taskwait46 47#pragma omp target48 {49 long long res = __llvm_omp_host_call(fn_ptr, NULL, 0);50 assert(res == -1 && "RPC call failed\n");51 }52 }53 54 // CHECK: PASS55 puts("PASS");56}57