Skip to content

Commit bdff136

Browse files
amitamd7amitamd7
committed
strided_update_offloading with lit-tests and clang-unittests added
1 parent f64d5df commit bdff136

File tree

9 files changed

+322
-1
lines changed

9 files changed

+322
-1
lines changed

‎clang/include/clang/ASTMatchers/ASTMatchers.h

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8713,6 +8713,33 @@ AST_MATCHER_P(OMPExecutableDirective, hasAnyClause,
87138713
Builder) != Clauses.end();
87148714
}
87158715

8716+
/// Matches any ``#pragma omp target update`` executable directive.
8717+
///
8718+
/// Given
8719+
///
8720+
/// \code
8721+
/// #pragma omp target update from(a)
8722+
/// #pragma omp target update to(b)
8723+
/// \endcode
8724+
///
8725+
/// ``ompTargetUpdateDirective()`` matches both ``omp target update from(a)``
8726+
/// and ``omp target update to(b)``.
8727+
extern const internal::VariadicDynCastAllOfMatcher<Stmt,
8728+
OMPTargetUpdateDirective>
8729+
ompTargetUpdateDirective;
8730+
8731+
/// Matches OpenMP ``from`` clause.
8732+
///
8733+
/// Given
8734+
///
8735+
/// \code
8736+
/// #pragma omp target update from(a)
8737+
/// \endcode
8738+
///
8739+
/// ``ompFromClause()`` matches ``from(a)``.
8740+
extern const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPFromClause>
8741+
ompFromClause;
8742+
87168743
/// Matches OpenMP ``default`` clause.
87178744
///
87188745
/// Given

