108 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2#include <assert.h>3#include <omp.h>4#include <stdio.h>5 6// ---------------------------------------------------------------------------7// Various definitions copied from OpenMP RTL8 9typedef struct {10 uint64_t Reserved;11 uint16_t Version;12 uint16_t Kind; // OpenMP==113 uint32_t Flags;14 void *Address;15 char *SymbolName;16 uint64_t Size;17 uint64_t Data;18 void *AuxAddr;19} __tgt_offload_entry;20 21enum OpenMPOffloadingDeclareTargetFlags {22 /// Mark the entry global as having a 'link' attribute.23 OMP_DECLARE_TARGET_LINK = 0x01,24 /// Mark the entry global as being an indirectly callable function.25 OMP_DECLARE_TARGET_INDIRECT = 0x08,26 /// This is an entry corresponding to a requirement to be registered.27 OMP_REGISTER_REQUIRES = 0x10,28 /// Mark the entry global as being an indirect vtable.29 OMP_DECLARE_TARGET_INDIRECT_VTABLE = 0x20,30};31 32#pragma omp begin declare variant match(device = {kind(gpu)})33// Provided by the runtime.34void *__llvm_omp_indirect_call_lookup(void *host_ptr);35#pragma omp declare target to(__llvm_omp_indirect_call_lookup) \36 device_type(nohost)37#pragma omp end declare variant38 39#pragma omp begin declare variant match(device = {kind(cpu)})40// We assume unified addressing on the CPU target.41void *__llvm_omp_indirect_call_lookup(void *host_ptr) { return host_ptr; }42#pragma omp end declare variant43 44#pragma omp begin declare target45void foo(int *i) { *i += 1; }46void bar(int *i) { *i += 10; }47void baz(int *i) { *i += 100; }48#pragma omp end declare target49 50typedef void (*fptr_t)(int *i);51 52// Dispatch Table - declare separately on host and device to avoid53// registering with the library; this also allows us to use separate54// names, which is convenient for debugging. This dispatchTable is55// intended to mimic what Clang emits for C++ vtables.56fptr_t dispatchTable[] = {foo, bar, baz};57#pragma omp begin declare target device_type(nohost)58fptr_t GPUdispatchTable[] = {foo, bar, baz};59fptr_t *GPUdispatchTablePtr = GPUdispatchTable;60#pragma omp end declare target61 62// Define "manual" OpenMP offload entries, where we emit Clang63// offloading entry structure definitions in the appropriate ELF64// section. This allows us to emulate the offloading entries that Clang would65// normally emit for us66 67__attribute__((weak, section("llvm_offload_entries"), aligned(8)))68const __tgt_offload_entry __offloading_entry[] = {{69 0ULL, // Reserved70 1, // Version71 1, // Kind72 OMP_DECLARE_TARGET_INDIRECT_VTABLE, // Flags73 &dispatchTable, // Address74 "GPUdispatchTablePtr", // SymbolName75 (size_t)(sizeof(dispatchTable)), // Size76 0ULL, // Data77 NULL // AuxAddr78}};79 80// Mimic how Clang emits vtable pointers for C++ classes81typedef struct {82 fptr_t *dispatchPtr;83} myClass;84 85// ---------------------------------------------------------------------------86int main() {87 myClass obj_foo = {dispatchTable + 0};88 myClass obj_bar = {dispatchTable + 1};89 myClass obj_baz = {dispatchTable + 2};90 int aaa = 0;91 92#pragma omp target map(aaa) map(to : obj_foo, obj_bar, obj_baz)93 {94 // Lookup95 fptr_t *foo_ptr = __llvm_omp_indirect_call_lookup(obj_foo.dispatchPtr);96 fptr_t *bar_ptr = __llvm_omp_indirect_call_lookup(obj_bar.dispatchPtr);97 fptr_t *baz_ptr = __llvm_omp_indirect_call_lookup(obj_baz.dispatchPtr);98 foo_ptr[0](&aaa);99 bar_ptr[0](&aaa);100 baz_ptr[0](&aaa);101 }102 103 assert(aaa == 111);104 // CHECK: PASS105 printf("PASS\n");106 return 0;107}108