diff --git a/clang/include/clang/ASTMatchers/ASTMatchers.h b/clang/include/clang/ASTMatchers/ASTMatchers.h index e4d605d165324..2b85d6c8c89ea 100644 --- a/clang/include/clang/ASTMatchers/ASTMatchers.h +++ b/clang/include/clang/ASTMatchers/ASTMatchers.h @@ -8713,6 +8713,33 @@ 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 + ompTargetUpdateDirective; + +/// Matches OpenMP ``from`` clause. +/// +/// Given +/// +/// \code +/// #pragma omp target update from(a) +/// \endcode +/// +/// ``ompFromClause()`` matches ``from(a)``. +extern const internal::VariadicDynCastAllOfMatcher + ompFromClause; + /// Matches OpenMP ``default`` clause. /// /// Given diff --git a/clang/lib/ASTMatchers/ASTMatchersInternal.cpp b/clang/lib/ASTMatchers/ASTMatchersInternal.cpp index 9cc50a656d37f..08738a91a4ca1 100644 --- a/clang/lib/ASTMatchers/ASTMatchersInternal.cpp +++ b/clang/lib/ASTMatchers/ASTMatchersInternal.cpp @@ -1121,6 +1121,10 @@ AST_TYPELOC_TRAVERSE_MATCHER_DEF( const internal::VariadicDynCastAllOfMatcher ompExecutableDirective; +const internal::VariadicDynCastAllOfMatcher + ompTargetUpdateDirective; +const internal::VariadicDynCastAllOfMatcher + ompFromClause; const internal::VariadicDynCastAllOfMatcher ompDefaultClause; const internal::VariadicDynCastAllOfMatcher diff --git a/clang/lib/ASTMatchers/Dynamic/Registry.cpp b/clang/lib/ASTMatchers/Dynamic/Registry.cpp index 562df715e08ae..dcb82b9c6a5a1 100644 --- a/clang/lib/ASTMatchers/Dynamic/Registry.cpp +++ b/clang/lib/ASTMatchers/Dynamic/Registry.cpp @@ -532,6 +532,8 @@ RegistryMaps::RegistryMaps() { REGISTER_MATCHER(ofKind); REGISTER_MATCHER(ompDefaultClause); REGISTER_MATCHER(ompExecutableDirective); + REGISTER_MATCHER(ompTargetUpdateDirective); + REGISTER_MATCHER(ompFromClause); REGISTER_MATCHER(on); REGISTER_MATCHER(onImplicitObjectArgument); REGISTER_MATCHER(opaqueValueExpr); diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 8ccc37ef98a74..60a87b5ce6eea 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7378,7 +7378,31 @@ class MappableExprsHandler { // dimension. uint64_t DimSize = 1; - bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous; + // Detects non-contiguous updates due to strided accesses. + // Sets the 'IsNonContiguous' flag so that the 'MapType' bits are set + // correctly when generating information to be passed to the runtime. The + // flag is set to true if any array section has a stride not equal to 1, or + // if the stride is not a constant expression (conservatively assumed + // non-contiguous). + bool IsNonContiguous = [&]() -> bool { + for (const auto &Component : Components) { + const auto *OASE = + dyn_cast(Component.getAssociatedExpression()); + if (OASE) { + const Expr *StrideExpr = OASE->getStride(); + if (StrideExpr) { + if (const auto Constant = + StrideExpr->getIntegerConstantExpr(CGF.getContext())) { + if (!Constant->isOne()) { + return true; + } + } + } + } + } + return false; + }(); + bool IsPrevMemberReference = false; bool IsPartialMapped = diff --git a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp index 49abe881eeabb..e1eae1e95ebf1 100644 --- a/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp +++ b/clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp @@ -4724,6 +4724,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 07450a0c59ec6..c7c963507e78f 100644 --- a/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp +++ b/clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp @@ -2779,6 +2779,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())); diff --git a/offload/test/offloading/strided_multiple_update.c b/offload/test/offloading/strided_multiple_update.c new file mode 100644 index 0000000000000..a3e8d10863aef --- /dev/null +++ b/offload/test/offloading/strided_multiple_update.c @@ -0,0 +1,62 @@ +// 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. + +// RUN: %libomptarget-compile-run-and-check-generic +#include +#include + +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; + } + } + + printf("original host array values:\n"); + printf("data1: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data1[i]); + printf("\ndata2: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data2[i]); + printf("\n\n"); + +#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]) + } + + printf("device array values after update from:\n"); + printf("data1: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data1[i]); + printf("\ndata2: "); + for (int i = 0; i < len; i++) + printf("%.1f ", data2[i]); + printf("\n\n"); + + // 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 +} diff --git a/offload/test/offloading/strided_partial_update.c b/offload/test/offloading/strided_partial_update.c new file mode 100644 index 0000000000000..15d477f2b9b78 --- /dev/null +++ b/offload/test/offloading/strided_partial_update.c @@ -0,0 +1,63 @@ +// 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 + +// RUN: %libomptarget-compile-run-and-check-generic +#include +#include + +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 + 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(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 + } + + printf("device array values after update from:\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 + + // 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 +} diff --git a/offload/test/offloading/strided_update.c b/offload/test/offloading/strided_update.c new file mode 100644 index 0000000000000..fe875b7fd55c9 --- /dev/null +++ b/offload/test/offloading/strided_update.c @@ -0,0 +1,54 @@ +// 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 + +// RUN: %libomptarget-compile-run-and-check-generic +#include +#include + +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 + 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(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 + + printf("from target array results:\n"); + for (int i = 0; i < len; i++) + printf("%f\n", data[i]); + printf("\n"); + + return 0; +}