Rebaseusers/pcc/spr/codegen-introduce-machinefunctiongetpreferredalignment

Created using spr 1.3.6-beta.1

13 files changed, 443 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(); }