[clang] [llvm] [Clang][OpenMP] Handle check for non-contiguous mapping in pointer-based array sections (PR #157443)
Amit Tiwari via llvm-commits
llvm-commits at lists.llvm.org
Mon Dec 22 22:36:17 PST 2025
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/157443
>From 1e0f4ff5e8e2aa7c3f1928850a89768564c6b342 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 8 Sep 2025 08:23:21 -0400
Subject: [PATCH 1/8] handle_array_pointer_var_check
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 10 ++++++----
1 file changed, 6 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 43ec1572aaf90..e2ee451a6efa8 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8341,10 +8341,12 @@ class MappableExprsHandler {
ElementType = CAT->getElementType().getTypePtr();
else if (VAT)
ElementType = VAT->getElementType().getTypePtr();
- else
- assert(&Component == &*Components.begin() &&
- "Only expect pointer (non CAT or VAT) when this is the "
- "first Component");
+ else if (&Component == &*Components.begin()) {
+ // Handle pointer-based array sections like data[a:b:c]
+ if (const auto *PtrType = Ty->getAs<PointerType>()) {
+ ElementType = PtrType->getPointeeType().getTypePtr();
+ }
+ }
// If ElementType is null, then it means the base is a pointer
// (neither CAT nor VAT) and we'll attempt to get ElementType again
// for next iteration.
>From 1d425fbebec0e1281eeeaf18fe5243a4029323dd Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 16 Dec 2025 01:46:04 -0500
Subject: [PATCH 2/8] merge conflict-resolve
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 37 +++++++++++++++++++--------
1 file changed, 26 insertions(+), 11 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index e2ee451a6efa8..c96703bc87267 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8342,14 +8342,25 @@ class MappableExprsHandler {
else if (VAT)
ElementType = VAT->getElementType().getTypePtr();
else if (&Component == &*Components.begin()) {
- // Handle pointer-based array sections like data[a:b:c]
+ // If the base is a raw pointer (e.g. T *data with data[a:b:c]),
+ // there was no earlier CAT/VAT/array handling to establish
+ // ElementType. Capture the pointee type now so that subsequent
+ // components (offset/length/stride) have a concrete element type to
+ // work with. This makes pointer-backed sections behave consistently
+ // with CAT/VAT/array bases.
if (const auto *PtrType = Ty->getAs<PointerType>()) {
ElementType = PtrType->getPointeeType().getTypePtr();
}
+ } else {
+ // Any component after the first should never have a raw pointer type;
+ // by this point. ElementType must already be known (set above or in
+ // prior array / CAT / VAT handling).
+ assert(!Ty->isPointerType() &&
+ "Non-first components should not be raw pointers");
}
- // If ElementType is null, then it means the base is a pointer
- // (neither CAT nor VAT) and we'll attempt to get ElementType again
- // for next iteration.
+
+ // At this stage, if ElementType was a base pointer and we are in the
+ // first iteration, it has been computed.
if (ElementType) {
// For the case that having pointer as base, we need to remove one
// level of indirection.
@@ -9099,12 +9110,17 @@ class MappableExprsHandler {
// If there is an entry in PartialStruct it means we have a struct with
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid()) {
- GroupUnionCurInfo.NonContigInfo.Dims.push_back(0);
- emitCombinedEntry(
- CurInfo, GroupUnionCurInfo.Types, PartialStruct, AttachInfo,
- /*IsMapThis*/ !VD, OMPBuilder, VD,
- /*OffsetForMemberOfFlag=*/CombinedInfo.BasePointers.size(),
- /*NotTargetParam=*/true);
+ // Prepend a synthetic dimension of length 1 to represent the
+ // aggregated struct object. Using 1 (not 0, as 0 produced an
+ // incorrect non-contiguous descriptor (DimSize==1), causing the
+ // non-contiguous motion clause path to be skipped.) is important:
+ // * It preserves the correct rank so targetDataUpdate() computes
+ // DimSize == 2 for cases like strided array sections originating
+ // from user-defined mappers (e.g. test with s.data[0:8:2]).
+ UnionCurInfo.NonContigInfo.Dims.insert(
+ UnionCurInfo.NonContigInfo.Dims.begin(), 1);
+ emitCombinedEntry(CombinedInfo, UnionCurInfo.Types, PartialStruct,
+ /*IsMapThis*/ !VD, OMPBuilder, VD);
}
// Append this group's results to the overall CurInfo in the correct
@@ -9112,7 +9128,6 @@ class MappableExprsHandler {
CurInfo.append(GroupUnionCurInfo);
if (AttachInfo.isValid())
emitAttachEntry(CGF, CurInfo, AttachInfo);
- }
// We need to append the results of this capture to what we already have.
CombinedInfo.append(CurInfo);
>From 4692e5a7a1b82e9036fb14cd2d3b986ee9e91dbe Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 24 Nov 2025 05:55:21 -0500
Subject: [PATCH 3/8] testcases_added
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 3 +-
clang/test/OpenMP/target_update_codegen.cpp | 12 +-
.../target_update_strided_ptr_messages_from.c | 44 ++++++
.../target_update_strided_ptr_messages_to.c | 44 ++++++
...pdate_strided_ptr_multiple_messages_from.c | 64 +++++++++
..._update_strided_ptr_multiple_messages_to.c | 64 +++++++++
...update_strided_ptr_partial_messages_from.c | 41 ++++++
...t_update_strided_ptr_partial_messages_to.c | 40 ++++++
...rget_update_strided_struct_messages_from.c | 39 +++++
...target_update_strided_struct_messages_to.c | 39 +++++
...te_strided_struct_multiple_messages_from.c | 30 ++++
...date_strided_struct_multiple_messages_to.c | 32 +++++
...ate_strided_struct_partial_messages_from.c | 31 ++++
...pdate_strided_struct_partial_messages_to.c | 31 ++++
.../strided_ptr_multiple_update_from.c | 114 +++++++++++++++
.../strided_ptr_multiple_update_to.c | 130 +++++++++++++++++
.../strided_ptr_partial_update_from.c | 68 +++++++++
.../strided_ptr_partial_update_to.c | 81 +++++++++++
offload/test/offloading/target_update_from.c | 73 ++++++++++
.../target_update_strided_struct_from.c | 86 +++++++++++
...rget_update_strided_struct_multiple_from.c | 130 +++++++++++++++++
...target_update_strided_struct_multiple_to.c | 136 ++++++++++++++++++
...arget_update_strided_struct_partial_from.c | 86 +++++++++++
.../target_update_strided_struct_partial_to.c | 85 +++++++++++
.../target_update_strided_struct_to.c | 98 +++++++++++++
offload/test/offloading/target_update_to.c | 81 +++++++++++
26 files changed, 1674 insertions(+), 8 deletions(-)
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_messages_to.c
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_messages_to.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_multiple_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_multiple_messages_to.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c
create mode 100644 offload/test/offloading/strided_ptr_multiple_update_from.c
create mode 100644 offload/test/offloading/strided_ptr_multiple_update_to.c
create mode 100644 offload/test/offloading/strided_ptr_partial_update_from.c
create mode 100644 offload/test/offloading/strided_ptr_partial_update_to.c
create mode 100644 offload/test/offloading/target_update_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_multiple_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_multiple_to.c
create mode 100644 offload/test/offloading/target_update_strided_struct_partial_from.c
create mode 100644 offload/test/offloading/target_update_strided_struct_partial_to.c
create mode 100644 offload/test/offloading/target_update_strided_struct_to.c
create mode 100644 offload/test/offloading/target_update_to.c
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index c96703bc87267..01cc6cc574644 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8348,9 +8348,8 @@ class MappableExprsHandler {
// components (offset/length/stride) have a concrete element type to
// work with. This makes pointer-backed sections behave consistently
// with CAT/VAT/array bases.
- if (const auto *PtrType = Ty->getAs<PointerType>()) {
+ if (const auto *PtrType = Ty->getAs<PointerType>())
ElementType = PtrType->getPointeeType().getTypePtr();
- }
} else {
// Any component after the first should never have a raw pointer type;
// by this point. ElementType must already be known (set above or in
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index 4d9c49ebcf03a..b0a6c8b2f31a2 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -1187,7 +1187,7 @@ void foo(int arg) {
// CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x ptr]]] }
// CK21: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
-// CK21: [[SIZE:@.+]] = private unnamed_addr constant [2 x i64] zeroinitializer
+// CK21: [[SIZE:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]
// CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 299067162755073]
struct ST {
@@ -1231,11 +1231,11 @@ struct ST {
// CK21: store i64 1, ptr [[COUNT_4]],
// CK21: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
// CK21: store i64 {{4|8}}, ptr [[STRIDE_4]],
- // CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPSZ:%.+]], ptr [[MTYPE]]{{.+}})
- // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
- // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK21-DAG: [[PTRS:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0
- // CK21-DAG: store ptr [[DIMS]], ptr [[PTRS]],
+ // CK21-DAG: [[GEPBP:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
+ // CK21-DAG: [[GEPP:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0
+ // CK21-DAG: [[GEPSZ:%.+]] = getelementptr inbounds [2 x i64], ptr %.offload_sizes, i32 0, i32 0
+ // CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP]], ptr [[GEPP]], ptr [[GEPSZ]], ptr @.offload_maptypes, ptr null, ptr null)
+ // CK21: ret void
#pragma omp target update to(dptr[0:2][1:3][0:4])
}
};
diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c
new file mode 100644
index 0000000000000..cd158ff10e4de
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+extern void *malloc(__SIZE_TYPE__);
+extern void free(void *);
+
+int main(int argc, char **argv) {
+ int len = 16;
+ double *data = (double *)malloc(len * sizeof(double));
+
+ // Valid strided array sections with FROM
+ #pragma omp target update from(data[0:8:2]) // OK - even indices
+ {}
+
+ #pragma omp target update from(data[1:4:3]) // OK - odd start with stride
+ {}
+
+ #pragma omp target update from(data[2:3:5]) // OK - large stride
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update from(data[0:8]) // OK - default stride
+ {}
+
+ #pragma omp target update from(data[4:len-4]) // OK - computed length
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update from(data[0:8: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'}}
+
+ #pragma omp target update from(data[1:5:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Syntax errors
+ #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'}}
+ {}
+
+ #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'}}
+ {}
+
+ free(data);
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c
new file mode 100644
index 0000000000000..64839a8384425
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+extern void *malloc(__SIZE_TYPE__);
+extern void free(void *);
+
+int main(int argc, char **argv) {
+ int len = 16;
+ double *data = (double *)malloc(len * sizeof(double));
+
+ // Valid strided array sections with TO
+ #pragma omp target update to(data[0:8:2]) // OK - even indices
+ {}
+
+ #pragma omp target update to(data[1:4:3]) // OK - odd start with stride
+ {}
+
+ #pragma omp target update to(data[2:3:5]) // OK - large stride
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update to(data[0:8]) // OK - default stride
+ {}
+
+ #pragma omp target update to(data[4:len-4]) // OK - computed length
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update to(data[0:8: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 to(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'}}
+
+ #pragma omp target update to(data[1:5:-2]) // expected-error {{section stride is evaluated to a non-positive value -2}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ // Syntax errors
+ #pragma omp target update to(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'}}
+ {}
+
+ #pragma omp target update to(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'}}
+ {}
+
+ free(data);
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c
new file mode 100644
index 0000000000000..2462ae1c840c3
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c
@@ -0,0 +1,64 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+extern void *malloc(__SIZE_TYPE__);
+extern void free(void *);
+
+double *data;
+double *data1;
+double *data2;
+
+int main(int argc, char **argv) {
+ int len = 12;
+
+ // Allocate memory for explicit pointers
+ data = (double *)malloc(len * sizeof(double));
+ data1 = (double *)malloc(len * sizeof(double));
+ data2 = (double *)malloc(len * sizeof(double));
+
+ // Valid multiple strided array sections
+ #pragma omp target update from(data1[0:6:2], data2[0:4:3]) // OK - different strides
+ {}
+
+ #pragma omp target update from(data1[1:2:3], data2[2:3:2]) // OK - with offsets
+ {}
+
+ // Mixed strided and regular sections
+ #pragma omp target update from(data1[0:len], data2[0:4:2]) // OK - mixed
+ {}
+
+ #pragma omp target update from(data1[1:3:2], data2[0:len]) // OK - reversed mix
+ {}
+
+ // Using the single data pointer with strides
+ #pragma omp target update from(data[0:4:2]) // OK - single pointer
+ {}
+
+ // 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}}
+
+ #pragma omp target update from(data1[0:3:-1], data2[0:2:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
+
+ #pragma omp target update from(data[0:4:0], data1[0:2:1]) // 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[1:4:stride1+1], data2[0:3:stride2-1]) // OK - expressions
+ {}
+
+ // Mix all three pointers
+ #pragma omp target update from(data[0:2:3], data1[1:3:2], data2[2:2:4]) // OK - three arrays
+ {}
+
+ // Syntax errors in multiple arrays
+ #pragma omp target update from(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+ #pragma omp target update from(data1[0:4:2:3], data2[0:3:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+ #pragma omp target update from(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+ free(data);
+ free(data1);
+ free(data2);
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c
new file mode 100644
index 0000000000000..d8cebefb02606
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c
@@ -0,0 +1,64 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+extern void *malloc(__SIZE_TYPE__);
+extern void free(void *);
+
+double *data;
+double *data1;
+double *data2;
+
+int main(int argc, char **argv) {
+ int len = 12;
+
+ // Allocate memory for explicit pointers
+ data = (double *)malloc(len * sizeof(double));
+ data1 = (double *)malloc(len * sizeof(double));
+ data2 = (double *)malloc(len * sizeof(double));
+
+ // Valid multiple strided array sections
+ #pragma omp target update to(data1[0:6:2], data2[0:4:3]) // OK - different strides
+ {}
+
+ #pragma omp target update to(data1[1:2:3], data2[2:3:2]) // OK - with offsets
+ {}
+
+ // Mixed strided and regular sections
+ #pragma omp target update to(data1[0:len], data2[0:4:2]) // OK - mixed
+ {}
+
+ #pragma omp target update to(data1[1:3:2], data2[0:len]) // OK - reversed mix
+ {}
+
+ // Using the single data pointer with strides
+ #pragma omp target update to(data[0:4:2]) // OK - single pointer
+ {}
+
+ // Invalid stride in one of multiple sections
+ #pragma omp target update to(data1[0:3:4], data2[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+
+ #pragma omp target update to(data1[0:3:-1], data2[0:2:2]) // expected-error {{section stride is evaluated to a non-positive value -1}}
+
+ #pragma omp target update to(data[0:4:0], data1[0:2:1]) // 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 to(data1[1:4:stride1+1], data2[0:3:stride2-1]) // OK - expressions
+ {}
+
+ // Mix all three pointers
+ #pragma omp target update to(data[0:2:3], data1[1:3:2], data2[2:2:4]) // OK - three arrays
+ {}
+
+ // Syntax errors in multiple arrays
+ #pragma omp target update to(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+ #pragma omp target update to(data1[0:4:2:3], data2[0:3:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+ #pragma omp target update to(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+ free(data);
+ free(data1);
+ free(data2);
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c
new file mode 100644
index 0000000000000..2f36dfa84cd1b
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+extern void *malloc(__SIZE_TYPE__);
+extern void free(void *);
+
+int main(int argc, char **argv) {
+ int len = 11;
+ double *data = (double *)malloc(len * sizeof(double));
+
+ // Valid partial strided sections with FROM
+ #pragma omp target update from(data[0:2:3]) // OK - partial coverage
+ {}
+
+ #pragma omp target update from(data[1:3:4]) // OK - offset with partial stride
+ {}
+
+ #pragma omp target update from(data[2:2:5]) // OK - large partial stride
+ {}
+
+ // Stride larger than remaining elements
+ #pragma omp target update from(data[0:2:10]) // OK - stride > array size
+ {}
+
+ #pragma omp target update from(data[0:3:len]) // OK - stride = len
+ {}
+
+ // Complex expressions
+ int offset = 1;
+ int stride = 2;
+
+ // Runtime-dependent invalid strides
+ #pragma omp target update from(data[0:4:offset-1]) // OK if offset > 1
+ {}
+
+ // Compile-time invalid strides
+ #pragma omp target update from(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ free(data);
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c
new file mode 100644
index 0000000000000..ce4b315bec1ae
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+extern void *malloc(__SIZE_TYPE__);
+extern void free(void *);
+
+int main(int argc, char **argv) {
+ int len = 11;
+ double *data = (double *)malloc(len * sizeof(double));
+
+ // Valid partial strided sections with TO
+ #pragma omp target update to(data[0:2:3]) // OK - partial coverage
+ {}
+
+ #pragma omp target update to(data[1:3:4]) // OK - offset with partial stride
+ {}
+
+ #pragma omp target update to(data[2:2:5]) // OK - large partial stride
+ {}
+
+ // Stride larger than remaining elements
+ #pragma omp target update to(data[0:2:10]) // OK - stride > array size
+ {}
+
+ #pragma omp target update to(data[0:3:len]) // OK - stride = len
+ {}
+
+ int offset = 1;
+ int stride = 2;
+
+ // Runtime-dependent invalid strides
+ #pragma omp target update to(data[0:4:offset-1]) // OK if offset > 1
+ {}
+
+ // Compile-time invalid strides
+ #pragma omp target update to(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+ free(data);
+ return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_struct_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_messages_from.c
new file mode 100644
index 0000000000000..b70a2453dc122
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_messages_from.c
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 8
+typedef struct {
+ double data[N];
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+
+ // Valid strided array sections
+ #pragma omp target update from(s.data[0:4:2]) // OK
+ {}
+
+ #pragma omp target update from(s.data[1:3:2]) // OK
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update from(s.data[0:4]) // OK
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update from(s.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(s.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(s.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(s.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;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_messages_to.c
new file mode 100644
index 0000000000000..39e774014fd44
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_messages_to.c
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 8
+typedef struct {
+ double data[N];
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+
+ // Valid strided array sections
+ #pragma omp target update to(s.data[0:4:2]) // OK
+ {}
+
+ #pragma omp target update to(s.data[1:3:2]) // OK
+ {}
+
+ // Missing stride (default = 1)
+ #pragma omp target update to(s.data[0:4]) // OK
+ {}
+
+ // Invalid stride expressions
+ #pragma omp target update to(s.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 to(s.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 to(s.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 to(s.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;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_multiple_messages_from.c
new file mode 100644
index 0000000000000..3005d5bb28702
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_multiple_messages_from.c
@@ -0,0 +1,30 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 12
+typedef struct {
+ double data[N];
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s1, s2;
+ s1.len = N;
+ s2.len = N;
+
+ // Valid multiple strided array sections
+ #pragma omp target update from(s1.data[0:4:2], s2.data[0:2:5]) // OK
+ {}
+
+ // Mixed strided and regular array sections
+ #pragma omp target update from(s1.data[0:s1.len], s2.data[0:4:2]) // OK
+ {}
+
+ // Invalid stride in one of multiple sections
+ #pragma omp target update from(s1.data[0:3:4], s2.data[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+
+ // Missing colon
+ #pragma omp target update from(s1.data[0:4:2], s2.data[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_multiple_messages_to.c
new file mode 100644
index 0000000000000..c22c903472ce2
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_multiple_messages_to.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
+
+#define N 12
+typedef struct {
+ double data[N];
+ int len;
+ } T;
+
+int main(int argc, char **argv) {
+ T s1, s2;
+ s1.len = N;
+ s2.len = N;
+
+ // Valid multiple strided array sections
+ #pragma omp target update to(s1.data[0:4:2], s2.data[0:2:5]) // OK
+ {}
+
+ // Mixed strided and regular array sections
+ #pragma omp target update to(s1.data[0:s1.len], s2.data[0:4:2]) // OK
+ {}
+
+ // Invalid stride in one of multiple sections
+ #pragma omp target update to(s1.data[0:3:4], s2.data[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
+ {}
+
+ // Missing colon
+ #pragma omp target update to(s1.data[0:4:2], s2.data[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+ {}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c
new file mode 100644
index 0000000000000..0ae5790680437
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_partial_messages_from.c
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// expected-no-diagnostics
+
+#define N 11
+typedef struct {
+ double data[N];
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+
+ // Valid partial strided updates
+ #pragma omp target update from(s.data[0:4:3]) // OK
+ {}
+
+ // Stride larger than length
+ #pragma omp target update from(s.data[0:2:10]) // OK
+ {}
+
+ // Valid: complex expressions
+ int offset = 1;
+
+ // Invalid stride expressions (runtime-dependent)
+ #pragma omp target update from(s.data[0:4:offset-1]) // OK if offset > 1
+ {}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c
new file mode 100644
index 0000000000000..c003bbadaca5c
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_partial_messages_to.c
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+// expected-no-diagnostics
+
+#define N 11
+typedef struct {
+ double data[N];
+ int len;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = N;
+
+ // Valid partial strided updates
+ #pragma omp target update to(s.data[0:4:3]) // OK
+ {}
+
+ // Stride larger than length
+ #pragma omp target update to(s.data[0:2:10]) // OK
+ {}
+
+ // Valid: complex expressions
+ int offset = 1;
+
+ // Potentially invalid stride expressions depending on runtime values
+ #pragma omp target update to(s.data[0:4:offset-1]) // OK if offset > 1
+ {}
+
+ return 0;
+}
diff --git a/offload/test/offloading/strided_ptr_multiple_update_from.c b/offload/test/offloading/strided_ptr_multiple_update_from.c
new file mode 100644
index 0000000000000..afb2f1fc8dd53
--- /dev/null
+++ b/offload/test/offloading/strided_ptr_multiple_update_from.c
@@ -0,0 +1,114 @@
+// This test checks that #pragma omp target update from(data1[0:6:2],
+// data2[0:4:3]) correctly updates strided sections covering full arrays
+// from the device to the host using dynamically allocated memory.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ int len = 12;
+ double *data1 = (double *)calloc(len, sizeof(double));
+ double *data2 = (double *)calloc(len, sizeof(double));
+
+// 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:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data2[i]);
+
+#pragma omp target data map(to : data1[0 : len], data2[0 : len])
+ {
+#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:6:2] covers all even indices: 0,2,4,6,8,10 (6 elements)
+// data2[0:4:3] covers every 3rd: 0,3,6,9 (4 elements)
+#pragma omp target update from(data1[0 : 6 : 2], data2[0 : 4 : 3])
+ }
+
+ printf("device array values after update from:\n");
+ printf("data1:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data2[i]);
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: data1:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+ // CHECK-NEXT: data2:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 30.0
+ // CHECK-NEXT: 40.0
+ // CHECK-NEXT: 50.0
+ // CHECK-NEXT: 60.0
+ // CHECK-NEXT: 70.0
+ // CHECK-NEXT: 80.0
+ // CHECK-NEXT: 90.0
+ // CHECK-NEXT: 100.0
+ // CHECK-NEXT: 110.0
+
+ // CHECK: device array values after update from:
+ // CHECK-NEXT: data1:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 12.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 16.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 11.0
+ // CHECK-NEXT: data2:
+ // CHECK-NEXT: 100.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 130.0
+ // CHECK-NEXT: 40.0
+ // CHECK-NEXT: 50.0
+ // CHECK-NEXT: 160.0
+ // CHECK-NEXT: 70.0
+ // CHECK-NEXT: 80.0
+ // CHECK-NEXT: 190.0
+ // CHECK-NEXT: 100.0
+ // CHECK-NEXT: 110.0
+
+ free(data1);
+ free(data2);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/strided_ptr_multiple_update_to.c b/offload/test/offloading/strided_ptr_multiple_update_to.c
new file mode 100644
index 0000000000000..2005117c83754
--- /dev/null
+++ b/offload/test/offloading/strided_ptr_multiple_update_to.c
@@ -0,0 +1,130 @@
+// This test checks that #pragma omp target update to(data1[0:6:2],
+// data2[0:4:3]) correctly updates strided sections covering full arrays
+// from the host to the device using dynamically allocated memory.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ int len = 12;
+ double *data1 = (double *)calloc(len, sizeof(double));
+ double *data2 = (double *)calloc(len, sizeof(double));
+
+ // Initialize host arrays
+ for (int i = 0; i < len; i++) {
+ data1[i] = i;
+ data2[i] = i * 10;
+ }
+
+ printf("original host array values:\n");
+ printf("data1:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data2[i]);
+
+#pragma omp target data map(tofrom : data1[0 : len], data2[0 : len])
+ {
+// Initialize device arrays to 20
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data1[i] = 20.0;
+ data2[i] = 20.0;
+ }
+ }
+
+ // data1: even indices
+ for (int i = 0; i < 6; i++) {
+ data1[i * 2] = 10.0;
+ }
+ // data2: every 3rd index
+ for (int i = 0; i < 4; i++) {
+ data2[i * 3] = 10.0;
+ }
+
+// data1[0:6:2] updates all even indices: 0,2,4,6,8,10 (6 elements)
+// data2[0:4:3] updates every 3rd: 0,3,6,9 (4 elements)
+#pragma omp target update to(data1[0 : 6 : 2], data2[0 : 4 : 3])
+
+// Verify on device by adding 5
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data1[i] += 5.0;
+ data2[i] += 5.0;
+ }
+ }
+ }
+
+ printf("device array values after update to:\n");
+ printf("data1:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data1[i]);
+ printf("data2:\n");
+ for (int i = 0; i < len; i++)
+ printf("%.1f\n", data2[i]);
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: data1:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+ // CHECK-NEXT: data2:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 30.0
+ // CHECK-NEXT: 40.0
+ // CHECK-NEXT: 50.0
+ // CHECK-NEXT: 60.0
+ // CHECK-NEXT: 70.0
+ // CHECK-NEXT: 80.0
+ // CHECK-NEXT: 90.0
+ // CHECK-NEXT: 100.0
+ // CHECK-NEXT: 110.0
+
+ // CHECK: device array values after update to:
+ // CHECK-NEXT: data1:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: data2:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+
+ free(data1);
+ free(data2);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/strided_ptr_partial_update_from.c b/offload/test/offloading/strided_ptr_partial_update_from.c
new file mode 100644
index 0000000000000..e2f4398b2fbd7
--- /dev/null
+++ b/offload/test/offloading/strided_ptr_partial_update_from.c
@@ -0,0 +1,68 @@
+// This test checks that #pragma omp target update from(data[0:2:3]) correctly
+// updates every third element (stride 3) from the device to the host, partially
+// across the array using dynamically allocated memory.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ int len = 11;
+ double *data = (double *)calloc(len, sizeof(double));
+
+#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])
+ {
+#pragma omp target
+ for (int i = 0; i < len; i++)
+ data[i] += i;
+
+#pragma omp target update from(data[0 : 2 : 3]) // indices 0,3 only
+ }
+
+ printf("device array values after update from:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 1.000000
+ // CHECK-NEXT: 2.000000
+ // CHECK-NEXT: 3.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 5.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 7.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 9.000000
+ // CHECK-NEXT: 10.000000
+
+ // CHECK: device array values after update from:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 1.000000
+ // CHECK-NEXT: 2.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 5.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 7.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 9.000000
+ // CHECK-NEXT: 10.000000
+
+ free(data);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/strided_ptr_partial_update_to.c b/offload/test/offloading/strided_ptr_partial_update_to.c
new file mode 100644
index 0000000000000..28bd27fd04366
--- /dev/null
+++ b/offload/test/offloading/strided_ptr_partial_update_to.c
@@ -0,0 +1,81 @@
+// This test checks that #pragma omp target update to(data[0:2:3]) correctly
+// updates every third element (stride 3) from the host to the device, partially
+// across the array using dynamically allocated memory.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ int len = 11;
+ double *data = (double *)calloc(len, sizeof(double));
+
+ // Initialize on host
+ 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(tofrom : data[0 : len])
+ {
+// Initialize device array to 20
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++)
+ data[i] = 20.0;
+ }
+
+ // Modify host data for elements that will be updated
+ data[0] = 10.0;
+ data[3] = 10.0;
+
+// This creates Host→Device transfer for indices 0,3 only
+#pragma omp target update to(data[0 : 2 : 3])
+
+// Verify on device by adding 5 to all elements
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++)
+ data[i] += 5.0;
+ }
+ }
+
+ printf("device array values after update to:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+ printf("\n");
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 1.000000
+ // CHECK-NEXT: 2.000000
+ // CHECK-NEXT: 3.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 5.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 7.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 9.000000
+ // CHECK-NEXT: 10.000000
+
+ // CHECK: device array values after update to:
+ // CHECK-NEXT: 15.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 15.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 25.000000
+ // CHECK-NEXT: 25.000000
+
+ free(data);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_from.c b/offload/test/offloading/target_update_from.c
new file mode 100644
index 0000000000000..089b498ada0db
--- /dev/null
+++ b/offload/test/offloading/target_update_from.c
@@ -0,0 +1,73 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update from" clause in OpenMP supports strided
+// sections. #pragma omp target update from(result[0:N/2:2]) updates every other
+// element from device
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 32
+
+int main() {
+ double *result = (double *)calloc(N, sizeof(double));
+
+ printf("initial host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+ printf("\n");
+
+#pragma omp target data map(to : result[0 : N])
+ {
+#pragma omp target map(alloc : result[0 : N])
+ for (int i = 0; i < N; i++)
+ result[i] += i;
+
+ // Update strided elements from device: even indices 0,2,4,...,30
+#pragma omp target update from(result[0 : 16 : 2])
+ }
+
+ printf("after target update from (even indices up to 30 updated):\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+ printf("\n");
+
+ // Expected: even indices i, odd indices 0
+ // CHECK: 0.000000
+ // CHECK: 0.000000
+ // CHECK: 2.000000
+ // CHECK: 0.000000
+ // CHECK: 4.000000
+ // CHECK: 0.000000
+ // CHECK: 6.000000
+ // CHECK: 0.000000
+ // CHECK: 8.000000
+ // CHECK: 0.000000
+ // CHECK: 10.000000
+ // CHECK: 0.000000
+ // CHECK: 12.000000
+ // CHECK: 0.000000
+ // CHECK: 14.000000
+ // CHECK: 0.000000
+ // CHECK: 16.000000
+ // CHECK: 0.000000
+ // CHECK: 18.000000
+ // CHECK: 0.000000
+ // CHECK: 20.000000
+ // CHECK: 0.000000
+ // CHECK: 22.000000
+ // CHECK: 0.000000
+ // CHECK: 24.000000
+ // CHECK: 0.000000
+ // CHECK: 26.000000
+ // CHECK: 0.000000
+ // CHECK: 28.000000
+ // CHECK: 0.000000
+ // CHECK: 30.000000
+ // CHECK: 0.000000
+ // CHECK-NOT: 1.000000
+ // CHECK-NOT: 3.000000
+ // CHECK-NOT: 31.000000
+
+ free(result);
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_strided_struct_from.c b/offload/test/offloading/target_update_strided_struct_from.c
new file mode 100644
index 0000000000000..1c5759b1864c5
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_from.c
@@ -0,0 +1,86 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update from" with user-defined mapper supports strided
+// sections using fixed-size arrays in structs.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+typedef struct {
+ double data[N];
+ size_t len;
+} T;
+
+#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s;
+ s.len = N;
+
+ for (int i = 0; i < N; i++) {
+ s.data[i] = i;
+ }
+
+ printf("original host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", s.data[i]);
+ printf("\n");
+
+#pragma omp target data map(mapper(custom), tofrom : s)
+ {
+// Execute on device with explicit mapper
+#pragma omp target map(mapper(custom), tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += i;
+ }
+ }
+
+// Update strided elements from device: indices 0,2,4,6,8,10,12,14
+#pragma omp target update from(s.data[0 : 8 : 2])
+ }
+
+ printf("from target array results:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", s.data[i]);
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 1.000000
+ // CHECK-NEXT: 2.000000
+ // CHECK-NEXT: 3.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 5.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 7.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 9.000000
+ // CHECK-NEXT: 10.000000
+ // CHECK-NEXT: 11.000000
+ // CHECK-NEXT: 12.000000
+ // CHECK-NEXT: 13.000000
+ // CHECK-NEXT: 14.000000
+ // CHECK-NEXT: 15.000000
+
+ // CHECK: from target array results:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 1.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 3.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 5.000000
+ // CHECK-NEXT: 12.000000
+ // CHECK-NEXT: 7.000000
+ // CHECK-NEXT: 16.000000
+ // CHECK-NEXT: 9.000000
+ // CHECK-NEXT: 20.000000
+ // CHECK-NEXT: 11.000000
+ // CHECK-NEXT: 24.000000
+ // CHECK-NEXT: 13.000000
+ // CHECK-NEXT: 28.000000
+ // CHECK-NEXT: 15.000000
+
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_strided_struct_multiple_from.c b/offload/test/offloading/target_update_strided_struct_multiple_from.c
new file mode 100644
index 0000000000000..95faf104826d5
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_multiple_from.c
@@ -0,0 +1,130 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update from(s1.data[0:6:2],
+// s2.data[0:4:3]) correctly updates strided sections covering the full arrays
+// from device to host.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define LEN 12
+
+typedef struct {
+ double data[LEN];
+ int len;
+} T;
+
+int main() {
+ T s1, s2;
+ s1.len = LEN;
+ s2.len = LEN;
+
+ // Initialize struct arrays on host with simple sequential values
+ for (int i = 0; i < LEN; i++) {
+ s1.data[i] = i;
+ s2.data[i] = i;
+ }
+
+ printf("original host struct array values:\n");
+ printf("s1.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s1.data[i]);
+ printf("s2.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s2.data[i]);
+
+#pragma omp target data map(tofrom : s1, s2)
+ {
+// Initialize all device values to 20
+#pragma omp target map(tofrom : s1, s2)
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.data[i] = 20.0;
+ s2.data[i] = 20.0;
+ }
+ }
+
+// Modify specific strided elements on device to 10
+#pragma omp target map(tofrom : s1, s2)
+ {
+ // s1: modify even indices (0,2,4,6,8,10)
+ for (int i = 0; i < 6; i++) {
+ s1.data[i * 2] = 10.0;
+ }
+ // s2: modify every 3rd index (0,3,6,9)
+ for (int i = 0; i < 4; i++) {
+ s2.data[i * 3] = 10.0;
+ }
+ }
+
+// s1.data[0:6:2] updates only even indices: 0,2,4,6,8,10
+// s2.data[0:4:3] updates only every 3rd: 0,3,6,9
+#pragma omp target update from(s1.data[0 : 6 : 2], s2.data[0 : 4 : 3])
+ }
+
+ printf("host struct array values after update from:\n");
+ printf("s1.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s1.data[i]);
+ printf("s2.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s2.data[i]);
+
+ // CHECK: original host struct array values:
+ // CHECK-NEXT: s1.data:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+ // CHECK-NEXT: s2.data:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+
+ // CHECK: host struct array values after update from:
+ // CHECK-NEXT: s1.data:
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: s2.data:
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 20.0
+
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_strided_struct_multiple_to.c b/offload/test/offloading/target_update_strided_struct_multiple_to.c
new file mode 100644
index 0000000000000..1c3646aefabec
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_multiple_to.c
@@ -0,0 +1,136 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update to(s1.data[0:6:2],
+// s2.data[0:4:3]) correctly updates strided sections covering the full arrays
+// from host to device.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define LEN 12
+
+typedef struct {
+ double data[LEN];
+ int len;
+} T;
+
+int main() {
+ T s1, s2;
+ s1.len = LEN;
+ s2.len = LEN;
+
+ // Initialize struct arrays on host with simple sequential values
+ for (int i = 0; i < LEN; i++) {
+ s1.data[i] = i;
+ s2.data[i] = i;
+ }
+
+ printf("original host struct array values:\n");
+ printf("s1.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s1.data[i]);
+ printf("s2.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s2.data[i]);
+ printf("\n");
+
+#pragma omp target data map(tofrom : s1, s2)
+ {
+// Initialize device struct arrays to 20
+#pragma omp target map(tofrom : s1, s2)
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.data[i] = 20.0;
+ s2.data[i] = 20.0;
+ }
+ }
+
+ // s1: even indices (0,2,4,6,8,10)
+ for (int i = 0; i < 6; i++) {
+ s1.data[i * 2] = 10.0;
+ }
+ // s2: every 3rd index (0,3,6,9)
+ for (int i = 0; i < 4; i++) {
+ s2.data[i * 3] = 10.0;
+ }
+
+// s1.data[0:6:2] updates all even indices: 0,2,4,6,8,10 (6 elements, stride 2)
+// s2.data[0:4:3] updates every 3rd: 0,3,6,9 (4 elements, stride 3)
+#pragma omp target update to(s1.data[0 : 6 : 2], s2.data[0 : 4 : 3])
+
+// Verify update on device by adding 5
+#pragma omp target map(tofrom : s1, s2)
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.data[i] += 5.0;
+ s2.data[i] += 5.0;
+ }
+ }
+ }
+
+ printf("device struct array values after update to:\n");
+ printf("s1.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s1.data[i]);
+ printf("s2.data:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s2.data[i]);
+
+ // CHECK: original host struct array values:
+ // CHECK-NEXT: s1.data:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+ // CHECK-NEXT: s2.data:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+
+ // CHECK: device struct array values after update to:
+ // CHECK-NEXT: s1.data:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: s2.data:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_strided_struct_partial_from.c b/offload/test/offloading/target_update_strided_struct_partial_from.c
new file mode 100644
index 0000000000000..d40fc85adc4f7
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_partial_from.c
@@ -0,0 +1,86 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update from(s.data[0:2:3]) correctly
+// updates every third element (stride 3) from the device to the host
+// using a struct with fixed-size array member.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define LEN 11
+
+typedef struct {
+ double data[LEN];
+ size_t len;
+} T;
+
+#pragma omp declare mapper(custom : T v) map(to : v, v.len, v.data[0 : v.len])
+
+int main() {
+ T s;
+ s.len = LEN;
+
+ // Initialize struct data on host
+ for (int i = 0; i < LEN; i++) {
+ s.data[i] = i;
+ }
+
+ printf("original host array values:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+#pragma omp target data map(mapper(custom), tofrom : s)
+ {
+// Execute on device with mapper
+#pragma omp target map(mapper(custom), tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = 20.0; // Set all to 20 on device
+ }
+ }
+
+// Modify specific elements on device (only first 2 stride positions)
+#pragma omp target map(mapper(custom), tofrom : s)
+ {
+ s.data[0] = 10.0;
+ s.data[3] = 10.0;
+ }
+
+// indices 0,3 only
+#pragma omp target update from(s.data[0 : 2 : 3])
+ }
+
+ printf("device array values after update from:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+
+ // CHECK: device array values after update from:
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_strided_struct_partial_to.c b/offload/test/offloading/target_update_strided_struct_partial_to.c
new file mode 100644
index 0000000000000..d0c8adc6b5ea0
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_partial_to.c
@@ -0,0 +1,85 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that #pragma omp target update to(s.data[0:2:3]) correctly
+// updates every third element (stride 3) from the host to the device
+// for struct member arrays.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define LEN 11
+
+typedef struct {
+ double data[LEN];
+ int len;
+} T;
+
+int main() {
+ T s;
+ s.len = LEN;
+
+ // Initialize struct array on host with simple sequential values
+ for (int i = 0; i < LEN; i++)
+ s.data[i] = i;
+
+ printf("original host struct array values:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+#pragma omp target data map(tofrom : s)
+ {
+// Initialize all elements on device to 20
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++)
+ s.data[i] = 20.0;
+ }
+
+ // Modify host struct data for elements that will be updated (set to 10)
+ s.data[0] = 10.0;
+ s.data[3] = 10.0;
+
+// indices 0,3 only
+#pragma omp target update to(s.data[0 : 2 : 3])
+
+// Verify on device by adding 5 to all elements
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++)
+ s.data[i] += 5.0;
+ }
+ }
+
+ printf("device struct array values after update to:\n");
+ for (int i = 0; i < LEN; i++)
+ printf("%.1f\n", s.data[i]);
+
+ // CHECK: original host struct array values:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+
+ // CHECK: device struct array values after update to:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 25.0
+
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_strided_struct_to.c b/offload/test/offloading/target_update_strided_struct_to.c
new file mode 100644
index 0000000000000..b3fe32bbaa345
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_to.c
@@ -0,0 +1,98 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update to" with struct member arrays supports strided
+// sections using fixed-size arrays in structs.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+typedef struct {
+ double data[N];
+ int len;
+} T;
+
+int main() {
+ T s;
+ s.len = N;
+
+ // Initialize struct array on host with simple sequential values
+ for (int i = 0; i < N; i++) {
+ s.data[i] = i;
+ }
+
+ printf("original host struct array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f\n", s.data[i]);
+ printf("\n");
+
+#pragma omp target data map(tofrom : s)
+ {
+// Initialize device struct array to 20
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = 20.0;
+ }
+ }
+
+ // Modify host struct data for strided elements (set to 10)
+ for (int i = 0; i < 8; i++) {
+ s.data[i * 2] = 10.0; // Set even indices to 10
+ }
+
+// indices 0,2,4,6,8,10,12,14
+#pragma omp target update to(s.data[0 : 8 : 2])
+
+// Execute on device - add 5 to verify update worked
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += 5.0;
+ }
+ }
+ }
+
+ printf("after target update to struct:\n");
+ for (int i = 0; i < N; i++)
+ printf("%.1f\n", s.data[i]);
+
+ // CHECK: original host struct array values:
+ // CHECK-NEXT: 0.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 3.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
+ // CHECK-NEXT: 6.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
+ // CHECK-NEXT: 9.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 11.0
+ // CHECK-NEXT: 12.0
+ // CHECK-NEXT: 13.0
+ // CHECK-NEXT: 14.0
+ // CHECK-NEXT: 15.0
+
+ // CHECK: after target update to struct:
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+ // CHECK-NEXT: 15.0
+ // CHECK-NEXT: 25.0
+
+ return 0;
+}
\ No newline at end of file
diff --git a/offload/test/offloading/target_update_to.c b/offload/test/offloading/target_update_to.c
new file mode 100644
index 0000000000000..ecfc3bb69dbf4
--- /dev/null
+++ b/offload/test/offloading/target_update_to.c
@@ -0,0 +1,81 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// This test checks that "update to" clause in OpenMP supports strided sections.
+// #pragma omp target update to(result[0:8:2]) updates every other element
+// (stride 2)
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16
+
+int main() {
+ double *result = (double *)calloc(N, sizeof(double));
+
+ // Initialize on host
+ for (int i = 0; i < N; i++) {
+ result[i] = i;
+ }
+
+ // Initial values
+ printf("original host array values:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+ printf("\n");
+
+#pragma omp target data map(tofrom : result[0 : N])
+ {
+// Update strided elements to device: indices 0,2,4,6
+#pragma omp target update to(result[0 : 8 : 2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++) {
+ result[i] += i;
+ }
+ }
+ }
+
+ printf("from target array results:\n");
+ for (int i = 0; i < N; i++)
+ printf("%f\n", result[i]);
+
+ // CHECK: original host array values:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 1.000000
+ // CHECK-NEXT: 2.000000
+ // CHECK-NEXT: 3.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 5.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 7.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 9.000000
+ // CHECK-NEXT: 10.000000
+ // CHECK-NEXT: 11.000000
+ // CHECK-NEXT: 12.000000
+ // CHECK-NEXT: 13.000000
+ // CHECK-NEXT: 14.000000
+ // CHECK-NEXT: 15.000000
+
+ // CHECK: from target array results:
+ // CHECK-NEXT: 0.000000
+ // CHECK-NEXT: 2.000000
+ // CHECK-NEXT: 4.000000
+ // CHECK-NEXT: 6.000000
+ // CHECK-NEXT: 8.000000
+ // CHECK-NEXT: 10.000000
+ // CHECK-NEXT: 12.000000
+ // CHECK-NEXT: 14.000000
+ // CHECK-NEXT: 16.000000
+ // CHECK-NEXT: 18.000000
+ // CHECK-NEXT: 20.000000
+ // CHECK-NEXT: 22.000000
+ // CHECK-NEXT: 24.000000
+ // CHECK-NEXT: 26.000000
+ // CHECK-NEXT: 28.000000
+ // CHECK-NEXT: 30.000000
+
+ free(result);
+ return 0;
+}
\ No newline at end of file
>From d168c7d7e45fc4861ee29854bcd359dc650f95fc Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Fri, 5 Dec 2025 05:24:38 -0500
Subject: [PATCH 4/8] fixed_extern
---
.../target_update_strided_ptr_messages_from.c | 6 +-----
.../target_update_strided_ptr_messages_to.c | 6 +-----
...pdate_strided_ptr_multiple_messages_from.c | 20 ++++---------------
..._update_strided_ptr_multiple_messages_to.c | 18 +++--------------
...update_strided_ptr_partial_messages_from.c | 6 +-----
...t_update_strided_ptr_partial_messages_to.c | 6 +-----
6 files changed, 11 insertions(+), 51 deletions(-)
diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c
index cd158ff10e4de..c7e8a9efafc9b 100644
--- a/clang/test/OpenMP/target_update_strided_ptr_messages_from.c
+++ b/clang/test/OpenMP/target_update_strided_ptr_messages_from.c
@@ -1,12 +1,9 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-extern void *malloc(__SIZE_TYPE__);
-extern void free(void *);
-
int main(int argc, char **argv) {
int len = 16;
- double *data = (double *)malloc(len * sizeof(double));
+ double *data;
// Valid strided array sections with FROM
#pragma omp target update from(data[0:8:2]) // OK - even indices
@@ -39,6 +36,5 @@ int main(int argc, char **argv) {
#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'}}
{}
- free(data);
return 0;
}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c
index 64839a8384425..47a549c9bc0dc 100644
--- a/clang/test/OpenMP/target_update_strided_ptr_messages_to.c
+++ b/clang/test/OpenMP/target_update_strided_ptr_messages_to.c
@@ -1,12 +1,9 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-extern void *malloc(__SIZE_TYPE__);
-extern void free(void *);
-
int main(int argc, char **argv) {
int len = 16;
- double *data = (double *)malloc(len * sizeof(double));
+ double *data;
// Valid strided array sections with TO
#pragma omp target update to(data[0:8:2]) // OK - even indices
@@ -39,6 +36,5 @@ int main(int argc, char **argv) {
#pragma omp target update to(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'}}
{}
- free(data);
return 0;
}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c
index 2462ae1c840c3..af5dfec1b7941 100644
--- a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c
+++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_from.c
@@ -1,20 +1,11 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-extern void *malloc(__SIZE_TYPE__);
-extern void free(void *);
-
-double *data;
-double *data1;
-double *data2;
-
int main(int argc, char **argv) {
int len = 12;
-
- // Allocate memory for explicit pointers
- data = (double *)malloc(len * sizeof(double));
- data1 = (double *)malloc(len * sizeof(double));
- data2 = (double *)malloc(len * sizeof(double));
+ double *data;
+ double *data1;
+ double *data2;
// Valid multiple strided array sections
#pragma omp target update from(data1[0:6:2], data2[0:4:3]) // OK - different strides
@@ -56,9 +47,6 @@ int main(int argc, char **argv) {
#pragma omp target update from(data1[0:4:2:3], data2[0:3:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
#pragma omp target update from(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
-
- free(data);
- free(data1);
- free(data2);
+
return 0;
}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c
index d8cebefb02606..08c7581598612 100644
--- a/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c
+++ b/clang/test/OpenMP/target_update_strided_ptr_multiple_messages_to.c
@@ -1,20 +1,11 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-extern void *malloc(__SIZE_TYPE__);
-extern void free(void *);
-
-double *data;
-double *data1;
-double *data2;
-
int main(int argc, char **argv) {
int len = 12;
-
- // Allocate memory for explicit pointers
- data = (double *)malloc(len * sizeof(double));
- data1 = (double *)malloc(len * sizeof(double));
- data2 = (double *)malloc(len * sizeof(double));
+ double *data;
+ double *data1;
+ double *data2;
// Valid multiple strided array sections
#pragma omp target update to(data1[0:6:2], data2[0:4:3]) // OK - different strides
@@ -57,8 +48,5 @@ int main(int argc, char **argv) {
#pragma omp target update to(data[0:4:2], data1[0:3:2:1], data2[0:2:3]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
- free(data);
- free(data1);
- free(data2);
return 0;
}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c
index 2f36dfa84cd1b..5937184d1f358 100644
--- a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c
+++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_from.c
@@ -1,12 +1,9 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-extern void *malloc(__SIZE_TYPE__);
-extern void free(void *);
-
int main(int argc, char **argv) {
int len = 11;
- double *data = (double *)malloc(len * sizeof(double));
+ double *data;
// Valid partial strided sections with FROM
#pragma omp target update from(data[0:2:3]) // OK - partial coverage
@@ -36,6 +33,5 @@ int main(int argc, char **argv) {
// Compile-time invalid strides
#pragma omp target update from(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
- free(data);
return 0;
}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c
index ce4b315bec1ae..848ef95509f1c 100644
--- a/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c
+++ b/clang/test/OpenMP/target_update_strided_ptr_partial_messages_to.c
@@ -1,12 +1,9 @@
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
-extern void *malloc(__SIZE_TYPE__);
-extern void free(void *);
-
int main(int argc, char **argv) {
int len = 11;
- double *data = (double *)malloc(len * sizeof(double));
+ double *data;
// Valid partial strided sections with TO
#pragma omp target update to(data[0:2:3]) // OK - partial coverage
@@ -35,6 +32,5 @@ int main(int argc, char **argv) {
// Compile-time invalid strides
#pragma omp target update to(data[1:2:-3]) // expected-error {{section stride is evaluated to a non-positive value -3}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
- free(data);
return 0;
}
\ No newline at end of file
>From b89534d6b73206d7ba9c2add7ceacbb18c1f396d Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 8 Dec 2025 05:24:04 -0500
Subject: [PATCH 5/8] codegen-fix
---
clang/test/OpenMP/target_update_codegen.cpp | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index b0a6c8b2f31a2..633784edd7bb4 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -1231,9 +1231,9 @@ struct ST {
// CK21: store i64 1, ptr [[COUNT_4]],
// CK21: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
// CK21: store i64 {{4|8}}, ptr [[STRIDE_4]],
- // CK21-DAG: [[GEPBP:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
- // CK21-DAG: [[GEPP:%.+]] = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i32 0, i32 0
- // CK21-DAG: [[GEPSZ:%.+]] = getelementptr inbounds [2 x i64], ptr %.offload_sizes, i32 0, i32 0
+ // CK21-DAG: [[GEPBP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_BASEPTRS:%.+]], i32 0, i32 0
+ // CK21-DAG: [[GEPP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_PTRS:%.+]], i32 0, i32 0
+ // CK21-DAG: [[GEPSZ:%.+]] = getelementptr inbounds [2 x i64], ptr [[OFFLOAD_SIZES:%.+]], i32 0, i32 0
// CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP]], ptr [[GEPP]], ptr [[GEPSZ]], ptr @.offload_maptypes, ptr null, ptr null)
// CK21: ret void
#pragma omp target update to(dptr[0:2][1:3][0:4])
>From 969c2804cf57ae09fd1fb675588577558c228d8c Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 16 Dec 2025 02:03:43 -0500
Subject: [PATCH 6/8] merge-conflict-changes
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 11 +++++++----
1 file changed, 7 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 01cc6cc574644..33f95348aca20 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9116,10 +9116,12 @@ class MappableExprsHandler {
// * It preserves the correct rank so targetDataUpdate() computes
// DimSize == 2 for cases like strided array sections originating
// from user-defined mappers (e.g. test with s.data[0:8:2]).
- UnionCurInfo.NonContigInfo.Dims.insert(
- UnionCurInfo.NonContigInfo.Dims.begin(), 1);
- emitCombinedEntry(CombinedInfo, UnionCurInfo.Types, PartialStruct,
- /*IsMapThis*/ !VD, OMPBuilder, VD);
+ GroupUnionCurInfo.NonContigInfo.Dims.insert(
+ GroupUnionCurInfo.NonContigInfo.Dims.begin(), 1);
+ emitCombinedEntry(CurInfo, GroupUnionCurInfo.Types, PartialStruct,
+ AttachInfo, /*IsMapThis=*/!VD, OMPBuilder, VD,
+ /*OffsetForMemberOfFlag=*/0,
+ /*NotTargetParams=*/true);
}
// Append this group's results to the overall CurInfo in the correct
@@ -9127,6 +9129,7 @@ class MappableExprsHandler {
CurInfo.append(GroupUnionCurInfo);
if (AttachInfo.isValid())
emitAttachEntry(CGF, CurInfo, AttachInfo);
+ }
// We need to append the results of this capture to what we already have.
CombinedInfo.append(CurInfo);
>From ae5ede8223e1f7b724611040184f87acdcb998bb Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Tue, 16 Dec 2025 06:06:36 -0500
Subject: [PATCH 7/8] openmp-codegen-fix
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 9 +++++----
1 file changed, 5 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 33f95348aca20..779aa230a91d9 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9118,10 +9118,11 @@ class MappableExprsHandler {
// from user-defined mappers (e.g. test with s.data[0:8:2]).
GroupUnionCurInfo.NonContigInfo.Dims.insert(
GroupUnionCurInfo.NonContigInfo.Dims.begin(), 1);
- emitCombinedEntry(CurInfo, GroupUnionCurInfo.Types, PartialStruct,
- AttachInfo, /*IsMapThis=*/!VD, OMPBuilder, VD,
- /*OffsetForMemberOfFlag=*/0,
- /*NotTargetParams=*/true);
+ emitCombinedEntry(
+ CurInfo, GroupUnionCurInfo.Types, PartialStruct, AttachInfo,
+ /*IsMapThis=*/!VD, OMPBuilder, VD,
+ /*OffsetForMemberOfFlag=*/CombinedInfo.BasePointers.size(),
+ /*NotTargetParams=*/true);
}
// Append this group's results to the overall CurInfo in the correct
>From 716bc4f28782e0ace141315498db67f7efccce5a Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 22 Dec 2025 10:43:27 -0500
Subject: [PATCH 8/8] testcases-updated
---
clang/test/OpenMP/target_update_codegen.cpp | 9 ++--
.../strided_ptr_multiple_update_from.c | 2 +-
.../strided_ptr_multiple_update_to.c | 2 +-
.../strided_ptr_partial_update_from.c | 2 +-
.../strided_ptr_partial_update_to.c | 2 +-
offload/test/offloading/target_update_from.c | 3 +-
.../target_update_strided_struct_from.c | 4 +-
...rget_update_strided_struct_multiple_from.c | 32 ++++++------
...target_update_strided_struct_multiple_to.c | 2 +-
...arget_update_strided_struct_partial_from.c | 4 +-
.../target_update_strided_struct_partial_to.c | 2 +-
.../target_update_strided_struct_to.c | 2 +-
offload/test/offloading/target_update_to.c | 51 ++++++++++++-------
13 files changed, 66 insertions(+), 51 deletions(-)
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index 633784edd7bb4..df238f97f608c 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -1231,10 +1231,11 @@ struct ST {
// CK21: store i64 1, ptr [[COUNT_4]],
// CK21: [[STRIDE_4:%.+]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR]], ptr [[DIM_4]], {{.+}} 0, {{.+}} 2
// CK21: store i64 {{4|8}}, ptr [[STRIDE_4]],
- // CK21-DAG: [[GEPBP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_BASEPTRS:%.+]], i32 0, i32 0
- // CK21-DAG: [[GEPP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOAD_PTRS:%.+]], i32 0, i32 0
- // CK21-DAG: [[GEPSZ:%.+]] = getelementptr inbounds [2 x i64], ptr [[OFFLOAD_SIZES:%.+]], i32 0, i32 0
- // CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP]], ptr [[GEPP]], ptr [[GEPSZ]], ptr @.offload_maptypes, ptr null, ptr null)
+ // CK21-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPSZ:%.+]], ptr [[MTYPE]]{{.+}})
+ // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+ // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK21-DAG: store ptr [[DIMS]], ptr [[PTRS1:%.+]],
+ // CK21-DAG: [[PTRS1]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1
// CK21: ret void
#pragma omp target update to(dptr[0:2][1:3][0:4])
}
diff --git a/offload/test/offloading/strided_ptr_multiple_update_from.c b/offload/test/offloading/strided_ptr_multiple_update_from.c
index afb2f1fc8dd53..b671a1e2b688d 100644
--- a/offload/test/offloading/strided_ptr_multiple_update_from.c
+++ b/offload/test/offloading/strided_ptr_multiple_update_from.c
@@ -111,4 +111,4 @@ int main() {
free(data1);
free(data2);
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/strided_ptr_multiple_update_to.c b/offload/test/offloading/strided_ptr_multiple_update_to.c
index 2005117c83754..7ae9189ead8e3 100644
--- a/offload/test/offloading/strided_ptr_multiple_update_to.c
+++ b/offload/test/offloading/strided_ptr_multiple_update_to.c
@@ -127,4 +127,4 @@ int main() {
free(data1);
free(data2);
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/strided_ptr_partial_update_from.c b/offload/test/offloading/strided_ptr_partial_update_from.c
index e2f4398b2fbd7..d30a671393002 100644
--- a/offload/test/offloading/strided_ptr_partial_update_from.c
+++ b/offload/test/offloading/strided_ptr_partial_update_from.c
@@ -65,4 +65,4 @@ int main() {
free(data);
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/strided_ptr_partial_update_to.c b/offload/test/offloading/strided_ptr_partial_update_to.c
index 28bd27fd04366..37c165351a760 100644
--- a/offload/test/offloading/strided_ptr_partial_update_to.c
+++ b/offload/test/offloading/strided_ptr_partial_update_to.c
@@ -78,4 +78,4 @@ int main() {
free(data);
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_from.c b/offload/test/offloading/target_update_from.c
index 089b498ada0db..4f84968f6d9a9 100644
--- a/offload/test/offloading/target_update_from.c
+++ b/offload/test/offloading/target_update_from.c
@@ -32,6 +32,7 @@ int main() {
printf("\n");
// Expected: even indices i, odd indices 0
+ // CHECK: after target update from
// CHECK: 0.000000
// CHECK: 0.000000
// CHECK: 2.000000
@@ -70,4 +71,4 @@ int main() {
free(result);
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_strided_struct_from.c b/offload/test/offloading/target_update_strided_struct_from.c
index 1c5759b1864c5..46448d10189f2 100644
--- a/offload/test/offloading/target_update_strided_struct_from.c
+++ b/offload/test/offloading/target_update_strided_struct_from.c
@@ -28,7 +28,7 @@ int main() {
printf("%f\n", s.data[i]);
printf("\n");
-#pragma omp target data map(mapper(custom), tofrom : s)
+#pragma omp target data map(mapper(custom), to : s)
{
// Execute on device with explicit mapper
#pragma omp target map(mapper(custom), tofrom : s)
@@ -83,4 +83,4 @@ int main() {
// CHECK-NEXT: 15.000000
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_strided_struct_multiple_from.c b/offload/test/offloading/target_update_strided_struct_multiple_from.c
index 95faf104826d5..fa57d35b86a4a 100644
--- a/offload/test/offloading/target_update_strided_struct_multiple_from.c
+++ b/offload/test/offloading/target_update_strided_struct_multiple_from.c
@@ -33,7 +33,7 @@ int main() {
for (int i = 0; i < LEN; i++)
printf("%.1f\n", s2.data[i]);
-#pragma omp target data map(tofrom : s1, s2)
+#pragma omp target data map(to : s1, s2)
{
// Initialize all device values to 20
#pragma omp target map(tofrom : s1, s2)
@@ -101,30 +101,30 @@ int main() {
// CHECK: host struct array values after update from:
// CHECK-NEXT: s1.data:
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 1.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 3.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 5.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 7.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 9.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 11.0
// CHECK-NEXT: s2.data:
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 1.0
+ // CHECK-NEXT: 2.0
+ // CHECK-NEXT: 10.0
+ // CHECK-NEXT: 4.0
+ // CHECK-NEXT: 5.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 7.0
+ // CHECK-NEXT: 8.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
- // CHECK-NEXT: 20.0
// CHECK-NEXT: 10.0
- // CHECK-NEXT: 20.0
- // CHECK-NEXT: 20.0
+ // CHECK-NEXT: 11.0
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_strided_struct_multiple_to.c b/offload/test/offloading/target_update_strided_struct_multiple_to.c
index 1c3646aefabec..69c4015153580 100644
--- a/offload/test/offloading/target_update_strided_struct_multiple_to.c
+++ b/offload/test/offloading/target_update_strided_struct_multiple_to.c
@@ -133,4 +133,4 @@ int main() {
// CHECK-NEXT: 25.0
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_strided_struct_partial_from.c b/offload/test/offloading/target_update_strided_struct_partial_from.c
index d40fc85adc4f7..c0ffc72f900a6 100644
--- a/offload/test/offloading/target_update_strided_struct_partial_from.c
+++ b/offload/test/offloading/target_update_strided_struct_partial_from.c
@@ -30,7 +30,7 @@ int main() {
printf("%.1f\n", s.data[i]);
printf("\n");
-#pragma omp target data map(mapper(custom), tofrom : s)
+#pragma omp target data map(mapper(custom), to : s)
{
// Execute on device with mapper
#pragma omp target map(mapper(custom), tofrom : s)
@@ -83,4 +83,4 @@ int main() {
// CHECK-NEXT: 10.0
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_strided_struct_partial_to.c b/offload/test/offloading/target_update_strided_struct_partial_to.c
index d0c8adc6b5ea0..018dcac7cdcf8 100644
--- a/offload/test/offloading/target_update_strided_struct_partial_to.c
+++ b/offload/test/offloading/target_update_strided_struct_partial_to.c
@@ -82,4 +82,4 @@ int main() {
// CHECK-NEXT: 25.0
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_strided_struct_to.c b/offload/test/offloading/target_update_strided_struct_to.c
index b3fe32bbaa345..90b3a5fe26989 100644
--- a/offload/test/offloading/target_update_strided_struct_to.c
+++ b/offload/test/offloading/target_update_strided_struct_to.c
@@ -95,4 +95,4 @@ int main() {
// CHECK-NEXT: 25.0
return 0;
-}
\ No newline at end of file
+}
diff --git a/offload/test/offloading/target_update_to.c b/offload/test/offloading/target_update_to.c
index ecfc3bb69dbf4..d45d3182a166d 100644
--- a/offload/test/offloading/target_update_to.c
+++ b/offload/test/offloading/target_update_to.c
@@ -23,9 +23,22 @@ int main() {
printf("%f\n", result[i]);
printf("\n");
-#pragma omp target data map(tofrom : result[0 : N])
+#pragma omp target data map(from : result[0 : N])
{
-// Update strided elements to device: indices 0,2,4,6
+// Initialize device array to 0
+#pragma omp target
+ {
+ for (int i = 0; i < N; i++) {
+ result[i] = 0.0;
+ }
+ }
+
+ // Modify host strided elements (even indices)
+ for (int i = 0; i < 8; i++) {
+ result[i * 2] = 100.0;
+ }
+
+// Update strided elements to device: indices 0,2,4,6,8,10,12,14
#pragma omp target update to(result[0 : 8 : 2])
#pragma omp target
@@ -59,23 +72,23 @@ int main() {
// CHECK-NEXT: 15.000000
// CHECK: from target array results:
- // CHECK-NEXT: 0.000000
- // CHECK-NEXT: 2.000000
- // CHECK-NEXT: 4.000000
- // CHECK-NEXT: 6.000000
- // CHECK-NEXT: 8.000000
- // CHECK-NEXT: 10.000000
- // CHECK-NEXT: 12.000000
- // CHECK-NEXT: 14.000000
- // CHECK-NEXT: 16.000000
- // CHECK-NEXT: 18.000000
- // CHECK-NEXT: 20.000000
- // CHECK-NEXT: 22.000000
- // CHECK-NEXT: 24.000000
- // CHECK-NEXT: 26.000000
- // CHECK-NEXT: 28.000000
- // CHECK-NEXT: 30.000000
+ // CHECK-NEXT: 100.000000
+ // CHECK-NEXT: 1.000000
+ // CHECK-NEXT: 102.000000
+ // CHECK-NEXT: 3.000000
+ // CHECK-NEXT: 104.000000
+ // CHECK-NEXT: 5.000000
+ // CHECK-NEXT: 106.000000
+ // CHECK-NEXT: 7.000000
+ // CHECK-NEXT: 108.000000
+ // CHECK-NEXT: 9.000000
+ // CHECK-NEXT: 110.000000
+ // CHECK-NEXT: 11.000000
+ // CHECK-NEXT: 112.000000
+ // CHECK-NEXT: 13.000000
+ // CHECK-NEXT: 114.000000
+ // CHECK-NEXT: 15.000000
free(result);
return 0;
-}
\ No newline at end of file
+}
More information about the llvm-commits
mailing list