[clang] [llvm] [Clang][OpenMP] Support expression semantics in target update fields with non-contiguous array sections (PR #176708)
Amit Tiwari via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 19 00:59:56 PST 2026
https://github.com/amitamd7 updated https://github.com/llvm/llvm-project/pull/176708
>From 87d7113b3e9979bc31b0ba7d959652aa3b57e700 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Thu, 4 Dec 2025 05:05:38 -0500
Subject: [PATCH 1/4] expression_semantics_patch
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 29 ++++++++++++++++-------
1 file changed, 21 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 8d7a207a91f5a..578e63351ca5f 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9773,16 +9773,29 @@ Error OpenMPIRBuilder::emitOffloadingArrays(
ConstantInt::get(Int64Ty, 0));
SmallBitVector RuntimeSizes(CombinedInfo.Sizes.size());
for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) {
+ bool IsNonContigEntry =
+ IsNonContiguous &&
+ (static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ CombinedInfo.Types[I] &
+ OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0);
+ // For NON_CONTIG entries ArgSizes must carry the dimension count
+ // (number of descriptor_dim records) – NOT the byte size expression.
+ // Variable subsection forms (e.g. 0:s.len/2:2) previously produced a
+ // non-constant size so we marked them runtime and stored the byte size,
+ // leading the runtime to treat it as DimSize and overrun descriptors.
+ if (IsNonContigEntry) {
+ // Dims must be long enough and positive.
+ assert(I < CombinedInfo.NonContigInfo.Dims.size() &&
+ "Induction variable is in-bounds with the NON_CONTIG Dims array");
+ const uint64_t DimCount = CombinedInfo.NonContigInfo.Dims[I];
+ assert(DimCount > 0 && "NON_CONTIG DimCount must be > 0");
+ ConstSizes[I] =
+ ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]);
+ continue;
+ }
if (auto *CI = dyn_cast<Constant>(CombinedInfo.Sizes[I])) {
if (!isa<ConstantExpr>(CI) && !isa<GlobalValue>(CI)) {
- if (IsNonContiguous &&
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- CombinedInfo.Types[I] &
- OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG))
- ConstSizes[I] =
- ConstantInt::get(Int64Ty, CombinedInfo.NonContigInfo.Dims[I]);
- else
- ConstSizes[I] = CI;
+ ConstSizes[I] = CI;
continue;
}
}
>From a9dcaede9ba9b50e1e281ecdeaba16e400d17ecb Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Thu, 4 Dec 2025 06:34:09 -0500
Subject: [PATCH 2/4] variable_stride_fix
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 16 ++++++++++++++++
1 file changed, 16 insertions(+)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 8981a0de6d0e4..027d1fb26bc97 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8033,12 +8033,28 @@ class MappableExprsHandler {
const Expr *StrideExpr = OASE->getStride();
if (!StrideExpr)
return false;
+
+ assert(StrideExpr->getType()->isIntegerType() &&
+ "Stride expression must be of integer type");
+
+ // If the stride is a variable (not a constant), it's non-contiguous.
+ const Expr *S = StrideExpr->IgnoreParenImpCasts();
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(S)) {
+ if (isa<VarDecl>(DRE->getDecl()) ||
+ isa<ParmVarDecl>(DRE->getDecl()))
+ return true;
+ }
+ if (isa<MemberExpr>(S) || isa<ArraySubscriptExpr>(S))
+ return true;
+ // If stride is not evaluatable as a constant, treat as
+ // non-contiguous.
const auto Constant =
StrideExpr->getIntegerConstantExpr(CGF.getContext());
if (!Constant)
return false;
+ // Treat non-unitary strides as non-contiguous.
return !Constant->isOne();
});
>From 74dcd62d5836cbb6820e3e35a79bf92203afe7c4 Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 19 Jan 2026 03:25:48 -0500
Subject: [PATCH 3/4] cleanup
---
...d_ptr_variable_count_and_stride_messages.c | 62 ++++
...date_strided_ptr_variable_count_messages.c | 57 ++++
...ate_strided_ptr_variable_stride_messages.c | 64 ++++
...truct_variable_count_and_stride_messages.c | 72 +++++
...pdate_variable_count_and_stride_messages.c | 85 +++++
.../strided_update_count_expression.c | 133 ++++++++
.../strided_update_count_expression_complex.c | 289 +++++++++++++++++
.../strided_update_count_expression_misc.c | 99 ++++++
..._update_multiple_arrays_count_expression.c | 161 ++++++++++
...d_update_multiple_arrays_variable_stride.c | 145 +++++++++
...strided_update_variable_count_and_stride.c | 136 ++++++++
.../strided_update_variable_stride.c | 135 ++++++++
.../strided_update_variable_stride_complex.c | 293 ++++++++++++++++++
.../strided_update_variable_stride_misc.c | 94 ++++++
.../target_update_ptr_count_expression.c | 99 ++++++
...get_update_ptr_variable_count_and_stride.c | 94 ++++++
.../target_update_ptr_variable_stride.c | 95 ++++++
...t_update_strided_struct_count_expression.c | 97 ++++++
...strided_struct_variable_count_and_stride.c | 96 ++++++
...et_update_strided_struct_variable_stride.c | 95 ++++++
20 files changed, 2401 insertions(+)
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c
create mode 100644 clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c
create mode 100644 clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c
create mode 100644 clang/test/OpenMP/target_update_variable_count_and_stride_messages.c
create mode 100644 offload/test/offloading/strided_update_count_expression.c
create mode 100644 offload/test/offloading/strided_update_count_expression_complex.c
create mode 100644 offload/test/offloading/strided_update_count_expression_misc.c
create mode 100644 offload/test/offloading/strided_update_multiple_arrays_count_expression.c
create mode 100644 offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
create mode 100644 offload/test/offloading/strided_update_variable_count_and_stride.c
create mode 100644 offload/test/offloading/strided_update_variable_stride.c
create mode 100644 offload/test/offloading/strided_update_variable_stride_complex.c
create mode 100644 offload/test/offloading/strided_update_variable_stride_misc.c
create mode 100644 offload/test/offloading/target_update_ptr_count_expression.c
create mode 100644 offload/test/offloading/target_update_ptr_variable_count_and_stride.c
create mode 100644 offload/test/offloading/target_update_ptr_variable_stride.c
create mode 100644 offload/test/offloading/target_update_strided_struct_count_expression.c
create mode 100644 offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c
create mode 100644 offload/test/offloading/target_update_strided_struct_variable_stride.c
diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c
new file mode 100644
index 0000000000000..932cd6b1c97bb
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_and_stride_messages.c
@@ -0,0 +1,62 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int count = 8;
+ int stride = 2;
+ int stride_large = 5;
+ double *data;
+
+ // Valid strided array sections with both variable count and variable stride (FROM)
+ #pragma omp target update from(data[0:count:stride]) // OK - both variable
+ {}
+
+ #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride
+ {}
+
+ #pragma omp target update from(data[0:count:stride_large]) // OK - variable count, different stride
+ {}
+
+ #pragma omp target update from(data[1:len-2:stride]) // OK - with offset, count expression
+ {}
+
+ #pragma omp target update from(data[0:count/2:stride*2]) // OK - both expressions
+ {}
+
+ #pragma omp target update from(data[0:(len+1)/2:stride+1]) // OK - complex expressions
+ {}
+
+ #pragma omp target update from(data[2:count-2:len/4]) // OK - all expressions
+ {}
+
+ // Edge cases
+ int stride_one = 1;
+ #pragma omp target update from(data[0:count:stride_one]) // OK - variable count, stride=1
+ {}
+
+ #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride
+ {}
+
+ // Invalid compile-time constant strides with variable count
+ #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'}}
+
+ #pragma omp target update from(data[0:len/2:-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:count:-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'}}
+
+ // Valid strided array sections with variable count and stride (TO)
+ #pragma omp target update to(data[0:count:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/2:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:count:stride*2]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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;
+}
diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c
new file mode 100644
index 0000000000000..23fba9c8bc84f
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_variable_count_messages.c
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int count = 8;
+ int divisor = 2;
+ double *data;
+
+ // Valid strided array sections with variable count expressions (FROM)
+ #pragma omp target update from(data[0:count:2]) // OK - variable count
+ {}
+
+ #pragma omp target update from(data[0:len/2:2]) // OK - count expression
+ {}
+
+ #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction
+ {}
+
+ #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression
+ {}
+
+ #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication
+ {}
+
+ #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo
+ {}
+
+ // Variable count with stride = 1 (contiguous)
+ #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride
+ {}
+
+ #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride
+ {}
+
+ // Invalid stride expressions with variable count
+ #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'}}
+
+ #pragma omp target update from(data[0:len/2:-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:count:-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'}}
+
+ // Valid strided array sections with variable count expressions (TO)
+ #pragma omp target update to(data[0:count:2]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/2:2]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len-4:3]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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;
+}
diff --git a/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.c
new file mode 100644
index 0000000000000..3f85ed0c48d66
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_ptr_variable_stride_messages.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
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int stride = 2;
+ int stride_large = 5;
+ double *data;
+
+ // Valid strided array sections with variable stride (FROM)
+ #pragma omp target update from(data[0:8:stride]) // OK - variable stride
+ {}
+
+ #pragma omp target update from(data[0:4:stride_large]) // OK - different variable stride
+ {}
+
+ #pragma omp target update from(data[1:6:stride]) // OK - with offset
+ {}
+
+ #pragma omp target update from(data[0:5:stride+1]) // OK - stride expression
+ {}
+
+ #pragma omp target update from(data[0:4:stride*2]) // OK - stride multiplication
+ {}
+
+ #pragma omp target update from(data[2:3:len/4]) // OK - stride from expression
+ {}
+
+ // Edge case: stride = 1 (should be contiguous, not non-contiguous)
+ int stride_one = 1;
+ #pragma omp target update from(data[0:8:stride_one]) // OK - stride=1 is contiguous
+ {}
+
+ // Invalid variable stride expressions
+ int zero_stride = 0;
+ int neg_stride = -1;
+
+ // Note: These are runtime checks, so no compile-time error
+ #pragma omp target update from(data[0:8:zero_stride]) // OK at compile-time (runtime will fail)
+ {}
+
+ #pragma omp target update from(data[0:4:neg_stride]) // OK at compile-time (runtime will fail)
+ {}
+
+ // Compile-time constant invalid strides
+ #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'}}
+
+ // Valid strided array sections with variable stride (TO)
+ #pragma omp target update to(data[0:8:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:5:stride+1]) // OK
+ {}
+
+ #pragma omp target update to(data[0:4:stride*2]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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'}}
+
+ return 0;
+}
diff --git a/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c
new file mode 100644
index 0000000000000..70775d5c8322c
--- /dev/null
+++ b/clang/test/OpenMP/target_update_strided_struct_variable_count_and_stride_messages.c
@@ -0,0 +1,72 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+#define N 20
+typedef struct {
+ double data[N];
+ int len;
+ int stride;
+} T;
+
+int main(int argc, char **argv) {
+ T s;
+ s.len = 16;
+ s.stride = 2;
+ int count = 8;
+ int ext_stride = 3;
+
+ // Valid strided struct member array sections with variable count/stride (FROM)
+ #pragma omp target update from(s.data[0:s.len/2:2]) // OK - member count expression
+ {}
+
+ #pragma omp target update from(s.data[0:count:s.stride]) // OK - external count, member stride
+ {}
+
+ #pragma omp target update from(s.data[0:s.len:ext_stride]) // OK - member count, external stride
+ {}
+
+ #pragma omp target update from(s.data[0:count:ext_stride]) // OK - both external
+ {}
+
+ #pragma omp target update from(s.data[0:s.len/2:s.stride]) // OK - both from struct
+ {}
+
+ #pragma omp target update from(s.data[1:(s.len-2)/2:s.stride]) // OK - complex count expression
+ {}
+
+ #pragma omp target update from(s.data[0:count*2:s.stride+1]) // OK - expressions for both
+ {}
+
+ // Edge cases
+ int stride_one = 1;
+ #pragma omp target update from(s.data[0:s.len:stride_one]) // OK - stride=1
+ {}
+
+ #pragma omp target update from(s.data[0:s.len/s.stride:s.stride]) // OK - count depends on stride
+ {}
+
+ // Invalid compile-time constant strides with variable count
+ #pragma omp target update from(s.data[0:s.len: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:count:-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(s.data[1:s.len/2:-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'}}
+
+ // Valid strided struct member array sections with variable count and stride (TO)
+ #pragma omp target update to(s.data[0:s.len/2:2]) // OK
+ {}
+
+ #pragma omp target update to(s.data[0:count:s.stride]) // OK
+ {}
+
+ #pragma omp target update to(s.data[0:s.len:ext_stride]) // OK
+ {}
+
+ #pragma omp target update to(s.data[0:count*2:s.stride+1]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(s.data[0:s.len: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;
+}
diff --git a/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c
new file mode 100644
index 0000000000000..0082539538a32
--- /dev/null
+++ b/clang/test/OpenMP/target_update_variable_count_and_stride_messages.c
@@ -0,0 +1,85 @@
+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
+
+int main(int argc, char **argv) {
+ int len = 16;
+ int count = 8;
+ int stride = 2;
+ int divisor = 2;
+ double data[100];
+
+ // Valid strided array sections with variable count expressions (FROM)
+ #pragma omp target update from(data[0:count:2]) // OK - variable count
+ {}
+
+ #pragma omp target update from(data[0:len/2:2]) // OK - count expression
+ {}
+
+ #pragma omp target update from(data[0:len-4:3]) // OK - count with subtraction
+ {}
+
+ #pragma omp target update from(data[1:(len+1)/2:2]) // OK - complex count expression
+ {}
+
+ #pragma omp target update from(data[0:count*2:3]) // OK - count multiplication
+ {}
+
+ #pragma omp target update from(data[2:len%divisor:2]) // OK - count with modulo
+ {}
+
+ // Variable stride with constant/variable count
+ #pragma omp target update from(data[0:10:stride]) // OK - constant count, variable stride
+ {}
+
+ #pragma omp target update from(data[0:count:stride]) // OK - both variable
+ {}
+
+ #pragma omp target update from(data[0:len/2:stride]) // OK - count expression, variable stride
+ {}
+
+ #pragma omp target update from(data[0:count:stride*2]) // OK - variable count, stride expression
+ {}
+
+ #pragma omp target update from(data[0:len/divisor:stride+1]) // OK - both expressions
+ {}
+
+ // Variable count with stride = 1 (contiguous)
+ #pragma omp target update from(data[0:count]) // OK - variable count, implicit stride
+ {}
+
+ #pragma omp target update from(data[0:len/divisor]) // OK - expression count, implicit stride
+ {}
+
+ // Edge cases
+ int stride_one = 1;
+ #pragma omp target update from(data[0:len:stride_one]) // OK - stride=1 variable
+ {}
+
+ #pragma omp target update from(data[0:len/stride:stride]) // OK - count depends on stride
+ {}
+
+ // Invalid stride expressions with variable count
+ #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'}}
+
+ #pragma omp target update from(data[0:len/2:-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:count:-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'}}
+
+ // Valid strided array sections with variable count expressions (TO)
+ #pragma omp target update to(data[0:count:2]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/2:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:count:stride]) // OK
+ {}
+
+ #pragma omp target update to(data[0:len/divisor:stride+1]) // OK
+ {}
+
+ // Invalid stride with TO
+ #pragma omp target update to(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;
+}
diff --git a/offload/test/offloading/strided_update_count_expression.c b/offload/test/offloading/strided_update_count_expression.c
new file mode 100644
index 0000000000000..a87da289a9154
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression.c
@@ -0,0 +1,133 @@
+// This test checks that "update from" and "update to" clauses in OpenMP are
+// supported when elements are updated in a non-contiguous manner with variable
+// count expression. Tests #pragma omp target update from/to(data[0:len/2:2])
+// where the count (len/2) is a variable expression, not a constant.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 10;
+ double data[len];
+
+ // ====================================================================
+ // TEST 1: Update FROM device (device -> host)
+ // ====================================================================
+
+#pragma omp target map(tofrom : len, data[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ }
+ }
+
+ printf("Test 1: Update FROM device\n");
+ printf("original host array values:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+#pragma omp target data map(to : len, data[0 : len])
+ {
+#pragma omp target
+ for (int i = 0; i < len; i++) {
+ data[i] += i;
+ }
+
+#pragma omp target update from(data[0 : len / 2 : 2])
+ }
+
+ printf("from target array results:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+ // ====================================================================
+ // TEST 2: Update TO device (host -> device)
+ // ====================================================================
+
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ }
+
+ printf("\nTest 2: Update TO device\n");
+ printf("original host array values:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+#pragma omp target data map(tofrom : len, data[0 : len])
+ {
+#pragma omp target
+ for (int i = 0; i < len; i++) {
+ data[i] = 20.0;
+ }
+
+ data[0] = 10.0;
+ data[2] = 10.0;
+ data[4] = 10.0;
+ data[6] = 10.0;
+ data[8] = 10.0;
+
+#pragma omp target update to(data[0 : len / 2 : 2])
+
+#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]);
+
+ return 0;
+}
+
+// CHECK: Test 1: Update FROM device
+// 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: 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: Test 2: Update TO device
+// 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: device array values after update to:
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 25.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 25.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 25.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 25.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 25.000000
diff --git a/offload/test/offloading/strided_update_count_expression_complex.c b/offload/test/offloading/strided_update_count_expression_complex.c
new file mode 100644
index 0000000000000..f9beef513da24
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression_complex.c
@@ -0,0 +1,289 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests non-contiguous array sections with complex expression-based count
+// scenarios including multiple struct arrays and non-zero offset.
+
+#include <omp.h>
+#include <stdio.h>
+
+struct Data {
+ int offset;
+ int len;
+ double arr[20];
+};
+
+int main() {
+ struct Data s1, s2;
+
+ // Test 1: Multiple arrays with different count expressions
+ s1.len = 10;
+ s2.len = 10;
+
+ // Initialize on device
+#pragma omp target map(tofrom : s1, s2)
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] = i;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] = i * 10;
+ }
+ }
+
+ // Test FROM: Update multiple struct arrays with complex count expressions
+#pragma omp target data map(to : s1, s2)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] += i;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] += i * 10;
+ }
+ }
+
+ // Complex count: (len-2)/2 and len*2/5
+#pragma omp target update from(s1.arr[0 : (s1.len - 2) / 2 : 2], \
+ s2.arr[0 : s2.len * 2 / 5 : 2])
+ }
+
+ printf("Test 1 - complex count expressions (from):\n");
+ printf("s1 results:\n");
+ for (int i = 0; i < s1.len; i++)
+ printf("%f\n", s1.arr[i]);
+
+ printf("s2 results:\n");
+ for (int i = 0; i < s2.len; i++)
+ printf("%f\n", s2.arr[i]);
+
+ // Reset for TO test
+#pragma omp target map(tofrom : s1, s2)
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] = i * 2;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] = i * 20;
+ }
+ }
+
+ // Modify host data
+ for (int i = 0; i < (s1.len - 2) / 2; i++) {
+ s1.arr[i * 2] = i + 100;
+ }
+ for (int i = 0; i < s2.len * 2 / 5; i++) {
+ s2.arr[i * 2] = i + 50;
+ }
+
+ // Test TO: Update with complex count expressions
+#pragma omp target data map(to : s1, s2)
+ {
+#pragma omp target update to(s1.arr[0 : (s1.len - 2) / 2 : 2], \
+ s2.arr[0 : s2.len * 2 / 5 : 2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] += 100;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] += 100;
+ }
+ }
+ }
+
+ printf("Test 1 - complex count expressions (to):\n");
+ printf("s1 results:\n");
+ for (int i = 0; i < s1.len; i++)
+ printf("%f\n", s1.arr[i]);
+
+ printf("s2 results:\n");
+ for (int i = 0; i < s2.len; i++)
+ printf("%f\n", s2.arr[i]);
+
+ // Test 2: Complex count with non-zero offset
+ s1.offset = 2;
+ s1.len = 10;
+ s2.offset = 1;
+ s2.len = 10;
+
+ // Initialize on device
+#pragma omp target map(tofrom : s1, s2)
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] = i;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] = i * 10;
+ }
+ }
+
+ // Test FROM: Complex count with offset
+#pragma omp target data map(to : s1, s2)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] += i;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] += i * 10;
+ }
+ }
+
+ // Count: (len-offset)/2 with stride 2
+#pragma omp target update from( \
+ s1.arr[s1.offset : (s1.len - s1.offset) / 2 : 2], \
+ s2.arr[s2.offset : (s2.len - s2.offset) / 2 : 2])
+ }
+
+ printf("Test 2 - complex count with offset (from):\n");
+ printf("s1 results:\n");
+ for (int i = 0; i < s1.len; i++)
+ printf("%f\n", s1.arr[i]);
+
+ printf("s2 results:\n");
+ for (int i = 0; i < s2.len; i++)
+ printf("%f\n", s2.arr[i]);
+
+ // Reset for TO test
+#pragma omp target map(tofrom : s1, s2)
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] = i * 2;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] = i * 20;
+ }
+ }
+
+ // Modify host data
+ for (int i = 0; i < (s1.len - s1.offset) / 2; i++) {
+ s1.arr[s1.offset + i * 2] = i + 100;
+ }
+ for (int i = 0; i < (s2.len - s2.offset) / 2; i++) {
+ s2.arr[s2.offset + i * 2] = i + 50;
+ }
+
+ // Test TO: Update with complex count and offset
+#pragma omp target data map(to : s1, s2)
+ {
+#pragma omp target update to( \
+ s1.arr[s1.offset : (s1.len - s1.offset) / 2 : 2], \
+ s2.arr[s2.offset : (s2.len - s2.offset) / 2 : 2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < s1.len; i++) {
+ s1.arr[i] += 100;
+ }
+ for (int i = 0; i < s2.len; i++) {
+ s2.arr[i] += 100;
+ }
+ }
+ }
+
+ printf("Test 2 - complex count with offset (to):\n");
+ printf("s1 results:\n");
+ for (int i = 0; i < s1.len; i++)
+ printf("%f\n", s1.arr[i]);
+
+ printf("s2 results:\n");
+ for (int i = 0; i < s2.len; i++)
+ printf("%f\n", s2.arr[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1 - complex count expressions (from):
+// CHECK: s1 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+// CHECK: s2 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 40.000000
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 70.000000
+// CHECK-NEXT: 80.000000
+// CHECK-NEXT: 90.000000
+// CHECK: Test 1 - complex count expressions (to):
+// CHECK: s1 results:
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 16.000000
+// CHECK-NEXT: 18.000000
+// CHECK: s2 results:
+// CHECK-NEXT: 50.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 51.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 52.000000
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 53.000000
+// CHECK-NEXT: 140.000000
+// CHECK-NEXT: 160.000000
+// CHECK-NEXT: 180.000000
+// CHECK: Test 2 - complex count with offset (from):
+// CHECK: s1 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 18.000000
+// CHECK: s2 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 40.000000
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 140.000000
+// CHECK-NEXT: 80.000000
+// CHECK-NEXT: 90.000000
+// CHECK: Test 2 - complex count with offset (to):
+// CHECK: s1 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 18.000000
+// CHECK: s2 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 50.000000
+// CHECK-NEXT: 40.000000
+// CHECK-NEXT: 51.000000
+// CHECK-NEXT: 80.000000
+// CHECK-NEXT: 52.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 53.000000
+// CHECK-NEXT: 160.000000
+// CHECK-NEXT: 180.000000
diff --git a/offload/test/offloading/strided_update_count_expression_misc.c b/offload/test/offloading/strided_update_count_expression_misc.c
new file mode 100644
index 0000000000000..0e93a6d7df2cb
--- /dev/null
+++ b/offload/test/offloading/strided_update_count_expression_misc.c
@@ -0,0 +1,99 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Miscellaneous tests for count expressions: tests modulo, large stride with
+// computed count, and boundary calculations to ensure expression semantics work
+// correctly.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // ====================================================================
+ // TEST 1: Modulo operation in count expression
+ // ====================================================================
+
+ int len1 = 10;
+ int divisor = 5;
+ double data1[len1];
+
+#pragma omp target map(tofrom : len1, divisor, data1[0 : len1])
+ {
+ for (int i = 0; i < len1; i++) {
+ data1[i] = i;
+ }
+ }
+
+#pragma omp target data map(to : len1, divisor, data1[0 : len1])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < len1; i++) {
+ data1[i] += i;
+ }
+ }
+
+ // data[0:10%5:2] = data[0:0:2] updates no indices (count=0)
+#pragma omp target update from(data1[0 : len1 % divisor : 2])
+ }
+
+ printf("Test 1: Modulo count expression\n");
+ for (int i = 0; i < len1; i++)
+ printf("%f\n", data1[i]);
+
+ // ====================================================================
+ // TEST 2: Large stride with computed count for boundary coverage
+ // ====================================================================
+
+ int len2 = 10;
+ int stride = 5;
+ double data2[len2];
+
+#pragma omp target map(tofrom : len2, stride, data2[0 : len2])
+ {
+ for (int i = 0; i < len2; i++) {
+ data2[i] = i;
+ }
+ }
+
+#pragma omp target data map(to : len2, stride, data2[0 : len2])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < len2; i++) {
+ data2[i] += i;
+ }
+ }
+
+ // data[0:(10+5-1)/5:5] = data[0:2:5] updates indices: 0, 5
+#pragma omp target update from(data2[0 : (len2 + stride - 1) / stride : stride])
+ }
+
+ printf("\nTest 2: Large stride count expression\n");
+ for (int i = 0; i < len2; i++)
+ printf("%f\n", data2[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1: Modulo count expression
+// 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: Test 2: Large stride count expression
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
diff --git a/offload/test/offloading/strided_update_multiple_arrays_count_expression.c b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
new file mode 100644
index 0000000000000..9449baa663f67
--- /dev/null
+++ b/offload/test/offloading/strided_update_multiple_arrays_count_expression.c
@@ -0,0 +1,161 @@
+// This test checks "update from" and "update to" with multiple arrays and
+// variable count expressions. Tests both: (1) multiple arrays in single update
+// clause with different count expressions, and (2) overlapping updates to the
+// same array with various count expressions.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int n1 = 10, n2 = 10;
+ double arr1[n1], arr2[n2];
+
+ // ====================================================================
+ // TEST 1: Update FROM - Multiple arrays in single update clause
+ // ====================================================================
+
+#pragma omp target map(tofrom : n1, n2, arr1[0 : n1], arr2[0 : n2])
+ {
+ for (int i = 0; i < n1; i++) {
+ arr1[i] = i;
+ }
+ for (int i = 0; i < n2; i++) {
+ arr2[i] = i * 10;
+ }
+ }
+
+ printf("Test 1: Update FROM - Multiple arrays\n");
+
+#pragma omp target data map(to : n1, n2, arr1[0 : n1], arr2[0 : n2])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < n1; i++) {
+ arr1[i] += i;
+ }
+ for (int i = 0; i < n2; i++) {
+ arr2[i] += 100;
+ }
+ }
+
+ // Update with different count expressions in single clause:
+ // arr1[0:n1/2:2] = arr1[0:5:2] updates indices 0,2,4,6,8
+ // arr2[0:n2/5:2] = arr2[0:2:2] updates indices 0,2
+#pragma omp target update from(arr1[0 : n1 / 2 : 2], arr2[0 : n2 / 5 : 2])
+ }
+
+ printf("from target arr1 results:\n");
+ for (int i = 0; i < n1; i++)
+ printf("%f\n", arr1[i]);
+
+ printf("\nfrom target arr2 results:\n");
+ for (int i = 0; i < n2; i++)
+ printf("%f\n", arr2[i]);
+
+ // ====================================================================
+ // TEST 2: Update TO - Multiple arrays in single update clause
+ // ====================================================================
+
+ for (int i = 0; i < n1; i++) {
+ arr1[i] = i;
+ }
+ for (int i = 0; i < n2; i++) {
+ arr2[i] = i * 10;
+ }
+
+ printf("\nTest 2: Update TO - Multiple arrays\n");
+
+#pragma omp target data map(tofrom : n1, n2, arr1[0 : n1], arr2[0 : n2])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < n1; i++) {
+ arr1[i] = 100.0;
+ }
+ for (int i = 0; i < n2; i++) {
+ arr2[i] = 20.0;
+ }
+ }
+
+ // Modify host
+ for (int i = 0; i < n1; i += 2) {
+ arr1[i] = 10.0;
+ }
+ for (int i = 0; i < n2; i += 2) {
+ arr2[i] = 5.0;
+ }
+
+#pragma omp target update to(arr1[0 : n1 / 2 : 2], arr2[0 : n2 / 5 : 2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < n1; i++) {
+ arr1[i] += 2.0;
+ }
+ for (int i = 0; i < n2; i++) {
+ arr2[i] += 2.0;
+ }
+ }
+ }
+
+ printf("device arr1 values after update to:\n");
+ for (int i = 0; i < n1; i++)
+ printf("%f\n", arr1[i]);
+
+ printf("\ndevice arr2 values after update to:\n");
+ for (int i = 0; i < n2; i++)
+ printf("%f\n", arr2[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1: Update FROM - Multiple arrays
+// CHECK: from target arr1 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: from target arr2 results:
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 30.000000
+// CHECK-NEXT: 40.000000
+// CHECK-NEXT: 50.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 70.000000
+// CHECK-NEXT: 80.000000
+// CHECK-NEXT: 90.000000
+
+// CHECK: Test 2: Update TO - Multiple arrays
+// CHECK: device arr1 values after update to:
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+
+// CHECK: device arr2 values after update to:
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 22.000000
diff --git a/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
new file mode 100644
index 0000000000000..68c3eca4ccc56
--- /dev/null
+++ b/offload/test/offloading/strided_update_multiple_arrays_variable_stride.c
@@ -0,0 +1,145 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests multiple arrays with different variable strides in single update
+// clause.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int stride1 = 2;
+ int stride2 = 2;
+ double data1[10], data2[10];
+
+ // ====================================================================
+ // TEST 1: Update FROM - Multiple arrays with variable strides
+ // ====================================================================
+
+#pragma omp target map(tofrom : stride1, stride2, data1[0 : 10], data2[0 : 10])
+ {
+ for (int i = 0; i < 10; i++) {
+ data1[i] = i;
+ data2[i] = i * 10;
+ }
+ }
+
+ printf("Test 1: Update FROM - Multiple arrays\n");
+
+#pragma omp target data map(to : stride1, stride2, data1[0 : 10], data2[0 : 10])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data1[i] += i;
+ data2[i] += 100;
+ }
+ }
+
+#pragma omp target update from(data1[0 : 5 : stride1], data2[0 : 5 : stride2])
+ }
+
+ printf("from target data1:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data1[i]);
+
+ printf("\nfrom target data2:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data2[i]);
+
+ // ====================================================================
+ // TEST 2: Update TO - Multiple arrays with variable strides
+ // ====================================================================
+
+ for (int i = 0; i < 10; i++) {
+ data1[i] = i;
+ data2[i] = i * 10;
+ }
+
+ printf("\nTest 2: Update TO - Multiple arrays\n");
+
+#pragma omp target data map(tofrom : stride1, stride2, data1[0 : 10], \
+ data2[0 : 10])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data1[i] = 100.0;
+ data2[i] = 20.0;
+ }
+ }
+
+ for (int i = 0; i < 10; i += 2) {
+ data1[i] = 10.0;
+ data2[i] = 5.0;
+ }
+
+#pragma omp target update to(data1[0 : 5 : stride1], data2[0 : 5 : stride2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data1[i] += 2.0;
+ data2[i] += 2.0;
+ }
+ }
+ }
+
+ printf("device data1 after update to:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data1[i]);
+
+ printf("\ndevice data2 after update to:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data2[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1: Update FROM - Multiple arrays
+// CHECK: from target data1:
+// 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: from target data2:
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 30.000000
+// CHECK-NEXT: 140.000000
+// CHECK-NEXT: 50.000000
+// CHECK-NEXT: 160.000000
+// CHECK-NEXT: 70.000000
+// CHECK-NEXT: 180.000000
+// CHECK-NEXT: 90.000000
+
+// CHECK: Test 2: Update TO - Multiple arrays
+// CHECK: device data1 after update to:
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 102.000000
+
+// CHECK: device data2 after update to:
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 22.000000
diff --git a/offload/test/offloading/strided_update_variable_count_and_stride.c b/offload/test/offloading/strided_update_variable_count_and_stride.c
new file mode 100644
index 0000000000000..36056ab64250a
--- /dev/null
+++ b/offload/test/offloading/strided_update_variable_count_and_stride.c
@@ -0,0 +1,136 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests combining variable count expression AND variable stride in array
+// sections.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int len = 10;
+ int stride = 2;
+ double data[len];
+
+ // ====================================================================
+ // TEST 1: Update FROM - Variable count and stride
+ // ====================================================================
+
+#pragma omp target map(tofrom : len, stride, data[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ }
+ }
+
+ printf("Test 1: Update FROM - Variable count and stride\n");
+ printf("original values:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+#pragma omp target data map(to : len, stride, data[0 : len])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] += i;
+ }
+ }
+
+#pragma omp target update from(data[0 : len / 2 : stride])
+ }
+
+ printf("from target results:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+ // ====================================================================
+ // TEST 2: Update TO - Variable count and stride
+ // ====================================================================
+
+ for (int i = 0; i < len; i++) {
+ data[i] = i;
+ }
+
+ printf("\nTest 2: Update TO - Variable count and stride\n");
+ printf("original values:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+#pragma omp target data map(tofrom : len, stride, data[0 : len])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] = 50.0;
+ }
+ }
+
+ for (int i = 0; i < len / 2; i++) {
+ data[i * stride] = 10.0;
+ }
+
+#pragma omp target update to(data[0 : len / 2 : stride])
+
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ data[i] += 5.0;
+ }
+ }
+ }
+
+ printf("device values after update to:\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", data[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1: Update FROM - Variable count and stride
+// CHECK: original 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: from target 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: Test 2: Update TO - Variable count and stride
+// CHECK: original 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: device values after update to:
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
diff --git a/offload/test/offloading/strided_update_variable_stride.c b/offload/test/offloading/strided_update_variable_stride.c
new file mode 100644
index 0000000000000..94723d91734a6
--- /dev/null
+++ b/offload/test/offloading/strided_update_variable_stride.c
@@ -0,0 +1,135 @@
+// This test checks "update from" and "update to" with variable stride.
+// Tests data[0:5:stride] where stride is a variable, making it non-contiguous.
+
+// RUN: %libomptarget-compile-run-and-check-generic
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int stride = 2;
+ double data[10];
+
+ // ====================================================================
+ // TEST 1: Update FROM device (device -> host)
+ // ====================================================================
+
+#pragma omp target map(tofrom : stride, data[0 : 10])
+ {
+ for (int i = 0; i < 10; i++) {
+ data[i] = i;
+ }
+ }
+
+ printf("Test 1: Update FROM device\n");
+ printf("original values:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data[i]);
+
+#pragma omp target data map(to : stride, data[0 : 10])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data[i] += i;
+ }
+ }
+
+#pragma omp target update from(data[0 : 5 : stride])
+ }
+
+ printf("from target results:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data[i]);
+
+ // ====================================================================
+ // TEST 2: Update TO device (host -> device)
+ // ====================================================================
+
+ for (int i = 0; i < 10; i++) {
+ data[i] = i;
+ }
+
+ printf("\nTest 2: Update TO device\n");
+ printf("original values:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data[i]);
+
+#pragma omp target data map(tofrom : stride, data[0 : 10])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data[i] = 50.0;
+ }
+ }
+
+ for (int i = 0; i < 10; i += 2) {
+ data[i] = 10.0;
+ }
+
+#pragma omp target update to(data[0 : 5 : stride])
+
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data[i] += 5.0;
+ }
+ }
+ }
+
+ printf("device values after update to:\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1: Update FROM device
+// CHECK: original 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: from target 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: Test 2: Update TO device
+// CHECK: original 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: device values after update to:
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 55.000000
diff --git a/offload/test/offloading/strided_update_variable_stride_complex.c b/offload/test/offloading/strided_update_variable_stride_complex.c
new file mode 100644
index 0000000000000..3c9857ec22178
--- /dev/null
+++ b/offload/test/offloading/strided_update_variable_stride_complex.c
@@ -0,0 +1,293 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests complex variable stride patterns with multiple arrays and offsets.
+
+#include <omp.h>
+#include <stdio.h>
+
+struct Data {
+ int offset;
+ int stride;
+ double arr[20];
+};
+
+int main() {
+ struct Data d1, d2;
+ int len1 = 10;
+ int len2 = 10;
+
+ // Test 1: Complex stride expressions
+ int base_stride = 1;
+ int multiplier = 2;
+ d1.stride = 2;
+ d2.stride = 3;
+
+ // Initialize on device
+#pragma omp target map(tofrom : d1, d2, base_stride, multiplier)
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] = i * 3;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] = i * 30;
+ }
+ }
+
+ // Test FROM: Complex stride expressions
+#pragma omp target data map(to : d1, d2, base_stride, multiplier)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] += i * 3;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] += i * 30;
+ }
+ }
+
+ // Stride expressions: base_stride*multiplier and (d2.stride+1)/2
+#pragma omp target update from(d1.arr[0 : 5 : base_stride * multiplier], \
+ d2.arr[0 : 3 : (d2.stride + 1) / 2])
+ }
+
+ printf("Test 1 - complex stride expressions (from):\n");
+ printf("d1 results (stride=%d*%d=%d):\n", base_stride, multiplier,
+ base_stride * multiplier);
+ for (int i = 0; i < len1; i++)
+ printf("%f\n", d1.arr[i]);
+
+ printf("d2 results (stride=(%d+1)/2=%d):\n", d2.stride, (d2.stride + 1) / 2);
+ for (int i = 0; i < len2; i++)
+ printf("%f\n", d2.arr[i]);
+
+ // Reset for TO test
+#pragma omp target map(tofrom : d1, d2)
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] = i * 4;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] = i * 40;
+ }
+ }
+
+ // Modify host data with stride expressions
+ int stride1 = base_stride * multiplier;
+ int stride2 = (d2.stride + 1) / 2;
+ for (int i = 0; i < 5; i++) {
+ d1.arr[i * stride1] = i + 200;
+ }
+ for (int i = 0; i < 3; i++) {
+ d2.arr[i * stride2] = i + 150;
+ }
+
+ // Test TO: Update with complex stride expressions
+#pragma omp target data map(to : d1, d2, base_stride, multiplier)
+ {
+#pragma omp target update to(d1.arr[0 : 5 : base_stride * multiplier], \
+ d2.arr[0 : 3 : (d2.stride + 1) / 2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] += 200;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] += 200;
+ }
+ }
+ }
+
+ printf("Test 1 - complex stride expressions (to):\n");
+ printf("d1 results (stride=%d*%d=%d):\n", base_stride, multiplier,
+ base_stride * multiplier);
+ for (int i = 0; i < len1; i++)
+ printf("%f\n", d1.arr[i]);
+
+ printf("d2 results (stride=(%d+1)/2=%d):\n", d2.stride, (d2.stride + 1) / 2);
+ for (int i = 0; i < len2; i++)
+ printf("%f\n", d2.arr[i]);
+
+ // Test 2: Variable stride with non-zero offset
+ d1.offset = 2;
+ d1.stride = 2;
+ d2.offset = 1;
+ d2.stride = 2;
+
+ // Initialize on device
+#pragma omp target map(tofrom : d1, d2, len1, len2)
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] = i;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] = i * 10;
+ }
+ }
+
+ // Test FROM: Variable stride with offset
+#pragma omp target data map(to : d1, d2, len1, len2)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] += i;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] += i * 10;
+ }
+ }
+
+#pragma omp target update from(d1.arr[d1.offset : 4 : d1.stride], \
+ d2.arr[d2.offset : 4 : d2.stride])
+ }
+
+ printf("Test 2 - variable stride with offset (from):\n");
+ printf("d1 results:\n");
+ for (int i = 0; i < len1; i++)
+ printf("%f\n", d1.arr[i]);
+
+ printf("d2 results:\n");
+ for (int i = 0; i < len2; i++)
+ printf("%f\n", d2.arr[i]);
+
+ // Reset for TO test
+#pragma omp target map(tofrom : d1, d2)
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] = i * 2;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] = i * 20;
+ }
+ }
+
+ // Modify host data
+ for (int i = 0; i < 4; i++) {
+ d1.arr[d1.offset + i * d1.stride] = i + 100;
+ }
+ for (int i = 0; i < 4; i++) {
+ d2.arr[d2.offset + i * d2.stride] = i + 50;
+ }
+
+ // Test TO: Update with variable stride and offset
+#pragma omp target data map(to : d1, d2)
+ {
+#pragma omp target update to(d1.arr[d1.offset : 4 : d1.stride], \
+ d2.arr[d2.offset : 4 : d2.stride])
+
+#pragma omp target
+ {
+ for (int i = 0; i < len1; i++) {
+ d1.arr[i] += 100;
+ }
+ for (int i = 0; i < len2; i++) {
+ d2.arr[i] += 100;
+ }
+ }
+ }
+
+ printf("Test 2 - variable stride with offset (to):\n");
+ printf("d1 results:\n");
+ for (int i = 0; i < len1; i++)
+ printf("%f\n", d1.arr[i]);
+
+ printf("d2 results:\n");
+ for (int i = 0; i < len2; i++)
+ printf("%f\n", d2.arr[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1 - complex stride expressions (from):
+// CHECK: d1 results (stride=1*2=2):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 18.000000
+// CHECK-NEXT: 24.000000
+// CHECK-NEXT: 15.000000
+// CHECK-NEXT: 18.000000
+// CHECK-NEXT: 21.000000
+// CHECK-NEXT: 24.000000
+// CHECK-NEXT: 27.000000
+// CHECK: d2 results (stride=(3+1)/2=2):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 90.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 150.000000
+// CHECK-NEXT: 180.000000
+// CHECK-NEXT: 210.000000
+// CHECK-NEXT: 240.000000
+// CHECK-NEXT: 270.000000
+// CHECK: Test 1 - complex stride expressions (to):
+// CHECK: d1 results (stride=1*2=2):
+// CHECK-NEXT: 200.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 201.000000
+// CHECK-NEXT: 12.000000
+// CHECK-NEXT: 202.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 203.000000
+// CHECK-NEXT: 28.000000
+// CHECK-NEXT: 204.000000
+// CHECK-NEXT: 36.000000
+// CHECK: d2 results (stride=(3+1)/2=2):
+// CHECK-NEXT: 150.000000
+// CHECK-NEXT: 40.000000
+// CHECK-NEXT: 151.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 152.000000
+// CHECK-NEXT: 200.000000
+// CHECK-NEXT: 240.000000
+// CHECK-NEXT: 280.000000
+// CHECK-NEXT: 320.000000
+// CHECK-NEXT: 360.000000
+// CHECK: Test 2 - variable stride with offset (from):
+// CHECK: d1 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 18.000000
+// CHECK: d2 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 20.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 40.000000
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 60.000000
+// CHECK-NEXT: 140.000000
+// CHECK-NEXT: 80.000000
+// CHECK-NEXT: 90.000000
+// CHECK: Test 2 - variable stride with offset (to):
+// CHECK: d1 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 18.000000
+// CHECK: d2 results:
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 50.000000
+// CHECK-NEXT: 40.000000
+// CHECK-NEXT: 51.000000
+// CHECK-NEXT: 80.000000
+// CHECK-NEXT: 52.000000
+// CHECK-NEXT: 120.000000
+// CHECK-NEXT: 53.000000
+// CHECK-NEXT: 160.000000
+// CHECK-NEXT: 180.000000
diff --git a/offload/test/offloading/strided_update_variable_stride_misc.c b/offload/test/offloading/strided_update_variable_stride_misc.c
new file mode 100644
index 0000000000000..d27ae0123bfa8
--- /dev/null
+++ b/offload/test/offloading/strided_update_variable_stride_misc.c
@@ -0,0 +1,94 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Miscellaneous variable stride tests: stride=1, stride=array_size, stride from
+// array subscript.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // ====================================================================
+ // TEST 1: Variable stride = 1 (contiguous, but detected as variable)
+ // ====================================================================
+
+ int stride_one = 1;
+ double data1[10];
+
+#pragma omp target map(tofrom : stride_one, data1[0 : 10])
+ {
+ for (int i = 0; i < 10; i++) {
+ data1[i] = i;
+ }
+ }
+
+#pragma omp target data map(to : stride_one, data1[0 : 10])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data1[i] += i;
+ }
+ }
+
+#pragma omp target update from(data1[0 : 10 : stride_one])
+ }
+
+ printf("Test 1: Variable stride = 1\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data1[i]);
+
+ // ====================================================================
+ // TEST 2: Variable stride = array size (only 2 elements)
+ // ====================================================================
+
+ int stride_large = 5;
+ double data2[10];
+
+#pragma omp target map(tofrom : stride_large, data2[0 : 10])
+ {
+ for (int i = 0; i < 10; i++) {
+ data2[i] = i;
+ }
+ }
+
+#pragma omp target data map(to : stride_large, data2[0 : 10])
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < 10; i++) {
+ data2[i] += i;
+ }
+ }
+
+#pragma omp target update from(data2[0 : 2 : stride_large])
+ }
+
+ printf("\nTest 2: Variable stride = 5\n");
+ for (int i = 0; i < 10; i++)
+ printf("%f\n", data2[i]);
+
+ return 0;
+}
+
+// CHECK: Test 1: Variable stride = 1
+// 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: Test 2: Variable stride = 5
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 1.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 3.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 7.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
diff --git a/offload/test/offloading/target_update_ptr_count_expression.c b/offload/test/offloading/target_update_ptr_count_expression.c
new file mode 100644
index 0000000000000..c4b9fd566d401
--- /dev/null
+++ b/offload/test/offloading/target_update_ptr_count_expression.c
@@ -0,0 +1,99 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests non-contiguous array sections with expression-based count on
+// heap-allocated pointer arrays with both FROM and TO directives.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ int len = 10;
+ double *result = (double *)malloc(len * sizeof(double));
+
+ // Initialize host array to zero
+ for (int i = 0; i < len; i++) {
+ result[i] = 0;
+ }
+
+ // Initialize on device
+#pragma omp target enter data map(to : len, result[0 : len])
+
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] = i;
+ }
+ }
+
+ // Test FROM: Modify on device, then update from device
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] += i * 10;
+ }
+ }
+
+ // Update from device with expression-based count: len/2 elements
+#pragma omp target update from(result[0 : len / 2 : 2])
+
+ printf("heap ptr count expression (from):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", result[i]);
+
+ // Test TO: Reset, modify host, update to device
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] = i * 2;
+ }
+ }
+
+ // Modify host data
+ for (int i = 0; i < len / 2; i++) {
+ result[i * 2] = i + 100;
+ }
+
+ // Update to device with expression-based count
+#pragma omp target update to(result[0 : len / 2 : 2])
+
+ // Read back full array
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] += 100;
+ }
+ }
+
+#pragma omp target update from(result[0 : len])
+
+ printf("heap ptr count expression (to):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", result[i]);
+
+#pragma omp target exit data map(delete : len, result[0 : len])
+ free(result);
+ return 0;
+}
+
+// CHECK: heap ptr count expression (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 44.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 66.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 88.000000
+// CHECK-NEXT: 0.000000
+// CHECK: heap ptr count expression (to):
+// CHECK-NEXT: 200.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 201.000000
+// CHECK-NEXT: 106.000000
+// CHECK-NEXT: 202.000000
+// CHECK-NEXT: 110.000000
+// CHECK-NEXT: 203.000000
+// CHECK-NEXT: 114.000000
+// CHECK-NEXT: 204.000000
+// CHECK-NEXT: 118.000000
diff --git a/offload/test/offloading/target_update_ptr_variable_count_and_stride.c b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c
new file mode 100644
index 0000000000000..1a28595969c69
--- /dev/null
+++ b/offload/test/offloading/target_update_ptr_variable_count_and_stride.c
@@ -0,0 +1,94 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests heap-allocated pointers with both variable count expression and
+// variable stride.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ int len = 10;
+ int stride = 2;
+ double *result = (double *)malloc(len * sizeof(double));
+
+ for (int i = 0; i < len; i++) {
+ result[i] = 0;
+ }
+
+#pragma omp target enter data map(to : len, stride, result[0 : len])
+
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] = i;
+ }
+ }
+
+ // Test FROM: Variable count and stride
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] += i * 10;
+ }
+ }
+
+#pragma omp target update from(result[0 : len / 2 : stride])
+
+ printf("heap ptr variable count and stride (from):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", result[i]);
+
+ // Test TO: Reset, modify host, update to device
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] = i * 2;
+ }
+ }
+
+ for (int i = 0; i < len / 2; i++) {
+ result[i * stride] = i + 100;
+ }
+
+#pragma omp target update to(result[0 : len / 2 : stride])
+
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] += 100;
+ }
+ }
+
+#pragma omp target update from(result[0 : len])
+
+ printf("heap ptr variable count and stride (to):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", result[i]);
+
+#pragma omp target exit data map(delete : len, stride, result[0 : len])
+ free(result);
+ return 0;
+}
+
+// CHECK: heap ptr variable count and stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 44.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 66.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 88.000000
+// CHECK-NEXT: 0.000000
+// CHECK: heap ptr variable count and stride (to):
+// CHECK-NEXT: 200.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 201.000000
+// CHECK-NEXT: 106.000000
+// CHECK-NEXT: 202.000000
+// CHECK-NEXT: 110.000000
+// CHECK-NEXT: 203.000000
+// CHECK-NEXT: 114.000000
+// CHECK-NEXT: 204.000000
+// CHECK-NEXT: 118.000000
diff --git a/offload/test/offloading/target_update_ptr_variable_stride.c b/offload/test/offloading/target_update_ptr_variable_stride.c
new file mode 100644
index 0000000000000..bea396065b760
--- /dev/null
+++ b/offload/test/offloading/target_update_ptr_variable_stride.c
@@ -0,0 +1,95 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests non-contiguous array sections with variable stride on heap-allocated
+// pointers.
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ int stride = 2;
+ int len = 10;
+ double *result = (double *)malloc(len * sizeof(double));
+
+ // Initialize
+ for (int i = 0; i < len; i++) {
+ result[i] = 0;
+ }
+
+#pragma omp target enter data map(to : stride, len, result[0 : len])
+
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] = i;
+ }
+ }
+
+ // Test FROM
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] += i * 10;
+ }
+ }
+
+#pragma omp target update from(result[0 : 5 : stride])
+
+ printf("heap ptr variable stride (from):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", result[i]);
+
+ // Test TO: Reset, modify host, update to device
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] = i * 2;
+ }
+ }
+
+ for (int i = 0; i < 5; i++) {
+ result[i * stride] = i + 100;
+ }
+
+#pragma omp target update to(result[0 : 5 : stride])
+
+#pragma omp target map(alloc : result[0 : len])
+ {
+ for (int i = 0; i < len; i++) {
+ result[i] += 100;
+ }
+ }
+
+#pragma omp target update from(result[0 : len])
+
+ printf("heap ptr variable stride (to):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", result[i]);
+
+#pragma omp target exit data map(delete : stride, len, result[0 : len])
+ free(result);
+ return 0;
+}
+
+// CHECK: heap ptr variable stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 22.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 44.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 66.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 88.000000
+// CHECK-NEXT: 0.000000
+// CHECK: heap ptr variable stride (to):
+// CHECK-NEXT: 200.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 201.000000
+// CHECK-NEXT: 106.000000
+// CHECK-NEXT: 202.000000
+// CHECK-NEXT: 110.000000
+// CHECK-NEXT: 203.000000
+// CHECK-NEXT: 114.000000
+// CHECK-NEXT: 204.000000
+// CHECK-NEXT: 118.000000
diff --git a/offload/test/offloading/target_update_strided_struct_count_expression.c b/offload/test/offloading/target_update_strided_struct_count_expression.c
new file mode 100644
index 0000000000000..1c1fd005c405f
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_count_expression.c
@@ -0,0 +1,97 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests non-contiguous array sections with expression-based count on struct
+// member arrays with both FROM and TO directives.
+
+#include <omp.h>
+#include <stdio.h>
+
+struct S {
+ int len;
+ double data[20];
+};
+
+int main() {
+ struct S s;
+ s.len = 10;
+
+ // Initialize on device
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = i;
+ }
+ }
+
+ // Test FROM: Modify on device, then update from device
+#pragma omp target data map(to : s)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += i * 10;
+ }
+ }
+
+ // Update from device with expression-based count: len/2 elements
+#pragma omp target update from(s.data[0 : s.len / 2 : 2])
+ }
+
+ printf("struct count expression (from):\n");
+ for (int i = 0; i < s.len; i++)
+ printf("%f\n", s.data[i]);
+
+ // Test TO: Reset, modify host, update to device
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = i * 2;
+ }
+ }
+
+ // Modify host data
+ for (int i = 0; i < s.len / 2; i++) {
+ s.data[i * 2] = i + 100;
+ }
+
+ // Update to device with expression-based count
+#pragma omp target data map(to : s)
+ {
+#pragma omp target update to(s.data[0 : s.len / 2 : 2])
+
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += 100;
+ }
+ }
+ }
+
+ printf("struct count expression (to):\n");
+ for (int i = 0; i < s.len; i++)
+ printf("%f\n", s.data[i]);
+
+ return 0;
+}
+
+// CHECK: struct count expression (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 11.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 33.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 77.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+// CHECK: struct count expression (to):
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 104.000000
+// CHECK-NEXT: 18.000000
diff --git a/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c
new file mode 100644
index 0000000000000..6daf10383e921
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_variable_count_and_stride.c
@@ -0,0 +1,96 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests struct member arrays with both variable count expression and variable
+// stride.
+
+#include <omp.h>
+#include <stdio.h>
+
+struct S {
+ int len;
+ int stride;
+ double data[20];
+};
+
+int main() {
+ struct S s;
+ s.len = 10;
+ s.stride = 2;
+
+ // Initialize
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = i;
+ }
+ }
+
+ // Test FROM: Variable count and stride
+#pragma omp target data map(to : s)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += i * 10;
+ }
+ }
+
+#pragma omp target update from(s.data[0 : s.len / 2 : s.stride])
+ }
+
+ printf("struct variable count and stride (from):\n");
+ for (int i = 0; i < s.len; i++)
+ printf("%f\n", s.data[i]);
+
+ // Test TO: Reset, modify host, update to device
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] = i * 2;
+ }
+ }
+
+ for (int i = 0; i < s.len / 2; i++) {
+ s.data[i * s.stride] = i + 100;
+ }
+
+#pragma omp target data map(to : s)
+ {
+#pragma omp target update to(s.data[0 : s.len / 2 : s.stride])
+
+#pragma omp target
+ {
+ for (int i = 0; i < s.len; i++) {
+ s.data[i] += 100;
+ }
+ }
+ }
+
+ printf("struct variable count and stride (to):\n");
+ for (int i = 0; i < s.len; i++)
+ printf("%f\n", s.data[i]);
+
+ return 0;
+}
+
+// CHECK: struct variable count and stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 11.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 33.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 77.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+// CHECK: struct variable count and stride (to):
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 104.000000
+// CHECK-NEXT: 18.000000
diff --git a/offload/test/offloading/target_update_strided_struct_variable_stride.c b/offload/test/offloading/target_update_strided_struct_variable_stride.c
new file mode 100644
index 0000000000000..4cd9da629ca93
--- /dev/null
+++ b/offload/test/offloading/target_update_strided_struct_variable_stride.c
@@ -0,0 +1,95 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+// Tests non-contiguous array sections with variable stride on struct member
+// arrays.
+
+#include <omp.h>
+#include <stdio.h>
+
+struct S {
+ int stride;
+ double data[20];
+};
+
+int main() {
+ struct S s;
+ s.stride = 2;
+ int len = 10;
+
+ // Initialize
+#pragma omp target map(tofrom : s, len)
+ {
+ for (int i = 0; i < len; i++) {
+ s.data[i] = i;
+ }
+ }
+
+ // Test FROM
+#pragma omp target data map(to : s, len)
+ {
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ s.data[i] += i * 10;
+ }
+ }
+
+#pragma omp target update from(s.data[0 : 5 : s.stride])
+ }
+
+ printf("struct variable stride (from):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", s.data[i]);
+
+ // Test TO: Reset, modify host, update to device
+#pragma omp target map(tofrom : s)
+ {
+ for (int i = 0; i < len; i++) {
+ s.data[i] = i * 2;
+ }
+ }
+
+ for (int i = 0; i < 5; i++) {
+ s.data[i * s.stride] = i + 100;
+ }
+
+#pragma omp target data map(to : s)
+ {
+#pragma omp target update to(s.data[0 : 5 : s.stride])
+
+#pragma omp target
+ {
+ for (int i = 0; i < len; i++) {
+ s.data[i] += 100;
+ }
+ }
+ }
+
+ printf("struct variable stride (to):\n");
+ for (int i = 0; i < len; i++)
+ printf("%f\n", s.data[i]);
+
+ return 0;
+}
+
+// CHECK: struct variable stride (from):
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 11.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 33.000000
+// CHECK-NEXT: 4.000000
+// CHECK-NEXT: 55.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 77.000000
+// CHECK-NEXT: 8.000000
+// CHECK-NEXT: 9.000000
+// CHECK: struct variable stride (to):
+// CHECK-NEXT: 100.000000
+// CHECK-NEXT: 2.000000
+// CHECK-NEXT: 101.000000
+// CHECK-NEXT: 6.000000
+// CHECK-NEXT: 102.000000
+// CHECK-NEXT: 10.000000
+// CHECK-NEXT: 103.000000
+// CHECK-NEXT: 14.000000
+// CHECK-NEXT: 104.000000
+// CHECK-NEXT: 18.000000
>From d9bca90cd26fbb20320abfa262901d4de660a54d Mon Sep 17 00:00:00 2001
From: amtiwari <amtiwari at amd.com>
Date: Mon, 19 Jan 2026 03:58:37 -0500
Subject: [PATCH 4/4] formatting
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 4 ++--
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 4 ++--
2 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 027d1fb26bc97..e6f5f00a86922 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8033,9 +8033,9 @@ class MappableExprsHandler {
const Expr *StrideExpr = OASE->getStride();
if (!StrideExpr)
return false;
-
+
assert(StrideExpr->getType()->isIntegerType() &&
- "Stride expression must be of integer type");
+ "Stride expression must be of integer type");
// If the stride is a variable (not a constant), it's non-contiguous.
const Expr *S = StrideExpr->IgnoreParenImpCasts();
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 578e63351ca5f..6f44a9649abeb 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9776,8 +9776,8 @@ Error OpenMPIRBuilder::emitOffloadingArrays(
bool IsNonContigEntry =
IsNonContiguous &&
(static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- CombinedInfo.Types[I] &
- OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0);
+ CombinedInfo.Types[I] &
+ OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG) != 0);
// For NON_CONTIG entries ArgSizes must carry the dimension count
// (number of descriptor_dim records) – NOT the byte size expression.
// Variable subsection forms (e.g. 0:s.len/2:2) previously produced a
More information about the cfe-commits
mailing list