[clang] [llvm] [Offload][OpenMP][Lit-tests] Added cpp tests for non-contiguous strided update (PR #156828)
Amit Tiwari via cfe-commits
cfe-commits at lists.llvm.org
Thu Sep 4 01:34:51 PDT 2025
https://github.com/amitamd7 created https://github.com/llvm/llvm-project/pull/156828
The patch ([#144635](https://github.com/llvm/llvm-project/pull/144635)) additionally fixes non-contiguous update in C plus plus. This PR adds offload cpp lit-tests for the same.
>From 656756343fb5a294533446712dac520bcb40c177 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Fri, 25 Jul 2025 02:45:34 -0400
Subject: [PATCH 1/2] support_for_target_directive_clang_unittests
---
clang/include/clang/ASTMatchers/ASTMatchers.h | 27 +++++++++
clang/lib/ASTMatchers/ASTMatchersInternal.cpp | 4 ++
clang/lib/ASTMatchers/Dynamic/Registry.cpp | 2 +
.../ASTMatchers/ASTMatchersNarrowingTest.cpp | 59 +++++++++++++++++++
.../ASTMatchers/ASTMatchersNodeTest.cpp | 26 ++++++++
5 files changed, 118 insertions(+)
diff --git a/clang/include/clang/ASTMatchers/ASTMatchers.h b/clang/include/clang/ASTMatchers/ASTMatchers.h
index f1d88a9523838..0ddf29095c548 100644
--- a/clang/include/clang/ASTMatchers/ASTMatchers.h
+++ b/clang/include/clang/ASTMatchers/ASTMatchers.h
@@ -8723,6 +8723,21 @@ AST_MATCHER_P(OMPExecutableDirective, hasAnyClause,
Builder) != Clauses.end();
}
+/// Matches any ``#pragma omp target update`` executable directive.
+///
+/// Given
+///
+/// \code
+/// #pragma omp target update from(a)
+/// #pragma omp target update to(b)
+/// \endcode
+///
+/// ``ompTargetUpdateDirective()`` matches both ``omp target update from(a)``
+/// and ``omp target update to(b)``.
+extern const internal::VariadicDynCastAllOfMatcher<Stmt,
+ OMPTargetUpdateDirective>
+ ompTargetUpdateDirective;
+
/// Matches OpenMP ``default`` clause.
///
/// Given
@@ -8836,6 +8851,18 @@ AST_MATCHER_P(OMPExecutableDirective, isAllowedToContainClauseKind,
Finder->getASTContext().getLangOpts().OpenMP);
}
+/// Matches OpenMP ``from`` clause.
+///
+/// Given
+///
+/// \code
+/// #pragma omp target update from(a)
+/// \endcode
+///
+/// ``ompFromClause()`` matches ``from(a)``.
+extern const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPFromClause>
+ ompFromClause;
+
//----------------------------------------------------------------------------//
// End OpenMP handling.
//----------------------------------------------------------------------------//
diff --git a/clang/lib/ASTMatchers/ASTMatchersInternal.cpp b/clang/lib/ASTMatchers/ASTMatchersInternal.cpp
index 653b3810cb68b..1b01be069d9aa 100644
--- a/clang/lib/ASTMatchers/ASTMatchersInternal.cpp
+++ b/clang/lib/ASTMatchers/ASTMatchersInternal.cpp
@@ -1124,8 +1124,12 @@ AST_TYPELOC_TRAVERSE_MATCHER_DEF(
const internal::VariadicDynCastAllOfMatcher<Stmt, OMPExecutableDirective>
ompExecutableDirective;
+const internal::VariadicDynCastAllOfMatcher<Stmt, OMPTargetUpdateDirective>
+ ompTargetUpdateDirective;
const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPDefaultClause>
ompDefaultClause;
+const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPFromClause>
+ ompFromClause;
const internal::VariadicDynCastAllOfMatcher<Decl, CXXDeductionGuideDecl>
cxxDeductionGuideDecl;
diff --git a/clang/lib/ASTMatchers/Dynamic/Registry.cpp b/clang/lib/ASTMatchers/Dynamic/Registry.cpp
index 48a7b91969aef..a268e97f7aa18 100644
--- a/clang/lib/ASTMatchers/Dynamic/Registry.cpp
+++ b/clang/lib/ASTMatchers/Dynamic/Registry.cpp
@@ -528,7 +528,9 @@ RegistryMaps::RegistryMaps() {
REGISTER_MATCHER(ofClass);
REGISTER_MATCHER(ofKind);
REGISTER_MATCHER(ompDefaultClause);
+ REGISTER_MATCHER(ompFromClause);
REGISTER_MATCHER(ompExecutableDirective);
+ REGISTER_MATCHER(ompTargetUpdateDirective);
REGISTER_MATCHER(on);
REGISTER_MATCHER(onImplicitObjectArgument);
REGISTER_MATCHER(opaqueValueExpr);
diff --git a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp
index 8a957864cdd12..2c816250d5183 100644
--- a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp
+++ b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp
@@ -4734,6 +4734,65 @@ void x() {
EXPECT_TRUE(matchesWithOpenMP(Source8, Matcher));
}
+TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_IsStandaloneDirective) {
+ auto Matcher = ompTargetUpdateDirective(isStandaloneDirective());
+
+ StringRef Source0 = R"(
+ void foo() {
+ int arr[8];
+ #pragma omp target update from(arr[0:8:2])
+ ;
+ }
+ )";
+ EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
+}
+
+TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_HasStructuredBlock) {
+ StringRef Source0 = R"(
+ void foo() {
+ int arr[8];
+ #pragma omp target update from(arr[0:8:2])
+ ;
+ }
+ )";
+ EXPECT_TRUE(notMatchesWithOpenMP(
+ Source0, ompTargetUpdateDirective(hasStructuredBlock(nullStmt()))));
+}
+
+TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_HasClause) {
+ auto Matcher = ompTargetUpdateDirective(hasAnyClause(anything()));
+
+ StringRef Source0 = R"(
+ void foo() {
+ int arr[8];
+ #pragma omp target update from(arr[0:8:2])
+ ;
+ }
+ )";
+ EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
+}
+
+TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_IsAllowedToContainClauseKind) {
+ auto Matcher = ompTargetUpdateDirective(
+ isAllowedToContainClauseKind(llvm::omp::OMPC_from));
+
+ StringRef Source0 = R"(
+ void x() {
+ ;
+ }
+ )";
+ EXPECT_TRUE(notMatchesWithOpenMP(Source0, Matcher));
+
+ StringRef Source1 = R"(
+ void foo() {
+ int arr[8];
+ #pragma omp target update from(arr[0:8:2])
+ ;
+ }
+ )";
+ EXPECT_TRUE(matchesWithOpenMP(Source1, Matcher));
+}
+
TEST_P(ASTMatchersTest, HasAnyBase_DirectBase) {
if (!GetParam().isCXX()) {
return;
diff --git a/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp
index d7df9cae01f33..0035de81b2fa4 100644
--- a/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp
+++ b/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp
@@ -2742,6 +2742,32 @@ void x() {
EXPECT_TRUE(notMatchesWithOpenMP(Source2, Matcher));
}
+TEST(ASTMatchersTestOpenMP, OMPTargetUpdateDirective) {
+ auto Matcher = stmt(ompTargetUpdateDirective());
+
+ StringRef Source0 = R"(
+ void foo() {
+ int arr[8];
+ #pragma omp target update from(arr[0:8:2])
+ ;
+ }
+ )";
+ EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
+}
+
+TEST(ASTMatchersTestOpenMP, OMPFromClause) {
+ auto Matcher = ompTargetUpdateDirective(hasAnyClause(ompFromClause()));
+
+ StringRef Source0 = R"(
+ void foo() {
+ int arr[8];
+ #pragma omp target update from(arr[0:8:2])
+ ;
+ }
+ )";
+ EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
+}
+
TEST(ASTMatchersTestOpenMP, OMPDefaultClause) {
auto Matcher = ompExecutableDirective(hasAnyClause(ompDefaultClause()));
>From ba0e9c9409f2e056d349f4e39d03ebbcf1b4e84f Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Wed, 3 Sep 2025 10:34:13 -0400
Subject: [PATCH 2/2] Added_non-contiguous_update_cpptest
---
.../offloading/strided_multiple_update.cpp | 67 ++++++++++++++++++
.../offloading/strided_partial_update.cpp | 68 +++++++++++++++++++
offload/test/offloading/strided_update.cpp | 61 +++++++++++++++++
3 files changed, 196 insertions(+)
create mode 100644 offload/test/offloading/strided_multiple_update.cpp
create mode 100644 offload/test/offloading/strided_partial_update.cpp
create mode 100644 offload/test/offloading/strided_update.cpp
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;
+}
More information about the cfe-commits
mailing list