diff options
Diffstat (limited to 'offload/test')
23 files changed, 988 insertions, 1 deletions
diff --git a/offload/test/mapping/declare_mapper_target_checks.cpp b/offload/test/mapping/declare_mapper_target_checks.cpp new file mode 100644 index 0000000..562e283 --- /dev/null +++ b/offload/test/mapping/declare_mapper_target_checks.cpp @@ -0,0 +1,145 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic +// XFAIL: intelgpu + +#include <omp.h> +#include <stdio.h> + +#define TRUE 1 +#define FALSE 0 + +struct TY1 { + int i1, i2, i3; + static constexpr auto name = "TY1"; +}; +struct TY2 { + int i1, i2, i3; + static constexpr auto name = "TY2"; +}; + +// TY1 is not mapped, TY2 is +#pragma omp declare mapper(TY2 t) map(to : t.i1) map(from : t.i3) + +struct TY3 { + TY2 n; + static constexpr auto name = "TY3"; +}; +struct TY4 { + int a; + TY2 n; + int b; + static constexpr auto name = "TY4"; +}; + +template <typename T> int testType() { + T t1[2], t2[3], t3[4]; + for (int i = 0; i < 2; i++) + t1[i].i1 = t3[i].i1 = 1; + +#pragma omp target map(tofrom : t1, t2, t3) + for (int i = 0; i < 2; i++) { + t1[i].i3 = t3[i].i3 = t1[i].i1; + t1[i].i1 = t3[i].i1 = 7; + } + + for (int i = 0; i < 2; i++) { + if (t1[i].i3 != 1) { + printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, + t1[i].i3, i, t1[i].i1); + return 1; + } + if (t3[i].i3 != 1) { + printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, + t3[i].i3, i, t3[i].i1); + return 1; + } + } + + int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); + int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); + int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); + + printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name, + pt0, pt1, pt2); + return pt0 + pt1 + pt2; +} + +template <typename T> int testTypeNestedPtr(T t1[2], T t2[3], T t3[4]) { + for (int i = 0; i < 2; i++) + t1[i].n.i1 = t3[i].n.i1 = 1; + +#pragma omp target map(tofrom : t1[0 : 2], t2[0 : 3], t3[0 : 4]) + for (int i = 0; i < 2; i++) { + t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1; + t1[i].n.i1 = t3[i].n.i1 = 7; + } + + for (int i = 0; i < 2; i++) { + if (t1[i].n.i3 != t1[i].n.i1) { + printf("failed %s-ptr. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, + t1[i].n.i3, i, t1[i].n.i1); + return 1; + } + if (t3[i].n.i3 != t3[i].n.i1) { + printf("failed %s-ptr. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, + t3[i].n.i3, i, t3[i].n.i1); + return 1; + } + } + + int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); + int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); + int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); + + printf("present check for %s-ptr: t1 %i, t2 %i, t3 %i, expected 3x 0\n", + T::name, pt0, pt1, pt2); + return pt0 + pt1 + pt2; +} + +template <typename T> int testTypeNested() { + T t1[2], t2[3], t3[4]; + testTypeNestedPtr(t1, t2, t3); + for (int i = 0; i < 2; i++) + t1[i].n.i1 = t3[i].n.i1 = 1; + +#pragma omp target map(tofrom : t1, t2, t3) + for (int i = 0; i < 2; i++) { + t1[i].n.i3 = t3[i].n.i3 = t1[i].n.i1; + t1[i].n.i1 = t3[i].n.i1 = 7; + } + + for (int i = 0; i < 2; i++) { + if (t1[i].n.i3 != t1[i].n.i1) { + printf("failed %s. t1[%d].i3 (%d) != t1[%d].i1 (%d)\n", T::name, i, + t1[i].n.i3, i, t1[i].n.i1); + return 1; + } + if (t3[i].n.i3 != t3[i].n.i1) { + printf("failed %s. t3[%d].i3 (%d) != t3[%d].i1 (%d)\n", T::name, i, + t3[i].n.i3, i, t3[i].n.i1); + return 1; + } + } + + int pt0 = omp_target_is_present(&t1[0], omp_get_default_device()); + int pt1 = omp_target_is_present(&t2[1], omp_get_default_device()); + int pt2 = omp_target_is_present(&t3[2], omp_get_default_device()); + + printf("present check for %s: t1 %i, t2 %i, t3 %i, expected 3x 0\n", T::name, + pt0, pt1, pt2); + return pt0 + pt1 + pt2; +} + +int main(int argc, char **argv) { + int r = 0; + r += testType<TY1>(); + // CHECK: present check for TY1: t1 0, t2 0, t3 0, expected 3x 0 + r += testType<TY2>(); + // CHECK: present check for TY2: t1 0, t2 0, t3 0, expected 3x 0 + r += testTypeNested<TY3>(); + // CHECK: present check for TY3-ptr: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY3: t1 0, t2 0, t3 0, expected 3x 0 + r += testTypeNested<TY4>(); + // CHECK: present check for TY4-ptr: t1 0, t2 0, t3 0, expected 3x 0 + // CHECK: present check for TY4: t1 0, t2 0, t3 0, expected 3x 0 + return r; +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp new file mode 100644 index 0000000..5c232d5 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp @@ -0,0 +1,32 @@ +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-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. + +#include <stdio.h> + +int x = 0; + +struct ST { + int *a = &x; + + void f1() { + printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(a) + printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f1(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp new file mode 100644 index 0000000..fca0eee --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp @@ -0,0 +1,28 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +#include <stdio.h> + +int x = 0; + +struct ST { + int *a = &x; + + void f1() { + printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_nullify : a) + printf("%p\n", a); // OFFLOAD-NEXT: (nil) + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f1(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp new file mode 100644 index 0000000..51944c5 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp @@ -0,0 +1,28 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +#include <stdio.h> + +int x = 0; + +struct ST { + int *a = &x; + + void f1() { + printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_preserve : a) + printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f1(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp new file mode 100644 index 0000000..59a8fac --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp @@ -0,0 +1,33 @@ +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-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. + +#include <stdio.h> + +int x = 0; +int *y = &x; + +struct ST { + int *&b = y; + + void f2() { + printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(b) + printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f2(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp new file mode 100644 index 0000000..65c7173 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp @@ -0,0 +1,29 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +#include <stdio.h> + +int x = 0; +int *y = &x; + +struct ST { + int *&b = y; + + void f2() { + printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_nullify : b) + printf("%p\n", b); // OFFLOAD-NEXT: (nil) + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f2(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp new file mode 100644 index 0000000..beeb752 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp @@ -0,0 +1,29 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +#include <stdio.h> + +int x = 0; +int *y = &x; + +struct ST { + int *&b = y; + + void f2() { + printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_preserve : b) + printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f2(); +} 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.cpp index 33a3634..5be209a 100644 --- 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.cpp @@ -1,4 +1,8 @@ -// RUN: %libomptarget-compilexx-run-and-check-generic +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic // Test that when a use_device_ptr lookup fails, the // privatized pointer retains its original value by diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp new file mode 100644 index 0000000..984744cd --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp @@ -0,0 +1,21 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +#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(fb_nullify : xp) + printf("%p\n", xp); // OFFLOAD-NEXT: (nil) + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp new file mode 100644 index 0000000..197704f --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp @@ -0,0 +1,21 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +#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(fb_preserve : xp) + printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp new file mode 100644 index 0000000..1060ed9 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp @@ -0,0 +1,26 @@ +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-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. + +#include <stdio.h> +int x; +int *xp = &x; +int *&xpr = xp; + +void f2() { + printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(xpr) + printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f2(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp new file mode 100644 index 0000000..7fa76dd6 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp @@ -0,0 +1,24 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +#include <stdio.h> +int x; +int *xp = &x; +int *&xpr = xp; + +void f2() { + printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]] + // FIXME: We won't get "nil" until we start privatizing xpr. +#pragma omp target data use_device_ptr(fb_nullify : xpr) + printf("%p\n", xpr); // EXPECTED-OFFLOAD-NEXT: (nil) + // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f2(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp new file mode 100644 index 0000000..e7f8bd4 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp @@ -0,0 +1,22 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +#include <stdio.h> +int x; +int *xp = &x; +int *&xpr = xp; + +void f2() { + printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_preserve : xpr) + printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f2(); } 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; +} |