63 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2 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 variant12 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 variant19 20typedef struct args_s {21 int thread_id;22 int block_id;23} args_t;24 25// CHECK-DAG: Thread: 0, Block: 026// CHECK-DAG: Thread: 1, Block: 027// CHECK-DAG: Thread: 0, Block: 128// CHECK-DAG: Thread: 1, Block: 129// CHECK-DAG: Thread: 0, Block: 230// CHECK-DAG: Thread: 1, Block: 231// CHECK-DAG: Thread: 0, Block: 332// CHECK-DAG: Thread: 1, Block: 333unsigned 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 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_ptr, &args, sizeof(args_t));54 if (res != 42)55#pragma omp atomic write56 failed = 1;57 }58 59 // CHECK: PASS60 if (!failed)61 printf("PASS\n");62}63