https://github.com/amitamd7 updated
https://github.com/llvm/llvm-project/pull/144635
>From 9afb5a64977ebc256eacefbbf591e45153ab967d Mon Sep 17 00:00:00 2001
From: amtiwari <amtiw...@amd.com>
Date: Mon, 16 Jun 2025 01:07:01 -0400
Subject: [PATCH] strided_update_offloading with lit-offload-tests,
lit-clang-tests and clang-unittests
---
clang/docs/OpenMPSupport.rst | 2 +-
clang/docs/ReleaseNotes.rst | 1 +
clang/include/clang/ASTMatchers/ASTMatchers.h | 27 ++++++++
clang/lib/ASTMatchers/ASTMatchersInternal.cpp | 4 ++
clang/lib/ASTMatchers/Dynamic/Registry.cpp | 2 +
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 30 ++++++++-
.../OpenMP/target_update_strided_messages.c | 38 +++++++++++
.../target_update_strided_multiple_messages.c | 46 ++++++++++++++
.../target_update_strided_partial_messages.c | 32 ++++++++++
.../ASTMatchers/ASTMatchersNarrowingTest.cpp | 59 +++++++++++++++++
.../ASTMatchers/ASTMatchersNodeTest.cpp | 26 ++++++++
.../test/offloading/strided_multiple_update.c | 62 ++++++++++++++++++
.../test/offloading/strided_partial_update.c | 63 +++++++++++++++++++
offload/test/offloading/strided_update.c | 54 ++++++++++++++++
14 files changed, 444 insertions(+), 2 deletions(-)
create mode 100644 clang/test/OpenMP/target_update_strided_messages.c
create mode 100644 clang/test/OpenMP/target_update_strided_multiple_messages.c
create mode 100644 clang/test/OpenMP/target_update_strided_partial_messages.c
create mode 100644 offload/test/offloading/strided_multiple_update.c
create mode 100644 offload/test/offloading/strided_partial_update.c
create mode 100644 offload/test/offloading/strided_update.c
diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst
index 986aaabe1eed4..200b2ff733957 100644
--- a/clang/docs/OpenMPSupport.rst
+++ b/clang/docs/OpenMPSupport.rst
@@ -191,7 +191,7 @@ implementation.
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | teams construct on the host device
| :good:`done` | r371553
|
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
-| device | support non-contiguous array sections for
target update | :good:`done` |
|
+| device | support non-contiguous array sections for
target update | :good:`done` |
https://github.com/llvm/llvm-project/pull/144635
|
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | pointer attachment
| :good:`done` |
|
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index bc8e157e2b1c0..ff6b01b4a7b3e 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -1139,6 +1139,7 @@ OpenMP Support
- An error is now emitted when OpenMP ``collapse`` and ``ordered`` clauses have
an argument larger than what can fit within a 64-bit integer.
- Added support for private variable reduction.
+- Fixed non-contiguous strided update in the ``omp target update`` directive
with the ``from_clause``
- Fixed mapping of arrays of structs containing nested structs with user
defined
mappers, by using compiler-generated default mappers for the outer structs
for
such maps.
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<Stmt,
+ OMPTargetUpdateDirective>
+ ompTargetUpdateDirective;
+
+/// Matches OpenMP ``from`` clause.
+///
+/// Given
+///
+/// \code
+/// #pragma omp target update from(a)
+/// \endcode
+///
+/// ``ompFromClause()`` matches ``from(a)``.
+extern const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPFromClause>
+ 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<Stmt, OMPExecutableDirective>
ompExecutableDirective;
+const internal::VariadicDynCastAllOfMatcher<Stmt, OMPTargetUpdateDirective>
+ ompTargetUpdateDirective;
+const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPFromClause>
+ ompFromClause;
const internal::VariadicDynCastAllOfMatcher<OMPClause, OMPDefaultClause>
ompDefaultClause;
const internal::VariadicDynCastAllOfMatcher<Decl, CXXDeductionGuideDecl>
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..eaced8a326fb6 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7378,7 +7378,35 @@ 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 =
+ CombinedInfo.NonContigInfo.IsNonContiguous || [&]() -> bool {
+ for (const auto &Component : Components) {
+ const auto *OASE =
+ dyn_cast<ArraySectionExpr>(Component.getAssociatedExpression());
+ if (!OASE)
+ continue;
+
+ const Expr *StrideExpr = OASE->getStride();
+ if (!StrideExpr)
+ continue;
+
+ const auto Constant =
+ StrideExpr->getIntegerConstantExpr(CGF.getContext());
+ if (!Constant)
+ continue;
+
+ if (!Constant->isOne())
+ return true;
+ }
+ return false;
+ }();
+
bool IsPrevMemberReference = false;
bool IsPartialMapped =
diff --git a/clang/test/OpenMP/target_update_strided_messages.c
b/clang/test/OpenMP/target_update_strided_messages.c
new file mode 100644
index 0000000000000..1f50af4e52805
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_messages.c
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+void foo(void) {}
+
+int main(int argc, char **argv) {
+ int len = 8;
+ double data[len];
+
+ // Valid strided array sections
+ #pragma omp target update from(data[0:4:2]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/2:2]) // OK
+ {}
+
+ #pragma omp target update from(data[1:3:2]) // OK
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update from(data[0:4]) // OK
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update from(data[0:4:0]) // expected-error {{section
stride is evaluated to a non-positive value 0}} expected-error {{expected at
least one 'to' clause or 'from' clause specified to '#pragma omp target
update'}}
+
+ #pragma omp target update from(data[0:4:-1]) // expected-error {{section
stride is evaluated to a non-positive value -1}} expected-error {{expected at
least one 'to' clause or 'from' clause specified to '#pragma omp target
update'}}
+
+ // Missing colon
+ #pragma omp target update from(data[0:4 2]) // expected-error {{expected
']'}} expected-note {{to match this '['}} expected-error {{expected at least
one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ // Too many colons
+ #pragma omp target update from(data[0:4:2:1]) // expected-error {{expected
']'}} expected-note {{to match this '['}} expected-error {{expected at least
one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ {}
+
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_multiple_messages.c
b/clang/test/OpenMP/target_update_strided_multiple_messages.c
new file mode 100644
index 0000000000000..361d4c66c362b
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_multiple_messages.c
@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+void foo(void) {}
+
+typedef struct {
+ int len;
+ double data[12];
+} S;
+
+int main(int argc, char **argv) {
+ int len = 12;
+ double data1[len], data2[len];
+ S s;
+
+ // Valid multiple strided array sections
+ #pragma omp target update from(data1[0:4:2], data2[0:2:5]) // OK
+ {}
+
+ #pragma omp target update to(data1[1:2:3], data2[2:3:2]) // OK
+ {}
+
+ // Mixed strided and regular array sections
+ #pragma omp target update from(data1[0:len], data2[0:4:2]) // OK
+ {}
+
+ // Struct member arrays with strides
+ #pragma omp target update from(s.data[0:4:2]) // OK
+ {}
+
+ #pragma omp target update from(s.data[0:s.len/2:2]) // OK
+ {}
+
+ // Invalid stride in one of multiple sections
+ #pragma omp target update from(data1[0:3:4], data2[0:2:0]) // expected-error
{{section stride is evaluated to a non-positive value 0}}
+
+ // Complex expressions in multiple arrays
+ int stride1 = 2, stride2 = 3;
+ #pragma omp target update from(data1[0:len/2:stride1],
data2[1:len/3:stride2]) // OK
+ {}
+
+ // Missing colon
+ #pragma omp target update from(data1[0:4:2], data2[0:3 4]) // expected-error
{{expected ']'}} expected-note {{to match this '['}}
+
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_partial_messages.c
b/clang/test/OpenMP/target_update_strided_partial_messages.c
new file mode 100644
index 0000000000000..6dc286c8a1161
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_partial_messages.c
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+void foo(void) {}
+
+int main(int argc, char **argv) {
+ int len = 11;
+ double data[len];
+
+ // Valid partial strided updates
+ #pragma omp target update from(data[0:4:3]) // OK
+ {}
+
+ // Stride larger than length
+ #pragma omp target update from(data[0:2:10]) // OK
+ {}
+
+ // Valid: complex expressions
+ int offset = 1;
+ int count = 3;
+ int stride = 2;
+ #pragma omp target update from(data[offset:count:stride]) // OK
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update from(data[0:4:offset-1]) // OK if offset > 1
+ {}
+
+ #pragma omp target update from(data[0:count:0]) // expected-error {{section
stride is evaluated to a non-positive value 0}} expected-error {{expected at
least one 'to' clause or 'from' clause specified to '#pragma omp target
update'}}
+
+ return 0;
+}
\ No newline at end of file
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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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 <omp.h>
+#include <stdio.h>
+
+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;
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