‎clang/lib/ASTMatchers/ASTMatchersInternal.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1121,6 +1121,10 @@ AST_TYPELOC_TRAVERSE_MATCHER_DEF(
11211121

11221122
const internal::VariadicDynCastAllOfMatcher<Stmt, OMPExecutableDirective>
11231123
ompExecutableDirective;
1124+
const internal::VariadicDynCastAllOfMatcher<Stmt, OMPTargetUpdateDirective>
1125+
ompTargetUpdateDirective;
1126+
const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPFromClause>
1127+
ompFromClause;
11241128
const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPDefaultClause>
11251129
ompDefaultClause;
11261130
const internal::VariadicDynCastAllOfMatcher<Decl, CXXDeductionGuideDecl>

‎clang/lib/ASTMatchers/Dynamic/Registry.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -532,6 +532,8 @@ RegistryMaps::RegistryMaps() {
532532
REGISTER_MATCHER(ofKind);
533533
REGISTER_MATCHER(ompDefaultClause);
534534
REGISTER_MATCHER(ompExecutableDirective);
535+
REGISTER_MATCHER(ompTargetUpdateDirective);
536+
REGISTER_MATCHER(ompFromClause);
535537
REGISTER_MATCHER(on);
536538
REGISTER_MATCHER(onImplicitObjectArgument);
537539
REGISTER_MATCHER(opaqueValueExpr);

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 25 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7378,7 +7378,31 @@ class MappableExprsHandler {
73787378
// dimension.
73797379
uint64_t DimSize = 1;
73807380

7381-
bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
7381+
// Detects non-contiguous updates due to strided accesses.
7382+
// Sets the 'IsNonContiguous' flag so that the 'MapType' bits are set
7383+
// correctly when generating information to be passed to the runtime. The
7384+
// flag is set to true if any array section has a stride not equal to 1, or
7385+
// if the stride is not a constant expression (conservatively assumed
7386+
// non-contiguous).
7387+
bool IsNonContiguous = [&]() -> bool {
7388+
for (const auto &Component : Components) {
7389+
const auto *OASE =
7390+
dyn_cast<ArraySectionExpr>(Component.getAssociatedExpression());
7391+
if (OASE) {
7392+
const Expr *StrideExpr = OASE->getStride();
7393+
if (StrideExpr) {
7394+
if (const auto Constant =
7395+
StrideExpr->getIntegerConstantExpr(CGF.getContext())) {
7396+
if (!Constant->isOne()) {
7397+
return true;
7398+
}
7399+
}
7400+
}
7401+
}
7402+
}
7403+
return false;
7404+
}();
7405+
73827406
bool IsPrevMemberReference = false;
73837407

73847408
bool IsPartialMapped =

‎clang/unittests/ASTMatchers/ASTMatchersNarrowingTest.cpp

Lines changed: 59 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4724,6 +4724,65 @@ void x() {
47244724
EXPECT_TRUE(matchesWithOpenMP(Source8, Matcher));
47254725
}
47264726

4727+
TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_IsStandaloneDirective) {
4728+
auto Matcher = ompTargetUpdateDirective(isStandaloneDirective());
4729+
4730+
StringRef Source0 = R"(
4731+
void foo() {
4732+
int arr[8];
4733+
#pragma omp target update from(arr[0:8:2])
4734+
;
4735+
}
4736+
)";
4737+
EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
4738+
}
4739+
4740+
TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_HasStructuredBlock) {
4741+
StringRef Source0 = R"(
4742+
void foo() {
4743+
int arr[8];
4744+
#pragma omp target update from(arr[0:8:2])
4745+
;
4746+
}
4747+
)";
4748+
EXPECT_TRUE(notMatchesWithOpenMP(
4749+
Source0, ompTargetUpdateDirective(hasStructuredBlock(nullStmt()))));
4750+
}
4751+
4752+
TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_HasClause) {
4753+
auto Matcher = ompTargetUpdateDirective(hasAnyClause(anything()));
4754+
4755+
StringRef Source0 = R"(
4756+
void foo() {
4757+
int arr[8];
4758+
#pragma omp target update from(arr[0:8:2])
4759+
;
4760+
}
4761+
)";
4762+
EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
4763+
}
4764+
4765+
TEST_P(ASTMatchersTest, OMPTargetUpdateDirective_IsAllowedToContainClauseKind) {
4766+
auto Matcher = ompTargetUpdateDirective(
4767+
isAllowedToContainClauseKind(llvm::omp::OMPC_from));
4768+
4769+
StringRef Source0 = R"(
4770+
void x() {
4771+
;
4772+
}
4773+
)";
4774+
EXPECT_TRUE(notMatchesWithOpenMP(Source0, Matcher));
4775+
4776+
StringRef Source1 = R"(
4777+
void foo() {
4778+
int arr[8];
4779+
#pragma omp target update from(arr[0:8:2])
4780+
;
4781+
}
4782+
)";
4783+
EXPECT_TRUE(matchesWithOpenMP(Source1, Matcher));
4784+
}
4785+
47274786
TEST_P(ASTMatchersTest, HasAnyBase_DirectBase) {
47284787
if (!GetParam().isCXX()) {
47294788
return;

‎clang/unittests/ASTMatchers/ASTMatchersNodeTest.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2779,6 +2779,32 @@ void x() {
27792779
EXPECT_TRUE(notMatchesWithOpenMP(Source2, Matcher));
27802780
}
27812781

2782+
TEST(ASTMatchersTestOpenMP, OMPTargetUpdateDirective) {
2783+
auto Matcher = stmt(ompTargetUpdateDirective());
2784+
2785+
StringRef Source0 = R"(
2786+
void foo() {
2787+
int arr[8];
2788+
#pragma omp target update from(arr[0:8:2])
2789+
;
2790+
}
2791+
)";
2792+
EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
2793+
}
2794+
2795+
TEST(ASTMatchersTestOpenMP, OMPFromClause) {
2796+
auto Matcher = ompTargetUpdateDirective(hasAnyClause(ompFromClause()));
2797+
2798+
StringRef Source0 = R"(
2799+
void foo() {
2800+
int arr[8];
2801+
#pragma omp target update from(arr[0:8:2])
2802+
;
2803+
}
2804+
)";
2805+
EXPECT_TRUE(matchesWithOpenMP(Source0, Matcher));
2806+
}
2807+
27822808
TEST(ASTMatchersTestOpenMP, OMPDefaultClause) {
27832809
auto Matcher = ompExecutableDirective(hasAnyClause(ompDefaultClause()));
27842810

‎offload/test/offloading/strided_multiple_update.c

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,62 @@
1+
// This test checks that #pragma omp target update from(data1[0:3:4],
2+
// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays
3+
// from the device to the host.
4+
5+
// RUN: %libomptarget-compile-run-and-check-generic
6+
#include <omp.h>
7+
#include <stdio.h>
8+
9+
int main() {
10+
int len = 12;
11+
double data1[len], data2[len];
12+
13+
// Initial values
14+
#pragma omp target map(tofrom : data1[0 : len], data2[0 : len])
15+
{
16+
for (int i = 0; i < len; i++) {
17+
data1[i] = i;
18+
data2[i] = i * 10;
19+
}
20+
}
21+
22+
printf("original host array values:\n");
23+
printf("data1: ");
24+
for (int i = 0; i < len; i++)
25+
printf("%.1f ", data1[i]);
26+
printf("\ndata2: ");
27+
for (int i = 0; i < len; i++)
28+
printf("%.1f ", data2[i]);
29+
printf("\n\n");
30+
31+
#pragma omp target data map(to : data1[0 : len], data2[0 : len])
32+
{
33+
// Modify arrays on device
34+
#pragma omp target
35+
{
36+
for (int i = 0; i < len; i++)
37+
data1[i] += i;
38+
for (int i = 0; i < len; i++)
39+
data2[i] += 100;
40+
}
41+
42+
// data1[0:3:4] // indices 0,4,8
43+
// data2[0:2:5] // indices 0,5
44+
#pragma omp target update from(data1[0 : 3 : 4], data2[0 : 2 : 5])
45+
}
46+
47+
printf("device array values after update from:\n");
48+
printf("data1: ");
49+
for (int i = 0; i < len; i++)
50+
printf("%.1f ", data1[i]);
51+
printf("\ndata2: ");
52+
for (int i = 0; i < len; i++)
53+
printf("%.1f ", data2[i]);
54+
printf("\n\n");
55+
56+
// 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
57+
// 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
58+
59+
// 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
60+
// CHECK: data2: 100.0 10.0 20.0 30.0 40.0 150.0 60.0 70.0 80.0 90.0 100.0
61+
// 110.0
62+
}

‎offload/test/offloading/strided_partial_update.c

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,63 @@
1+
// This test checks that #pragma omp target update from(data[0:4:3]) correctly
2+
// updates every third element (stride 3) from the device to the host, partially
3+
// across the array
4+
5+
// RUN: %libomptarget-compile-run-and-check-generic
6+
#include <omp.h>
7+
#include <stdio.h>
8+
9+
int main() {
10+
int len = 11;
11+
double data[len];
12+
13+
#pragma omp target map(tofrom : data[0 : len])
14+
{
15+
for (int i = 0; i < len; i++)
16+
data[i] = i;
17+
}
18+
19+
// Initial values
20+
printf("original host array values:\n");
21+
for (int i = 0; i < len; i++)
22+
printf("%f\n", data[i]);
23+
printf("\n");
24+
25+
#pragma omp target data map(to : data[0 : len])
26+
{
27+
// Modify arrays on device
28+
#pragma omp target
29+
for (int i = 0; i < len; i++)
30+
data[i] += i;
31+
32+
#pragma omp target update from(data[0 : 4 : 3]) // indices 0,3,6,9
33+
}
34+
35+
printf("device array values after update from:\n");
36+
for (int i = 0; i < len; i++)
37+
printf("%f\n", data[i]);
38+
printf("\n");
39+
40+
// CHECK: 0.000000
41+
// CHECK: 1.000000
42+
// CHECK: 2.000000
43+
// CHECK: 3.000000
44+
// CHECK: 4.000000
45+
// CHECK: 5.000000
46+
// CHECK: 6.000000
47+
// CHECK: 7.000000
48+
// CHECK: 8.000000
49+
// CHECK: 9.000000
50+
// CHECK: 10.000000
51+
52+
// CHECK: 0.000000
53+
// CHECK: 1.000000
54+
// CHECK: 2.000000
55+
// CHECK: 6.000000
56+
// CHECK: 4.000000
57+
// CHECK: 5.000000
58+
// CHECK: 12.000000
59+
// CHECK: 7.000000
60+
// CHECK: 8.000000
61+
// CHECK: 18.000000
62+
// CHECK: 10.000000
63+
}

‎offload/test/offloading/strided_update.c

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
// This test checks that "update from" clause in OpenMP is supported when the
2+
// elements are updated in a non-contiguous manner. This test checks that
3+
// #pragma omp target update from(data[0:4:2]) correctly updates only every
4+
// other element (stride 2) from the device to the host
5+
6+
// RUN: %libomptarget-compile-run-and-check-generic
7+
#include <omp.h>
8+
#include <stdio.h>
9+
10+
int main() {
11+
int len = 8;
12+
double data[len];
13+
#pragma omp target map(tofrom : len, data[0 : len])
14+
{
15+
for (int i = 0; i < len; i++) {
16+
data[i] = i;
17+
}
18+
}
19+
// Initial values
20+
printf("original host array values:\n");
21+
for (int i = 0; i < len; i++)
22+
printf("%f\n", data[i]);
23+
printf("\n");
24+
25+
#pragma omp target data map(to : len, data[0 : len])
26+
{
27+
// Modify arrays on device
28+
#pragma omp target
29+
for (int i = 0; i < len; i++) {
30+
data[i] += i;
31+
}
32+
33+
#pragma omp target update from(data[0 : 4 : 2])
34+
}
35+
// CHECK: 0.000000
36+
// CHECK: 1.000000
37+
// CHECK: 4.000000
38+
// CHECK: 3.000000
39+
// CHECK: 8.000000
40+
// CHECK: 5.000000
41+
// CHECK: 12.000000
42+
// CHECK: 7.000000
43+
// CHECK-NOT: 2.000000
44+
// CHECK-NOT: 6.000000
45+
// CHECK-NOT: 10.000000
46+
// CHECK-NOT: 14.000000
47+
48+
printf("from target array results:\n");
49+
for (int i = 0; i < len; i++)
50+
printf("%f\n", data[i]);
51+
printf("\n");
52+
53+
return 0;
54+
}

0 commit comments

Comments
 (0)