diff options
Diffstat (limited to 'offload/test')
19 files changed, 644 insertions, 8 deletions
diff --git a/offload/test/api/omp_device_uid.c b/offload/test/api/omp_device_uid.c new file mode 100644 index 0000000..2a41d8d --- /dev/null +++ b/offload/test/api/omp_device_uid.c @@ -0,0 +1,76 @@ +// RUN: %libomptarget-compile-run-and-check-generic + +#include <omp.h> +#include <stdio.h> +#include <string.h> + +int test_omp_device_uid(int device_num) { + const char *device_uid = omp_get_uid_from_device(device_num); + if (device_uid == NULL) { + printf("FAIL for device %d: omp_get_uid_from_device returned NULL\n", + device_num); + return 0; + } + + int device_num_from_uid = omp_get_device_from_uid(device_uid); + if (device_num_from_uid != device_num) { + printf( + "FAIL for device %d: omp_get_device_from_uid returned %d (UID: %s)\n", + device_num, device_num_from_uid, device_uid); + return 0; + } + + if (device_num == omp_get_initial_device()) + return 1; + + int success = 1; + +// Note that the following code may be executed on the host if the host is the +// device +#pragma omp target map(tofrom : success) device(device_num) + { + int device_num = omp_get_device_num(); + + // omp_get_uid_from_device() in the device runtime is a dummy function + // returning NULL + const char *device_uid = omp_get_uid_from_device(device_num); + + // omp_get_device_from_uid() in the device runtime is a dummy function + // returning omp_invalid_device. + int device_num_from_uid = omp_get_device_from_uid(device_uid); + + // Depending on whether we're executing on the device or the host, we either + // got NULL as the device UID or the correct device UID. Consequently, + // omp_get_device_from_uid() either returned omp_invalid_device or the + // correct device number (aka omp_get_initial_device()). + if (device_uid ? device_num_from_uid != device_num + : device_num_from_uid != omp_invalid_device) { + printf("FAIL for device %d (target): omp_get_device_from_uid returned %d " + "(UID: %s)\n", + device_num, device_num_from_uid, device_uid); + success = 0; + } + } + + return success; +} + +int main() { + int num_devices = omp_get_num_devices(); + int num_failed = 0; + // (also test initial device aka num_devices) + for (int i = 0; i < num_devices + 1; i++) { + if (!test_omp_device_uid(i)) { + printf("FAIL for device %d\n", i); + num_failed++; + } + } + if (num_failed) { + printf("FAIL\n"); + return 1; + } + printf("PASS\n"); + return 0; +} + +// CHECK: PASS diff --git a/offload/test/api/omp_indirect_call_table_manual.c b/offload/test/api/omp_indirect_call_table_manual.c new file mode 100644 index 0000000..e958d47 --- /dev/null +++ b/offload/test/api/omp_indirect_call_table_manual.c @@ -0,0 +1,107 @@ +// RUN: %libomptarget-compile-run-and-check-generic +#include <assert.h> +#include <omp.h> +#include <stdio.h> + +// --------------------------------------------------------------------------- +// Various definitions copied from OpenMP RTL + +typedef struct { + uint64_t Reserved; + uint16_t Version; + uint16_t Kind; // OpenMP==1 + uint32_t Flags; + void *Address; + char *SymbolName; + uint64_t Size; + uint64_t Data; + void *AuxAddr; +} __tgt_offload_entry; + +enum OpenMPOffloadingDeclareTargetFlags { + /// Mark the entry global as having a 'link' attribute. + OMP_DECLARE_TARGET_LINK = 0x01, + /// Mark the entry global as being an indirectly callable function. + OMP_DECLARE_TARGET_INDIRECT = 0x08, + /// This is an entry corresponding to a requirement to be registered. + OMP_REGISTER_REQUIRES = 0x10, + /// Mark the entry global as being an indirect vtable. + OMP_DECLARE_TARGET_INDIRECT_VTABLE = 0x20, +}; + +#pragma omp begin declare variant match(device = {kind(gpu)}) +// Provided by the runtime. +void *__llvm_omp_indirect_call_lookup(void *host_ptr); +#pragma omp declare target to(__llvm_omp_indirect_call_lookup) \ + device_type(nohost) +#pragma omp end declare variant + +#pragma omp begin declare variant match(device = {kind(cpu)}) +// We assume unified addressing on the CPU target. +void *__llvm_omp_indirect_call_lookup(void *host_ptr) { return host_ptr; } +#pragma omp end declare variant + +#pragma omp begin declare target +void foo(int *i) { *i += 1; } +void bar(int *i) { *i += 10; } +void baz(int *i) { *i += 100; } +#pragma omp end declare target + +typedef void (*fptr_t)(int *i); + +// Dispatch Table - declare separately on host and device to avoid +// registering with the library; this also allows us to use separate +// names, which is convenient for debugging. This dispatchTable is +// intended to mimic what Clang emits for C++ vtables. +fptr_t dispatchTable[] = {foo, bar, baz}; +#pragma omp begin declare target device_type(nohost) +fptr_t GPUdispatchTable[] = {foo, bar, baz}; +fptr_t *GPUdispatchTablePtr = GPUdispatchTable; +#pragma omp end declare target + +// Define "manual" OpenMP offload entries, where we emit Clang +// offloading entry structure definitions in the appropriate ELF +// section. This allows us to emulate the offloading entries that Clang would +// normally emit for us + +__attribute__((weak, section("llvm_offload_entries"), aligned(8))) +const __tgt_offload_entry __offloading_entry[] = {{ + 0ULL, // Reserved + 1, // Version + 1, // Kind + OMP_DECLARE_TARGET_INDIRECT_VTABLE, // Flags + &dispatchTable, // Address + "GPUdispatchTablePtr", // SymbolName + (size_t)(sizeof(dispatchTable)), // Size + 0ULL, // Data + NULL // AuxAddr +}}; + +// Mimic how Clang emits vtable pointers for C++ classes +typedef struct { + fptr_t *dispatchPtr; +} myClass; + +// --------------------------------------------------------------------------- +int main() { + myClass obj_foo = {dispatchTable + 0}; + myClass obj_bar = {dispatchTable + 1}; + myClass obj_baz = {dispatchTable + 2}; + int aaa = 0; + +#pragma omp target map(aaa) map(to : obj_foo, obj_bar, obj_baz) + { + // Lookup + fptr_t *foo_ptr = __llvm_omp_indirect_call_lookup(obj_foo.dispatchPtr); + fptr_t *bar_ptr = __llvm_omp_indirect_call_lookup(obj_bar.dispatchPtr); + fptr_t *baz_ptr = __llvm_omp_indirect_call_lookup(obj_baz.dispatchPtr); + foo_ptr[0](&aaa); + bar_ptr[0](&aaa); + baz_ptr[0](&aaa); + } + + assert(aaa == 111); + // CHECK: PASS + printf("PASS\n"); + return 0; +} diff --git a/offload/test/lit.site.cfg.in b/offload/test/lit.site.cfg.in index 00f4e2b..c8ba45c 100644 --- a/offload/test/lit.site.cfg.in +++ b/offload/test/lit.site.cfg.in @@ -1,6 +1,6 @@ @AUTO_GEN_COMMENT@ -config.bin_llvm_tools_dir = "@LLVM_RUNTIME_OUTPUT_INTDIR@" +config.bin_llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.test_c_compiler = "@OPENMP_TEST_C_COMPILER@" config.test_cxx_compiler = "@OPENMP_TEST_CXX_COMPILER@" config.test_fortran_compiler="@OPENMP_TEST_Fortran_COMPILER@" diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c new file mode 100644 index 0000000..4b67a3b --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c @@ -0,0 +1,24 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// Test that when a use_device_addr lookup fails, the +// list-item retains its original address by default. +// +// This is necessary because we must assume that the +// list-item is device-accessible, even if it was not +// previously mapped. + +// XFAIL: * + +#include <stdio.h> +int h[10]; +int *ph = &h[0]; + +void f1() { + printf("%p\n", &h[2]); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_addr(h[2]) + printf("%p\n", &h[2]); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +#pragma omp target data use_device_addr(ph[2]) + printf("%p\n", &ph[2]); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c new file mode 100644 index 0000000..4495a46b --- /dev/null +++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c @@ -0,0 +1,21 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// Test that when a use_device_addr lookup fails, the +// list-item retains its original address by default. +// +// This is necessary because we must assume that the +// list-item is device-accessible, even if it was not +// previously mapped. + +// XFAIL: * + +#include <stdio.h> +int x; + +void f1() { + printf("%p\n", &x); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_addr(x) + printf("%p\n", &x); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c new file mode 100644 index 0000000..e8fa3b6 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c @@ -0,0 +1,32 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value by +// default. +// +// This is necessary because we must assume that the +// pointee is device-accessible, even if it was not +// previously mapped. +// +// OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31: +// If a list item that appears in a use_device_ptr clause ... does not point to +// a mapped object, it must contain a valid device address for the target +// device, and the list item references are instead converted to references to a +// local device pointer that refers to this device address. +// +// Note: OpenMP 6.1 will have a way to change the +// fallback behavior: preserve or nullify. + +// XFAIL: * + +#include <stdio.h> +int x; +int *xp = &x; + +void f1() { + printf("%p\n", xp); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(xp) + printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/offloading/fortran/declare-target-to-allocatable-vars-in-target-with-update.f90 b/offload/test/offloading/fortran/declare-target-to-allocatable-vars-in-target-with-update.f90 new file mode 100644 index 0000000..727a08b --- /dev/null +++ b/offload/test/offloading/fortran/declare-target-to-allocatable-vars-in-target-with-update.f90 @@ -0,0 +1,41 @@ +! Test that checks an allocatable array can be marked implicit +! `declare target to` and functions without issue. +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic +module test + implicit none + integer, allocatable, dimension(:) :: alloca_arr + !$omp declare target(alloca_arr) +end module test + +program main + use test + implicit none + integer :: cycle, i + + allocate(alloca_arr(10)) + + do i = 1, 10 + alloca_arr(i) = 0 + end do + + !$omp target data map(to:alloca_arr) + do cycle = 1, 2 + !$omp target + do i = 1, 10 + alloca_arr(i) = alloca_arr(i) + i + end do + !$omp end target + + ! NOTE: Technically doesn't affect the results, but there is a + ! regression case that'll cause a runtime crash if this is + ! invoked more than once, so this checks for that. + !$omp target update from(alloca_arr) + end do + !$omp end target data + + print *, alloca_arr +end program + +! CHECK: 2 4 6 8 10 12 14 16 18 20 diff --git a/offload/test/offloading/fortran/declare-target-to-vars-target-region-and-update.f90 b/offload/test/offloading/fortran/declare-target-to-vars-target-region-and-update.f90 new file mode 100644 index 0000000..16433af --- /dev/null +++ b/offload/test/offloading/fortran/declare-target-to-vars-target-region-and-update.f90 @@ -0,0 +1,40 @@ +! Test the implicit `declare target to` interaction with `target update from` +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic +module test + implicit none + integer :: array(10) + !$omp declare target(array) +end module test + +PROGRAM main + use test + implicit none + integer :: i + + do i = 1, 10 + array(i) = 0 + end do + + !$omp target + do i = 1, 10 + array(i) = i + end do + !$omp end target + + !$omp target + do i = 1, 10 + array(i) = array(i) + i + end do + !$omp end target + + print *, array + + !$omp target update from(array) + + print *, array +END PROGRAM + +! CHECK: 0 0 0 0 0 0 0 0 0 0 +! CHECK: 2 4 6 8 10 12 14 16 18 20 diff --git a/offload/test/offloading/fortran/declare-target-to-zero-index-allocatable-target-map.f90 b/offload/test/offloading/fortran/declare-target-to-zero-index-allocatable-target-map.f90 new file mode 100644 index 0000000..0d650f6 --- /dev/null +++ b/offload/test/offloading/fortran/declare-target-to-zero-index-allocatable-target-map.f90 @@ -0,0 +1,30 @@ +! Test `declare target to` interaction with an allocatable with a non-default +! range +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic +module test_0 + real(4), allocatable :: zero_off(:) + !$omp declare target(zero_off) +end module test_0 + +program main + use test_0 + implicit none + + allocate(zero_off(0:10)) + + zero_off(0) = 30.0 + zero_off(1) = 40.0 + zero_off(10) = 25.0 + + !$omp target map(tofrom: zero_off) + zero_off(0) = zero_off(1) + !$omp end target + + print *, zero_off(0) + print *, zero_off(1) +end program + +! CHECK: 40. +! CHECK: 40. diff --git a/offload/test/offloading/fortran/dtype-member-overlap-map.f90 b/offload/test/offloading/fortran/dtype-member-overlap-map.f90 new file mode 100644 index 0000000..e457014 --- /dev/null +++ b/offload/test/offloading/fortran/dtype-member-overlap-map.f90 @@ -0,0 +1,56 @@ +! Basic offloading test checking the interaction of an overlapping +! member map. +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic +program main + implicit none + integer :: i + + type dtype2 + integer :: int + real :: float + end type dtype2 + + type dtype1 + character (LEN=30) :: characters + type(dtype2) :: internal_dtype2 + end type dtype1 + + type dtype + integer :: elements(10) + type(dtype1) :: internal_dtype + integer :: value + end type dtype + + type (dtype) :: single_dtype + + do i = 1, 10 + single_dtype%elements(i) = 0 + end do + + !$omp target map(tofrom: single_dtype%internal_dtype, single_dtype%internal_dtype%internal_dtype2%int) + single_dtype%internal_dtype%internal_dtype2%int = 123 + single_dtype%internal_dtype%characters(1:1) = "Z" + !$omp end target + + !$omp target map(to: single_dtype) map(tofrom: single_dtype%internal_dtype%internal_dtype2, single_dtype%value) + single_dtype%value = 20 + do i = 1, 10 + single_dtype%elements(i) = i + end do + single_dtype%internal_dtype%internal_dtype2%float = 32.0 + !$omp end target + + print *, single_dtype%value + print *, single_dtype%internal_dtype%internal_dtype2%float + print *, single_dtype%elements + print *, single_dtype%internal_dtype%internal_dtype2%int + print *, single_dtype%internal_dtype%characters(1:1) +end program main + +! CHECK: 20 +! CHECK: 32. +! CHECK: 0 0 0 0 0 0 0 0 0 0 +! CHECK: 123 +! CHECK: Z diff --git a/offload/test/offloading/fortran/implicit-derived-enter-exit.f90 b/offload/test/offloading/fortran/implicit-derived-enter-exit.f90 new file mode 100644 index 0000000..0c896e64 --- /dev/null +++ b/offload/test/offloading/fortran/implicit-derived-enter-exit.f90 @@ -0,0 +1,65 @@ +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-generic +! RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic + +module enter_exit_mapper_mod + implicit none + + type :: field_type + real, allocatable :: values(:) + end type field_type + + type :: tile_type + type(field_type) :: field + integer, allocatable :: neighbors(:) + end type tile_type + +contains + subroutine init_tile(tile) + type(tile_type), intent(inout) :: tile + integer :: j + + allocate(tile%field%values(4)) + allocate(tile%neighbors(4)) + do j = 1, 4 + tile%field%values(j) = 10.0 * j + tile%neighbors(j) = j + end do + end subroutine init_tile + +end module enter_exit_mapper_mod + +program implicit_enter_exit + use enter_exit_mapper_mod + implicit none + integer :: j + type(tile_type) :: tile + + call init_tile(tile) + + !$omp target enter data map(alloc: tile%field%values) + + !$omp target + do j = 1, size(tile%field%values) + tile%field%values(j) = 5.0 * j + end do + !$omp end target + + !$omp target exit data map(from: tile%field%values) + + do j = 1, size(tile%field%values) + if (tile%field%values(j) /= 5.0 * j) then + print *, "======= Test Failed! =======" + stop 1 + end if + if (tile%neighbors(j) /= j) then + print *, "======= Test Failed! =======" + stop 1 + end if + end do + + print *, "======= Test Passed! =======" +end program implicit_enter_exit + +! CHECK: ======= Test Passed! ======= diff --git a/offload/test/offloading/fortran/target-custom-reduction-derivedtype.f90 b/offload/test/offloading/fortran/target-custom-reduction-derivedtype.f90 new file mode 100644 index 0000000..cc390cf0 --- /dev/null +++ b/offload/test/offloading/fortran/target-custom-reduction-derivedtype.f90 @@ -0,0 +1,88 @@ +! Basic offloading test with custom OpenMP reduction on derived type +! REQUIRES: flang, amdgpu +! +! RUN: %libomptarget-compile-fortran-generic +! RUN: env LIBOMPTARGET_INFO=16 %libomptarget-run-generic 2>&1 | %fcheck-generic +module maxtype_mod + implicit none + + type maxtype + integer::sumval + integer::maxval + end type maxtype + +contains + + subroutine initme(x,n) + type(maxtype) :: x,n + x%sumval=0 + x%maxval=0 + end subroutine initme + + function mycombine(lhs, rhs) + type(maxtype) :: lhs, rhs + type(maxtype) :: mycombine + mycombine%sumval = lhs%sumval + rhs%sumval + mycombine%maxval = max(lhs%maxval, rhs%maxval) + end function mycombine + +end module maxtype_mod + +program main + use maxtype_mod + implicit none + + integer :: n = 100 + integer :: i + integer :: error = 0 + type(maxtype) :: x(100) + type(maxtype) :: res + integer :: expected_sum, expected_max + +!$omp declare reduction(red_add_max:maxtype:omp_out=mycombine(omp_out,omp_in)) initializer(initme(omp_priv,omp_orig)) + + ! Initialize array with test data + do i = 1, n + x(i)%sumval = i + x(i)%maxval = i + end do + + ! Initialize reduction variable + res%sumval = 0 + res%maxval = 0 + + ! Perform reduction in target region + !$omp target parallel do map(to:x) reduction(red_add_max:res) + do i = 1, n + res = mycombine(res, x(i)) + end do + !$omp end target parallel do + + ! Compute expected values + expected_sum = 0 + expected_max = 0 + do i = 1, n + expected_sum = expected_sum + i + expected_max = max(expected_max, i) + end do + + ! Check results + if (res%sumval /= expected_sum) then + error = 1 + endif + + if (res%maxval /= expected_max) then + error = 1 + endif + + if (error == 0) then + print *,"PASSED" + else + print *,"FAILED" + endif + +end program main + +! CHECK: "PluginInterface" device {{[0-9]+}} info: Launching kernel {{.*}} +! CHECK: PASSED + diff --git a/offload/test/offloading/fortran/target-is-device-ptr.f90 b/offload/test/offloading/fortran/target-is-device-ptr.f90 new file mode 100644 index 0000000..d6d8c02 --- /dev/null +++ b/offload/test/offloading/fortran/target-is-device-ptr.f90 @@ -0,0 +1,60 @@ +! Validate that a device pointer obtained via omp_get_mapped_ptr can be used +! inside a TARGET region with the is_device_ptr clause. +! REQUIRES: flang, amdgcn-amd-amdhsa + +! RUN: %libomptarget-compile-fortran-run-and-check-generic + +module mod + implicit none + integer, parameter :: n = 4 +contains + subroutine kernel(dptr) + use iso_c_binding, only : c_ptr, c_f_pointer + implicit none + + type(c_ptr) :: dptr + integer, dimension(:), pointer :: b + integer :: i + + b => null() + + !$omp target is_device_ptr(dptr) + call c_f_pointer(dptr, b, [n]) + do i = 1, n + b(i) = b(i) + 1 + end do + !$omp end target + end subroutine kernel +end module mod + +program is_device_ptr_target + use iso_c_binding, only : c_ptr, c_loc, c_f_pointer + use omp_lib, only: omp_get_default_device, omp_get_mapped_ptr + use mod, only: kernel, n + implicit none + + integer, dimension(n), target :: a + integer :: dev + type(c_ptr) :: dptr + + a = [2, 4, 6, 8] + print '("BEFORE:", I3)', a + + dev = omp_get_default_device() + + !$omp target data map(tofrom: a) + dptr = omp_get_mapped_ptr(c_loc(a), dev) + call kernel(dptr) + !$omp end target data + + print '("AFTER: ", I3)', a + + if (all(a == [3, 5, 7, 9])) then + print '("PASS")' + else + print '("FAIL ", I3)', a + end if + +end program is_device_ptr_target + +!CHECK: PASS diff --git a/offload/test/offloading/gpupgo/pgo_atomic_teams.c b/offload/test/offloading/gpupgo/pgo_atomic_teams.c index 42d8ae4..b3b72db 100644 --- a/offload/test/offloading/gpupgo/pgo_atomic_teams.c +++ b/offload/test/offloading/gpupgo/pgo_atomic_teams.c @@ -18,7 +18,6 @@ // REQUIRES: amdgpu // REQUIRES: pgo -// XFAIL: amdgpu int test1(int a) { return a / 2; } int test2(int a) { return a * 2; } diff --git a/offload/test/offloading/gpupgo/pgo_atomic_threads.c b/offload/test/offloading/gpupgo/pgo_atomic_threads.c index 09a4dc1..440a6b5 100644 --- a/offload/test/offloading/gpupgo/pgo_atomic_threads.c +++ b/offload/test/offloading/gpupgo/pgo_atomic_threads.c @@ -18,7 +18,6 @@ // REQUIRES: amdgpu // REQUIRES: pgo -// XFAIL: amdgpu int test1(int a) { return a / 2; } diff --git a/offload/test/offloading/gpupgo/pgo_device_and_host.c b/offload/test/offloading/gpupgo/pgo_device_and_host.c index c53e69a..3e95791 100644 --- a/offload/test/offloading/gpupgo/pgo_device_and_host.c +++ b/offload/test/offloading/gpupgo/pgo_device_and_host.c @@ -50,7 +50,6 @@ // REQUIRES: amdgpu // REQUIRES: pgo -// XFAIL: amdgpu int main() { int host_var = 0; diff --git a/offload/test/offloading/gpupgo/pgo_device_only.c b/offload/test/offloading/gpupgo/pgo_device_only.c index 644df6e..2939af61 100644 --- a/offload/test/offloading/gpupgo/pgo_device_only.c +++ b/offload/test/offloading/gpupgo/pgo_device_only.c @@ -16,7 +16,6 @@ // REQUIRES: amdgpu // REQUIRES: pgo -// XFAIL: amdgpu int test1(int a) { return a / 2; } int test2(int a) { return a * 2; } diff --git a/offload/test/offloading/shared_lib_fp_mapping.c b/offload/test/offloading/shared_lib_fp_mapping.c index c620344..e0af9b7 100644 --- a/offload/test/offloading/shared_lib_fp_mapping.c +++ b/offload/test/offloading/shared_lib_fp_mapping.c @@ -7,8 +7,8 @@ #include <stdio.h> -extern int func(); // Provided in liba.so, returns 42 -typedef int (*fp_t)(); +extern int func(void); // Provided in liba.so, returns 42 +typedef int (*fp_t)(void); int main() { int x = 0; diff --git a/offload/test/offloading/static_linking.c b/offload/test/offloading/static_linking.c index 7be95a1..273109e 100644 --- a/offload/test/offloading/static_linking.c +++ b/offload/test/offloading/static_linking.c @@ -14,7 +14,7 @@ int foo() { } #else #include <stdio.h> -int foo(); +int foo(void); int main() { int x = foo(); |