diff options
Diffstat (limited to 'offload/test/offloading')
| -rw-r--r-- | offload/test/offloading/fortran/default-mapper-derived-enter-data-teams-collapse.f90 | 53 | ||||
| -rw-r--r-- | offload/test/offloading/fortran/recursive-default-mapper.f90 | 40 | ||||
| -rw-r--r-- | offload/test/offloading/fortran/target-defaultmap-implicit-mapper.f90 | 49 | ||||
| -rw-r--r-- | offload/test/offloading/fortran/target-parameter-array.f90 | 131 | ||||
| -rw-r--r-- | offload/test/offloading/strided_multiple_update_from.c (renamed from offload/test/offloading/strided_multiple_update.c) | 0 | ||||
| -rw-r--r-- | offload/test/offloading/strided_multiple_update_to.c | 124 | ||||
| -rw-r--r-- | offload/test/offloading/strided_partial_update_from.c (renamed from offload/test/offloading/strided_partial_update.c) | 0 | ||||
| -rw-r--r-- | offload/test/offloading/strided_partial_update_to.c | 74 | ||||
| -rw-r--r-- | offload/test/offloading/strided_update_from.c (renamed from offload/test/offloading/strided_update.c) | 0 | ||||
| -rw-r--r-- | offload/test/offloading/strided_update_to.c | 74 |
10 files changed, 545 insertions, 0 deletions
diff --git a/offload/test/offloading/fortran/default-mapper-derived-enter-data-teams-collapse.f90 b/offload/test/offloading/fortran/default-mapper-derived-enter-data-teams-collapse.f90 new file mode 100644 index 0000000..6b87e81 --- /dev/null +++ b/offload/test/offloading/fortran/default-mapper-derived-enter-data-teams-collapse.f90 @@ -0,0 +1,53 @@ +! Regression test for default mappers on nested derived types with allocatable +! members when mapping a parent object and running an optimized target region. + +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-generic -O3 +! RUN: %libomptarget-run-generic | %fcheck-generic + +program test_default_mapper_enter_data_teams_collapse + implicit none + + type inner_type + real, allocatable :: data(:) + end type inner_type + + type outer_type + type(inner_type) :: inner + character(len=19) :: desc = ' ' + end type outer_type + + type(outer_type) :: obj + integer, parameter :: n = 10 + integer :: i, j + real :: expected, actual + + allocate(obj%inner%data(n)) + obj%inner%data = 0.0 + + !$omp target enter data map(to: obj) + + !$omp target teams distribute parallel do collapse(2) + do i = 1, n + do j = 1, n + obj%inner%data(i) = real(i) + end do + end do + !$omp end target teams distribute parallel do + + !$omp target exit data map(from: obj) + + expected = real(n * (n + 1)) / 2.0 + actual = sum(obj%inner%data) + + if (abs(actual - expected) < 1.0e-6) then + print *, "PASS" + else + print *, "FAIL", actual, expected + end if + + deallocate(obj%inner%data) +end program test_default_mapper_enter_data_teams_collapse + +! CHECK: PASS diff --git a/offload/test/offloading/fortran/recursive-default-mapper.f90 b/offload/test/offloading/fortran/recursive-default-mapper.f90 new file mode 100644 index 0000000..47b706d --- /dev/null +++ b/offload/test/offloading/fortran/recursive-default-mapper.f90 @@ -0,0 +1,40 @@ +! Offloading test for recursive default mapper emission +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic + +module recursive_mapper_mod + implicit none + + type :: inner + integer :: value + type(inner), pointer :: next + end type inner + + type :: outer + integer, allocatable :: arr(:) + type(inner), pointer :: head + end type outer + +contains + +end module recursive_mapper_mod + +program main + use recursive_mapper_mod + implicit none + + type(outer) :: o + + allocate(o%arr(2)) + o%arr = [1, 2] + + !$omp target map(tofrom: o) + o%arr(1) = o%arr(1) + 1 + o%arr(2) = o%arr(2) + 1 + !$omp end target + + print *, o%arr(1), o%arr(2) +end program main + +! CHECK: 2 3 diff --git a/offload/test/offloading/fortran/target-defaultmap-implicit-mapper.f90 b/offload/test/offloading/fortran/target-defaultmap-implicit-mapper.f90 new file mode 100644 index 0000000..d2d8f7a --- /dev/null +++ b/offload/test/offloading/fortran/target-defaultmap-implicit-mapper.f90 @@ -0,0 +1,49 @@ +! Offload test that ensures defaultmap(tofrom: scalar) does not suppress +! implicit default mapper generation for allocatable derived types. +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic +program defaultmap_implicit_mapper + implicit none + + type :: payload_t + integer, allocatable :: arr(:) + end type payload_t + + type(payload_t), allocatable :: obj + integer, parameter :: n = 8 + integer :: i + integer :: scalar + logical :: ok + + allocate(obj) + allocate(obj%arr(n)) + obj%arr = 1 + scalar = 2 + + !$omp target defaultmap(tofrom: scalar) + do i = 1, n + obj%arr(i) = obj%arr(i) + scalar + end do + scalar = 7 + !$omp end target + + ok = .true. + do i = 1, n + if (obj%arr(i) /= 3) ok = .false. + end do + if (scalar /= 7) ok = .false. + + if (ok) then + print *, "Test passed!" + else + print *, "Test failed!" + print *, obj%arr + print *, scalar + end if + + deallocate(obj%arr) + deallocate(obj) +end program defaultmap_implicit_mapper + +! CHECK: Test passed! diff --git a/offload/test/offloading/fortran/target-parameter-array.f90 b/offload/test/offloading/fortran/target-parameter-array.f90 new file mode 100644 index 0000000..b85fb06 --- /dev/null +++ b/offload/test/offloading/fortran/target-parameter-array.f90 @@ -0,0 +1,131 @@ +! Offload test for parameter (constant) arrays and character scalars accessed +! with dynamic indices/substrings in OpenMP target regions. + +! REQUIRES: flang, amdgpu + +! RUN: %libomptarget-compile-fortran-run-and-check-generic + +program test_parameter_mapping + implicit none + integer, parameter :: dp = selected_real_kind(15, 307) + logical :: all_tests_pass + + all_tests_pass = .true. + + ! Test 1: Parameter array with dynamic index + call test_param_array_dynamic_index(all_tests_pass) + + ! Test 2: Integer parameter array + call test_int_param_array(all_tests_pass) + + ! Test 3: Character scalar with dynamic substring + call test_char_substring(all_tests_pass) + + ! Test 4: Verify scalar parameters work (inlined) + call test_scalar_param(all_tests_pass) + + if (all_tests_pass) then + print *, "PASS" + else + print *, "FAIL" + endif + +contains + +! Test 1: Parameter array with dynamic index in target region +subroutine test_param_array_dynamic_index(test_pass) + logical, intent(inout) :: test_pass + real(dp), parameter :: const_array(3) = [1.0_dp, 2.0_dp, 3.0_dp] + integer :: idx + real(dp) :: result + real(dp), parameter :: expected = 2.0_dp + real(dp), parameter :: tolerance = 1.0e-10_dp + + idx = 2 + result = 0.0_dp + + !$omp target map(tofrom:result) map(to:idx) + ! Access parameter array with dynamic index + result = const_array(idx) + !$omp end target + + if (abs(result - expected) > tolerance) then + print *, "Test 1 FAILED: expected", expected, "got", result + test_pass = .false. + endif +end subroutine test_param_array_dynamic_index + +! Test 2: Integer parameter array with different indices +subroutine test_int_param_array(test_pass) + logical, intent(inout) :: test_pass + integer, parameter :: int_array(4) = [10, 20, 30, 40] + integer :: idx1, idx2 + integer :: result1, result2 + + idx1 = 1 + idx2 = 4 + result1 = 0 + result2 = 0 + + !$omp target map(tofrom:result1, result2) map(to:idx1, idx2) + ! Access parameter array with different dynamic indices + result1 = int_array(idx1) + result2 = int_array(idx2) + !$omp end target + + if (result1 /= 10 .or. result2 /= 40) then + print *, "Test 2 FAILED: expected 10, 40 got", result1, result2 + test_pass = .false. + endif +end subroutine test_int_param_array + +! Test 3: Character scalar parameter with dynamic substring access +subroutine test_char_substring(test_pass) + logical, intent(inout) :: test_pass + character(len=20), parameter :: char_scalar = "constant_string_data" + integer :: start_idx, end_idx + character(len=8) :: result + character(len=8), parameter :: expected = "string_d" + + start_idx = 10 + end_idx = 17 + result = "" + + !$omp target map(tofrom:result) map(to:start_idx, end_idx) + ! Dynamic substring access - character scalar must be mapped + result = char_scalar(start_idx:end_idx) + !$omp end target + + if (result /= expected) then + print *, "Test 3 FAILED: expected '", expected, "' got '", result, "'" + test_pass = .false. + endif +end subroutine test_char_substring + +! Test 4: Scalar parameter (can be inlined, no mapping needed) +subroutine test_scalar_param(test_pass) + logical, intent(inout) :: test_pass + integer, parameter :: scalar_const = 42 + real(dp), parameter :: real_const = 3.14159_dp + integer :: int_result + real(dp) :: real_result + real(dp), parameter :: tolerance = 1.0e-5_dp + + int_result = 0 + real_result = 0.0_dp + + !$omp target map(tofrom:int_result, real_result) + ! Scalar parameters should be inlined (no mapping needed) + int_result = scalar_const + real_result = real_const + !$omp end target + + if (int_result /= 42 .or. abs(real_result - real_const) > tolerance) then + print *, "Test 4 FAILED: expected 42, 3.14159 got", int_result, real_result + test_pass = .false. + endif +end subroutine test_scalar_param + +end program test_parameter_mapping + +! CHECK: PASS diff --git a/offload/test/offloading/strided_multiple_update.c b/offload/test/offloading/strided_multiple_update_from.c index 4f2df81..4f2df81 100644 --- a/offload/test/offloading/strided_multiple_update.c +++ b/offload/test/offloading/strided_multiple_update_from.c diff --git a/offload/test/offloading/strided_multiple_update_to.c b/offload/test/offloading/strided_multiple_update_to.c new file mode 100644 index 0000000..bb16d7a --- /dev/null +++ b/offload/test/offloading/strided_multiple_update_to.c @@ -0,0 +1,124 @@ +// This test checks that #pragma omp target update to(data1[0:3:4], +// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays +// from the host to the device. + +// RUN: %libomptarget-compile-run-and-check-generic +// XFAIL: intelgpu + +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 12; + double data1[len], data2[len]; + + // Initialize host arrays + for (int i = 0; i < len; i++) { + data1[i] = i; + data2[i] = i * 10; + } + + printf("original host array values:\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data2[i]); + + // CHECK: original host array values: + // CHECK-NEXT: data1: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 1.0 + // CHECK-NEXT: 2.0 + // CHECK-NEXT: 3.0 + // CHECK-NEXT: 4.0 + // CHECK-NEXT: 5.0 + // CHECK-NEXT: 6.0 + // CHECK-NEXT: 7.0 + // CHECK-NEXT: 8.0 + // CHECK-NEXT: 9.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 11.0 + // CHECK-NEXT: data2: + // CHECK-NEXT: 0.0 + // CHECK-NEXT: 10.0 + // CHECK-NEXT: 20.0 + // CHECK-NEXT: 30.0 + // CHECK-NEXT: 40.0 + // CHECK-NEXT: 50.0 + // CHECK-NEXT: 60.0 + // CHECK-NEXT: 70.0 + // CHECK-NEXT: 80.0 + // CHECK-NEXT: 90.0 + // CHECK-NEXT: 100.0 + // CHECK-NEXT: 110.0 + +#pragma omp target data map(tofrom : data1[0 : len], data2[0 : len]) + { + // Initialize device arrays to 20 +#pragma omp target + { + for (int i = 0; i < len; i++) { + data1[i] = 20.0; + data2[i] = 20.0; + } + } + + // Modify host arrays for strided elements + data1[0] = 10.0; + data1[4] = 10.0; + data1[8] = 10.0; + data2[0] = 10.0; + data2[5] = 10.0; + + // data1[0:3:4] // indices 0,4,8 + // data2[0:2:5] // indices 0,5 +#pragma omp target update to(data1[0 : 3 : 4], data2[0 : 2 : 5]) + + // Verify on device by adding 5 +#pragma omp target + { + for (int i = 0; i < len; i++) + data1[i] += 5.0; + for (int i = 0; i < len; i++) + data2[i] += 5.0; + } + } + + printf("device array values after update to:\n"); + printf("data1:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data1[i]); + printf("data2:\n"); + for (int i = 0; i < len; i++) + printf("%.1f\n", data2[i]); + + // CHECK: device array values after update to: + // CHECK-NEXT: data1: + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: data2: + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 15.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 + // CHECK-NEXT: 25.0 +} diff --git a/offload/test/offloading/strided_partial_update.c b/offload/test/offloading/strided_partial_update_from.c index 4a2977a..4a2977a 100644 --- a/offload/test/offloading/strided_partial_update.c +++ b/offload/test/offloading/strided_partial_update_from.c diff --git a/offload/test/offloading/strided_partial_update_to.c b/offload/test/offloading/strided_partial_update_to.c new file mode 100644 index 0000000..f9c960f --- /dev/null +++ b/offload/test/offloading/strided_partial_update_to.c @@ -0,0 +1,74 @@ +// This test checks that #pragma omp target update to(data[0:4:3]) correctly +// updates every third element (stride 3) from the host to the device, partially +// across the array + +// RUN: %libomptarget-compile-run-and-check-generic +// XFAIL: intelgpu + +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 11; + double data[len]; + + // Initialize on host + for (int i = 0; i < len; i++) + data[i] = i; + + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + // CHECK: 0.000000 + // CHECK: 1.000000 + // CHECK: 2.000000 + // CHECK: 3.000000 + // CHECK: 4.000000 + // CHECK: 5.000000 + // CHECK: 6.000000 + // CHECK: 7.000000 + // CHECK: 8.000000 + // CHECK: 9.000000 + // CHECK: 10.000000 + +#pragma omp target data map(tofrom : data[0 : len]) + { + // Initialize device array to 20 +#pragma omp target + for (int i = 0; i < len; i++) + data[i] = 20.0; + + // Modify host data for strided elements + data[0] = 10.0; + data[3] = 10.0; + data[6] = 10.0; + data[9] = 10.0; + +#pragma omp target update to(data[0 : 4 : 3]) // indices 0,3,6,9 + + // Verify on device by adding 5 +#pragma omp target + for (int i = 0; i < len; i++) + data[i] += 5.0; + } + + printf("device array values after update to:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 +} diff --git a/offload/test/offloading/strided_update.c b/offload/test/offloading/strided_update_from.c index 9910bed..9910bed 100644 --- a/offload/test/offloading/strided_update.c +++ b/offload/test/offloading/strided_update_from.c diff --git a/offload/test/offloading/strided_update_to.c b/offload/test/offloading/strided_update_to.c new file mode 100644 index 0000000..eca20ab --- /dev/null +++ b/offload/test/offloading/strided_update_to.c @@ -0,0 +1,74 @@ +// This test checks that "update to" clause in OpenMP is supported when the +// elements are updated in a non-contiguous manner. This test checks that +// #pragma omp target update to(data[0:4:2]) correctly updates only every +// other element (stride 2) from the host to the device + +// RUN: %libomptarget-compile-run-and-check-generic +// XFAIL: intelgpu + +#include <omp.h> +#include <stdio.h> + +int main() { + int len = 8; + double data[len]; + + // Initialize on host + for (int i = 0; i < len; i++) { + data[i] = i; + } + + // Initial values + printf("original host array values:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + +#pragma omp target data map(tofrom : len, data[0 : len]) + { + // Initialize device to 20 +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] = 20.0; + } + + // Modify host for strided elements + data[0] = 10.0; + data[2] = 10.0; + data[4] = 10.0; + data[6] = 10.0; + +#pragma omp target update to(data[0 : 4 : 2]) + + // Verify on device by adding 5 +#pragma omp target + for (int i = 0; i < len; i++) { + data[i] += 5.0; + } + } + + // CHECK: 0.000000 + // CHECK: 1.000000 + // CHECK: 2.000000 + // CHECK: 3.000000 + // CHECK: 4.000000 + // CHECK: 5.000000 + // CHECK: 6.000000 + // CHECK: 7.000000 + + printf("device array values after update to:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + // CHECK: 15.000000 + // CHECK: 25.000000 + + return 0; +} |