5aea02a163
`__llvm_omp_indirect_call_lookup` takes in and returns a function pointer, so make sure the types are correct, which includes the correct address space. The FE was recently changed to generate the correct code [here](https://github.com/llvm/llvm-project/pull/192470). With this change, three function pointer tests start passing. Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
322 lines
11 KiB
C
322 lines
11 KiB
C
// RUN: %libomptarget-compile-run-and-check-generic
|
|
// REQUIRES: gpu
|
|
|
|
#include <assert.h>
|
|
#include <omp.h>
|
|
#include <stdio.h>
|
|
|
|
#define TEST_VAL 5
|
|
|
|
#pragma omp declare target indirect
|
|
int direct_arg(int x) { return 2 * x; }
|
|
int indirect_base_arg(int x) { return -1 * x; }
|
|
int direct() { return TEST_VAL; }
|
|
int indirect_base() { return -1 * TEST_VAL; }
|
|
#pragma omp end declare target
|
|
|
|
struct indirect_stru {
|
|
int buffer;
|
|
int (*indirect1)();
|
|
int (*indirect0)(int);
|
|
};
|
|
typedef struct {
|
|
int buffer;
|
|
int (*indirect1_ptr)();
|
|
int (*indirect0_ptr)(int);
|
|
} indirect_stru_mapped;
|
|
|
|
#pragma omp declare mapper(indirect_stru_mapped s) \
|
|
map(s, s.indirect0_ptr, s.indirect1_ptr)
|
|
|
|
struct nested_stru {
|
|
struct indirect_stru inner;
|
|
};
|
|
|
|
struct indirect_stru global_indirect_val = {.indirect0 = indirect_base_arg,
|
|
.indirect1 = indirect_base};
|
|
indirect_stru_mapped global_mapped_val = {.indirect0_ptr = indirect_base_arg,
|
|
.indirect1_ptr = indirect_base};
|
|
struct nested_stru global_nested_val = {
|
|
.inner = {.buffer = 0,
|
|
.indirect0 = indirect_base_arg,
|
|
.indirect1 = indirect_base}};
|
|
|
|
void test_nested_struct() {
|
|
int indirect0_ret, indirect1_ret;
|
|
int indirect0_base = indirect_base_arg(TEST_VAL);
|
|
int indirect1_base = indirect_base();
|
|
|
|
#pragma omp target map(global_nested_val, global_nested_val.inner.indirect0, \
|
|
global_nested_val.inner.indirect1) \
|
|
map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = global_nested_val.inner.indirect0(TEST_VAL);
|
|
indirect1_ret = global_nested_val.inner.indirect1();
|
|
}
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 returned incorrect value on device");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 returned incorrect value on device");
|
|
|
|
// Change function pointers and re-test
|
|
global_nested_val.inner.indirect0 = direct_arg;
|
|
global_nested_val.inner.indirect1 = direct;
|
|
|
|
indirect0_base = direct_arg(TEST_VAL);
|
|
indirect1_base = direct();
|
|
|
|
#pragma omp target map(global_nested_val, global_nested_val.inner.indirect0, \
|
|
global_nested_val.inner.indirect1) \
|
|
map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = global_nested_val.inner.indirect0(TEST_VAL);
|
|
indirect1_ret = global_nested_val.inner.indirect1();
|
|
}
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 returned incorrect value on device after update");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 returned incorrect value on device after update");
|
|
}
|
|
|
|
void test_global_struct_explicit_mapping() {
|
|
int indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
|
|
int indirect0_base = indirect_base_arg(TEST_VAL);
|
|
|
|
int indirect1_ret = global_indirect_val.indirect1();
|
|
int indirect1_base = indirect_base();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(global_indirect_val, global_indirect_val.indirect1, \
|
|
global_indirect_val.indirect0) \
|
|
map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
|
|
indirect1_ret = global_indirect_val.indirect1();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
|
|
global_indirect_val.indirect0 = direct_arg;
|
|
global_indirect_val.indirect1 = direct;
|
|
|
|
indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
|
|
indirect0_base = direct_arg(TEST_VAL);
|
|
|
|
indirect1_ret = global_indirect_val.indirect1();
|
|
indirect1_base = direct();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(global_indirect_val, global_indirect_val.indirect0, \
|
|
global_indirect_val.indirect1) \
|
|
map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = global_indirect_val.indirect0(TEST_VAL);
|
|
indirect1_ret = global_indirect_val.indirect1();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
}
|
|
|
|
void test_local_struct_explicit_mapping() {
|
|
struct indirect_stru local_indirect_val;
|
|
local_indirect_val.indirect0 = indirect_base_arg;
|
|
local_indirect_val.indirect1 = indirect_base;
|
|
|
|
int indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
|
|
int indirect0_base = indirect_base_arg(TEST_VAL);
|
|
|
|
int indirect1_ret = local_indirect_val.indirect1();
|
|
int indirect1_base = indirect_base();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(local_indirect_val, local_indirect_val.indirect1, \
|
|
local_indirect_val.indirect0) \
|
|
map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
|
|
indirect1_ret = local_indirect_val.indirect1();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
|
|
local_indirect_val.indirect0 = direct_arg;
|
|
local_indirect_val.indirect1 = direct;
|
|
|
|
indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
|
|
indirect0_base = direct_arg(TEST_VAL);
|
|
|
|
indirect1_ret = local_indirect_val.indirect1();
|
|
indirect1_base = direct();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(local_indirect_val, local_indirect_val.indirect0, \
|
|
local_indirect_val.indirect1) \
|
|
map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = local_indirect_val.indirect0(TEST_VAL);
|
|
indirect1_ret = local_indirect_val.indirect1();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
}
|
|
|
|
void test_global_struct_user_mapper() {
|
|
int indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
|
|
int indirect0_base = indirect_base_arg(TEST_VAL);
|
|
|
|
int indirect1_ret = global_mapped_val.indirect1_ptr();
|
|
int indirect1_base = indirect_base();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
|
|
indirect1_ret = global_mapped_val.indirect1_ptr();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
|
|
global_mapped_val.indirect0_ptr = direct_arg;
|
|
global_mapped_val.indirect1_ptr = direct;
|
|
|
|
indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
|
|
indirect0_base = direct_arg(TEST_VAL);
|
|
|
|
indirect1_ret = global_mapped_val.indirect1_ptr();
|
|
indirect1_base = direct();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = global_mapped_val.indirect0_ptr(TEST_VAL);
|
|
indirect1_ret = global_mapped_val.indirect1_ptr();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
}
|
|
|
|
void test_local_struct_user_mapper() {
|
|
indirect_stru_mapped local_mapped_val;
|
|
local_mapped_val.indirect0_ptr = indirect_base_arg;
|
|
local_mapped_val.indirect1_ptr = indirect_base;
|
|
|
|
int indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
|
|
int indirect0_base = indirect_base_arg(TEST_VAL);
|
|
|
|
int indirect1_ret = local_mapped_val.indirect1_ptr();
|
|
int indirect1_base = indirect_base();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
|
|
indirect1_ret = local_mapped_val.indirect1_ptr();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
|
|
local_mapped_val.indirect0_ptr = direct_arg;
|
|
local_mapped_val.indirect1_ptr = direct;
|
|
|
|
indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
|
|
indirect0_base = direct_arg(TEST_VAL);
|
|
|
|
indirect1_ret = local_mapped_val.indirect1_ptr();
|
|
indirect1_base = direct();
|
|
|
|
assert(indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on host");
|
|
assert(indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on host");
|
|
|
|
#pragma omp target map(from : indirect0_ret, indirect1_ret)
|
|
{
|
|
indirect0_ret = local_mapped_val.indirect0_ptr(TEST_VAL);
|
|
indirect1_ret = local_mapped_val.indirect1_ptr();
|
|
}
|
|
|
|
assert(
|
|
indirect0_ret == indirect0_base &&
|
|
"Error: indirect0 function pointer returned incorrect value on device");
|
|
assert(
|
|
indirect1_ret == indirect1_base &&
|
|
"Error: indirect1 function pointer returned incorrect value on device");
|
|
}
|
|
|
|
int main() {
|
|
test_nested_struct();
|
|
test_global_struct_explicit_mapping();
|
|
test_local_struct_explicit_mapping();
|
|
test_global_struct_user_mapper();
|
|
test_local_struct_user_mapper();
|
|
|
|
// CHECK: PASS
|
|
printf("PASS\n");
|
|
return 0;
|
|
}
|