[Offload][OpenMP][Lit-tests] Added cpp tests for non-contiguous strided update

[amitamd7]

The patch (#144635) additionally fixes non-contiguous update in C plus plus. This PR adds offload cpp lit-tests for the same.

[llvmbot]

@llvm/pr-subscribers-offload

Author: Amit Tiwari (amitamd7)

Changes

The patch (#144635) additionally fixes non-contiguous update in C plus plus. This PR adds offload cpp lit-tests for the same.

---

Full diff: https://github.com/llvm/llvm-project/pull/156828.diff

3 Files Affected:

• (added) offload/test/offloading/strided_multiple_update.cpp (+67)
• (added) offload/test/offloading/strided_partial_update.cpp (+68)
• (added) offload/test/offloading/strided_update.cpp (+61)

diff --git a/offload/test/offloading/strided_multiple_update.cpp b/offload/test/offloading/strided_multiple_update.cpp
new file mode 100644
index 0000000000000..0661ac8e6a64f
--- /dev/null
+++ b/offload/test/offloading/strided_multiple_update.cpp
@@ -0,0 +1,67 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+// This test checks that #pragma omp target update from(data1[0:3:4],
+// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays
+// from the device to the host.
+
+#include <iomanip>
+#include <iostream>
+#include <omp.h>
+
+using namespace std;
+
+int main() {
+  int len = 12;
+  double data1[len], data2[len];
+
+// Initial values
+#pragma omp target map(tofrom : data1[0 : len], data2[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      data1[i] = i;
+      data2[i] = i * 10;
+    }
+  }
+
+  cout << "original host array values:" << endl;
+  cout << "data1: ";
+  for (int i = 0; i < len; i++)
+    cout << fixed << setprecision(1) << data1[i] << " ";
+  cout << endl << "data2: ";
+  for (int i = 0; i < len; i++)
+    cout << fixed << setprecision(1) << data2[i] << " ";
+  cout << endl << endl;
+
+#pragma omp target data map(to : data1[0 : len], data2[0 : len])
+  {
+// Modify arrays on device
+#pragma omp target
+    {
+      for (int i = 0; i < len; i++)
+        data1[i] += i;
+      for (int i = 0; i < len; i++)
+        data2[i] += 100;
+    }
+
+// data1[0:3:4]  // indices 0,4,8
+// data2[0:2:5]  // indices 0,5
+#pragma omp target update from(data1[0 : 3 : 4], data2[0 : 2 : 5])
+  }
+
+  cout << "device array values after update from:" << endl;
+  cout << "data1: ";
+  for (int i = 0; i < len; i++)
+    cout << fixed << setprecision(1) << data1[i] << " ";
+  cout << endl << "data2: ";
+  for (int i = 0; i < len; i++)
+    cout << fixed << setprecision(1) << data2[i] << " ";
+  cout << endl << endl;
+
+  // CHECK: data1: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0
+  // CHECK: data2: 0.0 10.0 20.0 30.0 40.0 50.0 60.0 70.0 80.0 90.0 100.0 110.0
+
+  // CHECK: data1: 0.0 1.0 2.0 3.0 8.0 5.0 6.0 7.0 16.0 9.0 10.0 11.0
+  // CHECK: data2: 100.0 10.0 20.0 30.0 40.0 150.0 60.0 70.0 80.0 90.0 100.0
+  // 110.0
+
+  return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/strided_partial_update.cpp b/offload/test/offloading/strided_partial_update.cpp
new file mode 100644
index 0000000000000..8151034c145eb
--- /dev/null
+++ b/offload/test/offloading/strided_partial_update.cpp
@@ -0,0 +1,68 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+// This test checks that #pragma omp target update from(data[0:4:3]) correctly
+// updates every third element (stride 3) from the device to the host, partially
+// across the array
+
+#include <iomanip>
+#include <iostream>
+#include <omp.h>
+
+using namespace std;
+
+int main() {
+  int len = 11;
+  double data[len];
+
+#pragma omp target map(tofrom : data[0 : len])
+  {
+    for (int i = 0; i < len; i++)
+      data[i] = i;
+  }
+
+  // Initial values
+  cout << "original host array values:" << endl;
+  for (int i = 0; i < len; i++)
+    cout << fixed << setprecision(6) << data[i] << endl;
+  cout << endl;
+
+#pragma omp target data map(to : data[0 : len])
+  {
+// Modify arrays on device
+#pragma omp target
+    for (int i = 0; i < len; i++)
+      data[i] += i;
+
+#pragma omp target update from(data[0 : 4 : 3]) // indices 0,3,6,9
+  }
+
+  cout << "device array values after update from:" << endl;
+  for (int i = 0; i < len; i++)
+    cout << fixed << setprecision(6) << data[i] << endl;
+  cout << endl;
+
+  // 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
+
+  // CHECK: 0.000000
+  // CHECK: 1.000000
+  // CHECK: 2.000000
+  // CHECK: 6.000000
+  // CHECK: 4.000000
+  // CHECK: 5.000000
+  // CHECK: 12.000000
+  // CHECK: 7.000000
+  // CHECK: 8.000000
+  // CHECK: 18.000000
+  // CHECK: 10.000000
+
+  return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/strided_update.cpp b/offload/test/offloading/strided_update.cpp
new file mode 100644
index 0000000000000..3c4d9515b0d90
--- /dev/null
+++ b/offload/test/offloading/strided_update.cpp
@@ -0,0 +1,61 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+// This test checks that "update from" clause in OpenMP is supported when the
+// elements are updated in a non-contiguous manner. This test checks that
+// #pragma omp target update from(data[0:4:2]) correctly updates only every
+// other element (stride 2) from the device to the host
+
+#include <iomanip>
+#include <iostream>
+#include <omp.h>
+
+using namespace std;
+
+int main() {
+  int len = 8;
+  double data[len];
+
+#pragma omp target map(tofrom : len, data[0 : len])
+  {
+    for (int i = 0; i < len; i++) {
+      data[i] = i;
+    }
+  }
+
+  // Initial values
+  cout << "original host array values:" << endl;
+  for (int i = 0; i < len; i++) {
+    cout << fixed << setprecision(6) << data[i] << endl;
+  }
+  cout << endl;
+
+#pragma omp target data map(to : len, data[0 : len])
+  {
+// Modify arrays on device
+#pragma omp target
+    for (int i = 0; i < len; i++) {
+      data[i] += i;
+    }
+
+#pragma omp target update from(data[0 : 4 : 2])
+  }
+
+  // CHECK: 0.000000
+  // CHECK: 1.000000
+  // CHECK: 4.000000
+  // CHECK: 3.000000
+  // CHECK: 8.000000
+  // CHECK: 5.000000
+  // CHECK: 12.000000
+  // CHECK: 7.000000
+  // CHECK-NOT: 2.000000
+  // CHECK-NOT: 6.000000
+  // CHECK-NOT: 10.000000
+  // CHECK-NOT: 14.000000
+  cout << "from target array results:" << endl;
+  for (int i = 0; i < len; i++) {
+    cout << fixed << setprecision(6) << data[i] << endl;
+  }
+  cout << endl;
+
+  return 0;
+}