Rebaseusers/pcc/spr/codegen-introduce-machinefunctiongetpreferredalignment

Created using spr 1.3.6-beta.1

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;
+}