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